[FFmpeg-devel] [PATCH 1/2] libavutil/libavfilter: opencl wrapper based on comments on 20130325
Wei Gao
highgod0401 at gmail.com
Tue Mar 26 04:44:30 CET 2013
Hi,
Stefano Sabatini, thank you for your reply, some questions and explanations
as follows
Thanks
Best regards
2013/3/26 Stefano Sabatini <stefasab at gmail.com>
> On date Monday 2013-03-25 19:00:24 +0800, Wei Gao encoded:
> >
>
> > From fb53b21efb819336e30eb50bc4408b3c80111e82 Mon Sep 17 00:00:00 2001
> > From: highgod0401 <highgod0401 at gmail.com>
> > Date: Mon, 25 Mar 2013 18:52:51 +0800
> > Subject: [PATCH 1/2] opencl wrapper based on comments on 20130325
> >
> > ---
> > configure | 4 +
> > libavutil/Makefile | 3 +
> > libavutil/opencl.c | 967
> +++++++++++++++++++++++++++++++++++++++++++++++++++++
> > libavutil/opencl.h | 219 ++++++++++++
> > 4 files changed, 1193 insertions(+)
> > create mode 100644 libavutil/opencl.c
> > create mode 100644 libavutil/opencl.h
> >
> > diff --git a/configure b/configure
> > index 8443db4..9c42a85 100755
> > --- a/configure
> > +++ b/configure
> > @@ -233,6 +233,7 @@ External library support:
> > --enable-libxvid enable Xvid encoding via xvidcore,
> > native MPEG-4/Xvid encoder exists [no]
> > --enable-openal enable OpenAL 1.1 capture support [no]
> > + --enable-opencl enable OpenCL code
> > --enable-openssl enable openssl [no]
> > --enable-x11grab enable X11 grabbing [no]
> > --enable-zlib enable zlib [autodetect]
> > @@ -1178,6 +1179,7 @@ EXTERNAL_LIBRARY_LIST="
> > libxavs
> > libxvid
> > openal
> > + opencl
> > openssl
> > x11grab
> > zlib
> > @@ -3982,6 +3984,7 @@ enabled openal && { { for al_libs in
> "${OPENAL_LIBS}" "-lopenal" "-lOpenAL32
> > die "ERROR: openal not found"; } &&
> > { check_cpp_condition "AL/al.h"
> "defined(AL_VERSION_1_1)" ||
> > die "ERROR: openal must be installed and
> version must be 1.1 or compatible"; }
> > +enabled opencl && require2 opencl CL/cl.h clEnqueueNDRangeKernel
> -lOpenCL
> > enabled openssl && { check_lib openssl/ssl.h SSL_library_init -lssl
> -lcrypto ||
> > check_lib openssl/ssl.h SSL_library_init
> -lssl32 -leay32 ||
> > check_lib openssl/ssl.h SSL_library_init -lssl
> -lcrypto -lws2_32 -lgdi32 ||
> > @@ -4350,6 +4353,7 @@ echo "network support ${network-no}"
> > echo "threading support ${thread_type-no}"
> > echo "safe bitstream reader ${safe_bitstream_reader-no}"
> > echo "SDL support ${sdl-no}"
> > +echo "opencl enabled ${opencl-no}"
> > echo "texi2html enabled ${texi2html-no}"
> > echo "perl enabled ${perl-no}"
> > echo "pod2man enabled ${pod2man-no}"
> > diff --git a/libavutil/Makefile b/libavutil/Makefile
> > index 103ce5e..6375e10 100644
> > --- a/libavutil/Makefile
> > +++ b/libavutil/Makefile
> > @@ -52,6 +52,8 @@ HEADERS = adler32.h
> \
> >
> > HEADERS-$(CONFIG_LZO) += lzo.h
> >
> > +HEADERS-$(CONFIG_OPENCL) += opencl.h
> > +
> > ARCH_HEADERS = bswap.h
> \
> > intmath.h
> \
> > intreadwrite.h
> \
> > @@ -115,6 +117,7 @@ SKIPHEADERS-$(HAVE_MACHINE_RW_BARRIER) +=
> atomic_suncc.h
> > SKIPHEADERS-$(HAVE_MEMORYBARRIER) += atomic_win32.h
> > SKIPHEADERS-$(HAVE_SYNC_VAL_COMPARE_AND_SWAP) += atomic_gcc.h
> >
> > +OBJS-$(CONFIG_OPENCL) += opencl.o
> > TESTPROGS = adler32
> \
> > aes
> \
> > atomic
> \
> > diff --git a/libavutil/opencl.c b/libavutil/opencl.c
> > new file mode 100644
> > index 0000000..20f07de
> > --- /dev/null
> > +++ b/libavutil/opencl.c
> > @@ -0,0 +1,967 @@
> > +/*
> > + * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
> > + * Copyright (C) 2012 Li Cao <li at multicorewareinc.com>
> > + * Copyright (C) 2012 Wei Gao <weigao at multicorewareinc.com>
> > + *
> > + * This file is part of FFmpeg.
> > + *
> > + * FFmpeg is free software; you can redistribute it and/or
> > + * modify it under the terms of the GNU Lesser General Public
> > + * License as published by the Free Software Foundation; either
> > + * version 2.1 of the License, or (at your option) any later version.
> > + *
> > + * FFmpeg is distributed in the hope that it will be useful,
> > + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
> > + * Lesser General Public License for more details.
> > + *
> > + * You should have received a copy of the GNU Lesser General Public
> > + * License along with FFmpeg; if not, write to the Free Software
> > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
> 02110-1301 USA
> > + */
> > +
> > +#include "opencl.h"
> > +#include "avstring.h"
> > +#include "log.h"
> > +#include "avassert.h"
> > +
>
> > +#define MAX_KERNEL_STRING_LEN 64
>
> MAX_KERNEL_NAME_LEN may be more internally consistent (since it is
> used with kernel_name variable).
>
> > +#define MAX_FILE_NUM 50
> > +#define MAX_KERNEL_NUM 200
>
> > +#define MAX_FILTER_NAME_LEN 64
> > +#define MAX_FILTER_NUM 200
>
> Both unused.
>
> > +#define MAX_KERNEL_SRCFILE_LEN 256
> > +
> > +
> > +typedef struct OpenCLEnv {
> > + cl_platform_id platform;
> > + cl_context context;
> > + cl_device_id devices;
> > + cl_command_queue command_queue;
> > +} OpenCLEnv;
> > +
> > +typedef struct GPUEnv {
> > + cl_platform_id platform;
> > + cl_device_type device_type;
> > + cl_context context;
> > + cl_device_id *device_ids;
> > + cl_device_id device_id;
> > + cl_command_queue command_queue;
>
> OpenCLEnv fields are all contained here, so maybe you could just put
> an OpenCLEnv field here.
>
> > + cl_kernel kernels[MAX_FILE_NUM];
> > + cl_program program;
> > + char kernel_srcfile_name[MAX_KERNEL_SRCFILE_LEN];
>
> > + char kernel_names[MAX_KERNEL_NUM][MAX_KERNEL_STRING_LEN+1];
> > + av_opencl_kernel_function kernel_functions[MAX_KERNEL_NUM];
> > + const char *kernel_code[MAX_KERNEL_NUM];
>
> > + int kernel_count;
> > + int runtime_kernel_count;
> > + int is_user_created; // 1: the opencl env is created by user and
> use AVOpenCLExternalInfo to pass to ffmpeg ,0:created by opencl wrapper
> > + uint8_t *temp_buffer;
> > + int temp_buffer_size;
> > +} GPUEnv;
> > +
> > +typedef struct OpenclErrorMsg {
> > + int err_code;
> > + const char *err_str;
> > +} OpenclErrorMsg;
> > +
> > +static OpenclErrorMsg opencl_err_msg[] = {
> > + {CL_DEVICE_NOT_FOUND, "DEVICE NOT
> FOUND"},
> > + {CL_DEVICE_NOT_AVAILABLE, "DEVICE NOT
> AVAILABLE"},
> > + {CL_COMPILER_NOT_AVAILABLE, "COMPILER
> NOT AVAILABLE"},
> > + {CL_MEM_OBJECT_ALLOCATION_FAILURE, "MEM OBJECT
> ALLOCATION FAILURE"},
> > + {CL_OUT_OF_RESOURCES, "OUT OF
> RESOURCES"},
> > + {CL_OUT_OF_HOST_MEMORY, "OUT OF
> HOST MEMORY"},
> > + {CL_PROFILING_INFO_NOT_AVAILABLE, "PROFILING
> INFO NOT AVAILABLE"},
> > + {CL_MEM_COPY_OVERLAP, "MEM COPY
> OVERLAP"},
> > + {CL_IMAGE_FORMAT_MISMATCH, "IMAGE
> FORMAT MISMATCH"},
> > + {CL_IMAGE_FORMAT_NOT_SUPPORTED, "IMAGE
> FORMAT NOT_SUPPORTED"},
> > + {CL_BUILD_PROGRAM_FAILURE, "BUILD
> PROGRAM FAILURE"},
> > + {CL_MAP_FAILURE, "MAP
> FAILURE"},
> > + {CL_MISALIGNED_SUB_BUFFER_OFFSET, "MISALIGNED
> SUB BUFFER OFFSET"},
> > + {CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST, "EXEC
> STATUS ERROR FOR EVENTS IN WAIT LIST"},
> > + {CL_COMPILE_PROGRAM_FAILURE, "COMPILE
> PROGRAM FAILURE"},
> > + {CL_LINKER_NOT_AVAILABLE, "LINKER NOT
> AVAILABLE"},
> > + {CL_LINK_PROGRAM_FAILURE, "LINK
> PROGRAM FAILURE"},
> > + {CL_DEVICE_PARTITION_FAILED, "DEVICE
> PARTITION FAILED"},
> > + {CL_KERNEL_ARG_INFO_NOT_AVAILABLE, "KERNEL ARG
> INFO NOT AVAILABLE"},
> > + {CL_INVALID_VALUE, "INVALID
> VALUE"},
> > + {CL_INVALID_DEVICE_TYPE, "INVALID
> DEVICE TYPE"},
> > + {CL_INVALID_PLATFORM, "INVALID
> PLATFORM"},
> > + {CL_INVALID_DEVICE, "INVALID
> DEVICE"},
> > + {CL_INVALID_CONTEXT, "INVALID
> CONTEXT"},
> > + {CL_INVALID_QUEUE_PROPERTIES, "INVALID
> QUEUE PROPERTIES"},
> > + {CL_INVALID_COMMAND_QUEUE, "INVALID
> COMMAND QUEUE"},
> > + {CL_INVALID_HOST_PTR, "INVALID
> HOST PTR"},
> > + {CL_INVALID_MEM_OBJECT, "INVALID
> MEM OBJECT"},
> > + {CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, "INVALID
> IMAGE FORMAT DESCRIPTOR"},
> > + {CL_INVALID_IMAGE_SIZE, "INVALID
> IMAGE SIZE"},
> > + {CL_INVALID_SAMPLER, "INVALID
> SAMPLER"},
> > + {CL_INVALID_BINARY, "INVALID
> BINARY"},
> > + {CL_INVALID_BUILD_OPTIONS, "INVALID
> BUILD OPTIONS"},
> > + {CL_INVALID_PROGRAM, "INVALID
> PROGRAM"},
> > + {CL_INVALID_PROGRAM_EXECUTABLE, "INVALID
> PROGRAM EXECUTABLE"},
> > + {CL_INVALID_KERNEL_NAME, "INVALID
> KERNEL NAME"},
> > + {CL_INVALID_KERNEL_DEFINITION, "INVALID
> KERNEL DEFINITION"},
> > + {CL_INVALID_KERNEL, "INVALID
> KERNEL"},
> > + {CL_INVALID_ARG_INDEX, "INVALID
> ARG INDEX"},
> > + {CL_INVALID_ARG_VALUE, "INVALID
> ARG VALUE"},
> > + {CL_INVALID_ARG_SIZE, "INVALID
> ARG_SIZE"},
> > + {CL_INVALID_KERNEL_ARGS, "INVALID
> KERNEL ARGS"},
> > + {CL_INVALID_WORK_DIMENSION, "INVALID
> WORK DIMENSION"},
> > + {CL_INVALID_WORK_GROUP_SIZE, "INVALID
> WORK GROUP SIZE"},
> > + {CL_INVALID_WORK_ITEM_SIZE, "INVALID
> WORK ITEM SIZE"},
> > + {CL_INVALID_GLOBAL_OFFSET, "INVALID
> GLOBAL OFFSET"},
> > + {CL_INVALID_EVENT_WAIT_LIST, "INVALID
> EVENT WAIT LIST"},
> > + {CL_INVALID_EVENT, "INVALID
> EVENT"},
> > + {CL_INVALID_OPERATION, "INVALID
> OPERATION"},
> > + {CL_INVALID_GL_OBJECT, "INVALID GL
> OBJECT"},
> > + {CL_INVALID_BUFFER_SIZE, "INVALID
> BUFFER SIZE"},
> > + {CL_INVALID_MIP_LEVEL, "INVALID
> MIP LEVEL"},
> > + {CL_INVALID_GLOBAL_WORK_SIZE, "INVALID
> GLOBAL WORK SIZE"},
> > + {CL_INVALID_PROPERTY, "INVALID
> PROPERTY"},
> > + {CL_INVALID_IMAGE_DESCRIPTOR, "INVALID
> IMAGE DESCRIPTOR"},
> > + {CL_INVALID_COMPILER_OPTIONS, "INVALID
> COMPILER OPTIONS"},
> > + {CL_INVALID_LINKER_OPTIONS, "INVALID
> LINKER OPTIONS"},
> > + {CL_INVALID_DEVICE_PARTITION_COUNT, "INVALID
> DEVICE PARTITION COUNT"},
> > +};
> > +
>
> > +typedef struct OpenclUtils {
> > + const AVClass *class;
> > + int log_offset;
> > + void *log_ctx;
> > +} OpenclUtils;
>
> As I mentioned on IRC, we could use this context to set some options
> (kernel and log file paths and build options).
>
> > +
> > +static const AVClass openclutils_class = {"OPENCLUTILS",
> av_default_item_name,
> > + NULL,
> LIBAVUTIL_VERSION_INT,
> > +
> offsetof(OpenclUtils, log_offset),
> > +
> offsetof(OpenclUtils, log_ctx)};
> > +static OpenclUtils openclutils = {&openclutils_class};
> > +static GPUEnv gpu_env;
> > +static int isinited;
> > +
> > +int av_opencl_register_kernel(const char *kernel_name, const char
> *kernel_code)
> > +{
> > + if (gpu_env.kernel_count < MAX_KERNEL_NUM) {
> > + if (strlen(kernel_name) <= MAX_KERNEL_STRING_LEN) {
> > + gpu_env.kernel_code[gpu_env.kernel_count] = kernel_code;
> > + av_strlcpy(gpu_env.kernel_names[gpu_env.kernel_count],
> kernel_name, MAX_KERNEL_STRING_LEN+1);
> > + gpu_env.kernel_count++;
> > + } else {
> > + av_log(&openclutils, AV_LOG_ERROR, "Registered kernel name
> %s is too long\n", kernel_name);
> > + return AVERROR(EINVAL);
> > + }
> > + } else {
>
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not register kernel
> with name '%s', maximum number of registered kernels %d already reached\n",
> kernel_name, MAX_KERNEL_NUM);
>
> Nit: split overly long line
>
> > + return AVERROR(EINVAL);
> > + }
> > + return 0;
> > +}
> > +static const char *opencl_errstr(int status)
> > +{
> > + int i;
> > + for (i = 0; i < sizeof(opencl_err_msg); i++) {
> > + if (opencl_err_msg[i].err_code == status)
> > + return opencl_err_msg[i].err_str;
> > + }
> > + return "unknown error";
> > +}
> > +
>
> > +static int access_binaries(cl_device_id *device_ids, int num_devices,
> const char *cl_file_name, FILE **fhandle,
> > + size_t *binary_sizes, char **binaries,
> int write)
> > +{
> > + FILE *fd = NULL;
> > + int status;
> > + char filename[1024] = {0};
> > + char cl_name[1024] = {0};
>
> > + char devicename[1024] = {0};
>
> device_name for readability/consistency
>
> > + int i;
> > + for (i = 0; i < num_devices; i++) {
> > + if (device_ids[i] != 0) {
>
> > + status = clGetDeviceInfo(device_ids[i],
> > + CL_DEVICE_NAME,
> > + sizeof(devicename),
> > + devicename,
> > + NULL);
>
> Style nit: this could be:
> status = clGetDeviceInfo(device_ids[i], CL_DEVICE_NAME,
> sizeof(devicename), devicename, NULL);
>
> and save three lines, which may help readability
> (feel free to ignore comments marked as "nit").
>
> > + if (status == CL_SUCCESS) {
> > + av_strlcpy(cl_name, cl_file_name, sizeof(cl_name));
> > + snprintf(filename, sizeof(filename), "./%s-%s.bin",
> cl_name, devicename);
>
> > + if (!write) {
>
> nit: as a norm better to avoid negative logic (which creates problem
> to the human reader), especially if it costs nothing, you can do:
> if (write) {
> WRITE BLOCK
> } else {
> READ BLOCK
> }
>
> > + fd = fopen(filename, "rb");
> > + if (fd) {
> > + if (fhandle)
> > + *fhandle = fd;
> > + break;
> > + }
> > + } else {
> > + if (binary_sizes[i] != 0) {
> > + fd = fopen(filename, "wb");
> > + if (!fd) {
>
> > + av_log(&openclutils, AV_LOG_ERROR, "Could
> not write file '%s'\n", filename);
> > + return AVERROR(errno);
>
> int err = AVERROR(errno);
> av_log(...);
> return err;
>
> The call to av_log() may perform other syscalls, changing the value of
> the global errno.
>
> > + }
> > + fwrite(binaries[i], sizeof(char),
> binary_sizes[i], fd);
> > + fclose(fd);
> > + }
> > + }
> > + } else {
>
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get
> OpenCL device information: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + }
> > + }
> > + return 0;
> > +}
>
> About the whole logic. This is accessing some files in the local
> path. This should be avoided. We have several choice:
> - we avoid this access, if it can be avoided, in the first version
> - we specify the filepath in a sort of global context
>
> In the latter case we need to understand how to create such global
> context.
>
> can you agree that we just compile every time? I thinks the most important
thing in this step is make the opencl code run.
> > +
> > +static int check_generated_binary(cl_context context, const char
> *cl_file_name, FILE **fhandle)
> > +{
> > + int ret = 0;
> > + cl_int status;
> > + size_t num_devices;
> > + cl_device_id *device_ids;
>
> > + status = clGetContextInfo(context,
> > + CL_CONTEXT_NUM_DEVICES,
> > + sizeof(num_devices),
> > + &num_devices,
> > + NULL);
>
> same note here, you could save three lines here
>
> > + if (status != CL_SUCCESS){
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get OpenCL
> context number of devices: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > +
> > + device_ids = av_malloc(sizeof(cl_device_id) * num_devices);
> > + if (!device_ids)
> > + return AVERROR(ENOMEM);
> > +
> > + /* grab the handles to all of the devices in the context. */
> > + status = clGetContextInfo(context,
> > + CL_CONTEXT_DEVICES,
> > + sizeof(cl_device_id) * num_devices,
> > + device_ids,
> > + NULL);
> > + if (status != CL_SUCCESS){
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get OpenCL
> context devices: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + ret = access_binaries(device_ids, num_devices,cl_file_name,
> fhandle, NULL, NULL, 0);
> > +end:
> > + av_free(device_ids);
> > + return ret;
> > +}
> > +
>
> > +static int generate_bin_from_kernel_source(cl_program program, const
> char *cl_file_name)
> > +{
> > + int i = 0;
> > + cl_int status;
> > + size_t *binary_sizes = NULL;
> > + size_t num_devices;
> > + cl_device_id *device_ids = NULL;
> > + char **binaries = NULL;
> > + int ret = 0;
> > + status = clGetProgramInfo(program,
> > + CL_PROGRAM_NUM_DEVICES,
> > + sizeof(num_devices),
> > + &num_devices,
> > + NULL);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get OpenCL
> program information: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + device_ids = av_mallocz(sizeof(cl_device_id) * num_devices);
> > + if (!device_ids)
> > + return AVERROR(ENOMEM);
> > + /* grab the handles to all of the devices in the program. */
> > + status = clGetProgramInfo(program,
> > + CL_PROGRAM_DEVICES,
> > + sizeof(cl_device_id) * num_devices,
> > + device_ids,
> > + NULL);
> > +
> > + /* figure out the sizes of each of the binaries. */
> > + binary_sizes = av_mallocz(sizeof(size_t) * num_devices);
> > + if (!binary_sizes) {
> > + ret = AVERROR(ENOMEM);
> > + goto end;
> > + }
> > +
> > + status = clGetProgramInfo(program,
> > + CL_PROGRAM_BINARY_SIZES,
> > + sizeof(size_t) * num_devices,
> > + binary_sizes, NULL);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get OpenCL
> program info: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + /* copy over all of the generated binaries. */
> > + binaries = av_mallocz(sizeof(char *) * num_devices);
> > + if (!binaries) {
> > + ret = AVERROR(ENOMEM);
> > + goto end;
> > + }
>
> > + for (i = 0; i < num_devices; i++) {
> > + if (binary_sizes[i] != 0) {
> > + binaries[i] = av_mallocz(sizeof(char) * binary_sizes[i]);
>
> sizeof(char) == 1, same below
>
> > + if (!binaries[i]) {
> > + ret = AVERROR(ENOMEM);
> > + goto end;
> > + }
> > + }
> > + }
> > +
> > + status = clGetProgramInfo(program,
> > + CL_PROGRAM_BINARIES,
> > + sizeof(char *) * num_devices,
> > + binaries,
> > + NULL);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get OpenCL
> program info: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + ret = access_binaries(device_ids, num_devices, cl_file_name, NULL,
> binary_sizes, binaries, 1);
> > +end:
> > + // Release all resouces and memory
> > + if (binaries) {
> > + for (i = 0;i < num_devices;i++ ) {
> > + av_free(binaries[i]);
> > + }
> > + }
> > + av_free(binaries);
> > + av_free(binary_sizes);
> > + av_free(device_ids);
> > + return ret;
> > +}
> > +
>
> > +int av_opencl_create_kernel(const char *kernel_name, AVOpenCLKernelEnv
> *env)
> > +{
> > + int status;
> > + env->kernel = clCreateKernel(gpu_env.program, kernel_name, &status);
>
> Does the user need to check in case the kernel was already set?
>
recently, I do it in user code to ensure the kernel created only one
time.the create kernel operation in the initial stage,
>
> > + env->context = gpu_env.context;
> > + env->command_queue = gpu_env.command_queue;
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not create OpenCL
> kernel: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + return 0;
> > +}
> > +
>
> > +void av_opencl_release_kernel(AVOpenCLKernelEnv *env)
> > +{
> > + int status = clReleaseKernel(env->kernel);
>
> will it crash if env->kernel == NULL?
>
> Should be probably a no-op in that case.
>
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not release kernel:
> %s\n", opencl_errstr(status));
> > + }
> > +}
>
> > +
> > +static int init_opencl_env(GPUEnv *gpu_env, AVOpenCLExternalInfo
> *ext_opencl_info)
> > +{
> > + size_t device_length;
> > + cl_int status;
> > + cl_uint num_platforms, num_devices;
> > + cl_platform_id *platform_ids = NULL;
> > + cl_context_properties cps[3];
> > + char platform_name[100];
> > + int i;
> > + int ret = 0;
> > + cl_device_type device_type[] = {CL_DEVICE_TYPE_GPU,
> CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_DEFAULT};
> > + if (ext_opencl_info) {
> > + if (gpu_env->is_user_created)
> > + return 0;
> > + gpu_env->platform = ext_opencl_info->platform;
> > + gpu_env->is_user_created = 1;
> > + gpu_env->command_queue = ext_opencl_info->command_queue;
> > + gpu_env->context = ext_opencl_info->context;
> > + gpu_env->device_ids = ext_opencl_info->device_ids;
> > + gpu_env->device_id = ext_opencl_info->device_id;
> > + gpu_env->device_type = ext_opencl_info->device_type;
> > + } else {
> > + if (!gpu_env->is_user_created) {
> > + status = clGetPlatformIDs(0, NULL, &num_platforms);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get
> OpenCL platform ids: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + if (num_platforms > 0) {
> > + platform_ids = av_mallocz(num_platforms *
> sizeof(cl_platform_id));
> > + if (!platform_ids) {
> > + ret = AVERROR(ENOMEM);
> > + goto end;
> > + }
> > + status = clGetPlatformIDs(num_platforms, platform_ids,
> NULL);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get
> OpenCL platform ids: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + for (i = 0; i < num_platforms; i++) {
> > + status = clGetPlatformInfo(platform_ids[i],
> CL_PLATFORM_VENDOR,
> > + sizeof(platform_name),
> platform_name,
> > + NULL);
> > +
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not
> get OpenCL platform info: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + gpu_env->platform = platform_ids[i];
> > + status = clGetDeviceIDs(gpu_env->platform /*
> platform */,
> > + CL_DEVICE_TYPE_GPU /*
> device_type */,
> > + 0 /* num_entries */,
> > + NULL /* devices */,
> > + &num_devices);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not
> get OpenCL device number:%s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + if (num_devices == 0) {
> > + //find CPU device
> > + status = clGetDeviceIDs(gpu_env->platform, /*
> platform */
> > + CL_DEVICE_TYPE_CPU, /*
> device_type */
> > + 0, /* num_entries */
> > + NULL, /* devices */
> > + &num_devices);
> > + }
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not
> get OpenCL device ids: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + if(num_devices)
>
> if_(...)
>
> > + break;
> > +
> > + }
> > + }
> > + if (!gpu_env->platform) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get
> OpenCL platforms\n");
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > +
> > + /*
> > + * Use available platform.
> > + */
> > + av_log(&openclutils, AV_LOG_VERBOSE, "Platform Name: %s\n",
> platform_name);
> > + cps[0] = CL_CONTEXT_PLATFORM;
> > + cps[1] = (cl_context_properties)gpu_env->platform;
> > + cps[2] = 0;
> > + /* Check for GPU. */
> > +
> > + for (i = 0; i < sizeof(device_type); i++) {
> > + gpu_env->device_type = device_type[i];
> > + gpu_env->context = clCreateContextFromType(cps,
> > +
> gpu_env->device_type,
> > + NULL,
> > + NULL,
> > + &status);
>
> > + if ((gpu_env->context) && (status == CL_SUCCESS))
>
> you can avoid double parentheses
>
> > + break;
> > + }
>
> > + if ((!gpu_env->context) || (status != CL_SUCCESS)) {
>
> same here
>
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get
> OpenCL context from device type: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
>
> > + /* Detect OpenCL devices. */
> > + /* First, get the size of device list data */
> > + status = clGetContextInfo(gpu_env->context,
> CL_CONTEXT_DEVICES,
> > + 0, NULL, &device_length);
> > + if ((status != CL_SUCCESS) || (device_length == 0)) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get
> OpenCL device length: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + /* Now allocate memory for device list based on the size we
> got earlier */
> > + gpu_env->device_ids = av_mallocz(device_length);
> > + if (!gpu_env->device_ids) {
> > + ret = AVERROR(ENOMEM);
> > + goto end;
> > + }
>
> > + /* Now, get the device list data */
> > + status = clGetContextInfo(gpu_env->context,
> CL_CONTEXT_DEVICES, device_length,
> > + gpu_env->device_ids, NULL);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get
> OpenCL context info: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + /* Create OpenCL command queue. */
> > + gpu_env->command_queue =
> clCreateCommandQueue(gpu_env->context,
> > +
> gpu_env->device_ids[0],
> > + 0, &status);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not create
> OpenCL command queue: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + }
> > + }
> > +end:
> > + av_free(platform_ids);
> > + return ret;
> > +}
> > +
> > +
> > +int av_opencl_register_kernel_function(const char *kernel_name,
> av_opencl_kernel_function function)
> > +{
> > + int i;
>
> > + if (function == NULL) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Input function should not
> be NULL\n");
> > + return AVERROR(EINVAL);
> > + }
>
> or just specify in the docs that it should not be NULL, and blame the
> user in case of crash
>
> > + for (i = 0; i < gpu_env.kernel_count; i++) {
> > + if (av_strcasecmp(kernel_name, gpu_env.kernel_names[i]) == 0) {
> > + gpu_env.kernel_functions[i] = function;
> > + gpu_env.runtime_kernel_count++;
> > + return 0;
> > + }
> > + }
>
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not find the kernle:
> %s\n", kernel_name);
>
> typo, also the feedback could be improved, e.g. with:
> "Could not find a kernel with name '%s', cannot register function\n"
>
> > + return AVERROR(EINVAL);
> > +}
> > +
>
> > +static int compile_kernel_file(const char *filename, GPUEnv *gpu_env,
> const char *build_option)
> > +{
> > + cl_int status;
> > + size_t file_size = 0;
> > + char *source_str = NULL;
> > + char *buildlog = NULL;
> > + int b_error = 0, binary_status, binary_existed;
> > + char *binary = NULL;
> > + char *temp;
> > + size_t num_devices;
> > + cl_device_id *device_ids = NULL;
> > + FILE *fd = NULL;
> > + int ret = 0;
> > + int i;
> > +
> > + if (av_strcasecmp(gpu_env->kernel_srcfile_name, filename) == 0)
> > + if (gpu_env->program)
> > + return ret;
> > + ret = check_generated_binary(gpu_env->context, filename, &fd);
> > + if (ret)
> > + return ret;
> > + binary_existed = !!fd;
> > + if (binary_existed) {
> > + status = clGetContextInfo(gpu_env->context,
> > + CL_CONTEXT_NUM_DEVICES,
> > + sizeof(num_devices),
> > + &num_devices,
> > + NULL);
> > + if(status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not release
> OpenCL device number: %s\n",
> > + opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + device_ids = av_mallocz(sizeof(cl_device_id) * num_devices);
> > + if (!device_ids) {
> > + return AVERROR(ENOMEM);
> > + }
>
> > + b_error |= fseek(fd, 0, SEEK_END) < 0;
> > + b_error |= (file_size = ftell(fd)) <= 0;
> > + b_error |= fseek(fd, 0, SEEK_SET) < 0;
> > + if (b_error) {
> > + ret = AVERROR(errno);
> > + goto end;
> > + }
>
> > + binary = av_mallocz(file_size);
> > + if (!binary) {
> > + ret = AVERROR(ENOMEM);
> > + goto end;
> > + }
> > +
> > + b_error |= fread(binary, 1, file_size, fd) != file_size;
> > + fclose(fd);
> > + fd = NULL;
> > + /* grab the handles to all of the devices in the context. */
> > + status = clGetContextInfo(gpu_env->context,
> > + CL_CONTEXT_DEVICES,
> > + sizeof(cl_device_id) * num_devices,
> > + device_ids,
> > + NULL);
> > + if(status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get OpenCL
> device ids: %s\n",
> > + opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + gpu_env->program = clCreateProgramWithBinary(gpu_env->context,
> > + num_devices,
> > + device_ids,
> > + &file_size,
> > + (const
> uint8_t**)&binary,
> > + &binary_status,
> > + &status);
> > + if(status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not create program
> with binary: %s\n",
> > + opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > +
> > + } else {
> > + for (i = 0; i < gpu_env->kernel_count; i++) {
> > + file_size += strlen(gpu_env->kernel_code[i]);
> > + }
> > + source_str = av_mallocz(file_size + 1);
> > + if (!source_str) {
> > + return AVERROR(ENOMEM);
> > + }
> > + temp = source_str;
> > + for (i = 0; i < gpu_env->kernel_count; i++) {
> > + memcpy(temp, gpu_env->kernel_code[i],
> strlen(gpu_env->kernel_code[i]));
> > + temp += strlen(gpu_env->kernel_code[i]);
> > + }
> > + /* create a CL program using the kernel source */
> > + gpu_env->program = clCreateProgramWithSource(gpu_env->context,
> > + 1,
> > + (const char
> **)(&source_str),
> > + &file_size,
> > + &status);
> > + if(status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not create OpenCL
> program with source code: %s\n",
> > + opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + }
> > +
> > + if ((!gpu_env->program) || (status != CL_SUCCESS)) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not create OpenCL
> program: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > +
> > +
> > + /* create a cl program executable for all the devices specified */
> > + if (!gpu_env->is_user_created)
> > + status = clBuildProgram(gpu_env->program, 1,
> gpu_env->device_ids,
> > + build_option, NULL, NULL);
> > + else
> > + status = clBuildProgram(gpu_env->program, 1,
> &(gpu_env->device_id),
> > + build_option, NULL, NULL);
> > +
> > + if (status != CL_SUCCESS) {
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not compile
> OpenCL kernel: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + }
> > + if (!gpu_env->is_user_created)
> > + status = clGetProgramBuildInfo(gpu_env->program,
> > + gpu_env->device_ids[0],
> > + CL_PROGRAM_BUILD_LOG, 0,
> NULL, &file_size);
> > + else
> > + status = clGetProgramBuildInfo(gpu_env->program,
> > + gpu_env->device_id,
> > + CL_PROGRAM_BUILD_LOG, 0,
> NULL, &file_size);
> > +
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get OpenCL
> build info length: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + buildlog = av_mallocz(file_size);
> > + if (!buildlog) {
> > + ret = AVERROR(ENOMEM);
> > + goto end;
> > + }
> > + if (!gpu_env->is_user_created)
> > + status = clGetProgramBuildInfo(gpu_env->program,
> gpu_env->device_ids[0],
> > + CL_PROGRAM_BUILD_LOG,
> file_size, buildlog, &file_size);
> > + else
> > + status = clGetProgramBuildInfo(gpu_env->program,
> gpu_env->device_id,
> > + CL_PROGRAM_BUILD_LOG,
> file_size, buildlog, &file_size);
> > +
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not get OpenCL
> build info: %s\n", opencl_errstr(status));
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + fd = fopen("kernel-build.log", "w+");
> > + if (fd) {
> > + fwrite(buildlog, sizeof(char), file_size, fd);
> > + fclose(fd);
> > + }
> > + ret = AVERROR_EXTERNAL;
> > + goto end;
> > + }
> > + av_strlcpy(gpu_env->kernel_srcfile_name, filename,
> MAX_KERNEL_SRCFILE_LEN);
> > + if (binary_existed == 0) {
> > + ret = generate_bin_from_kernel_source(gpu_env->program,
> filename);
> > + if (ret) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not generate
> binary file from kernel source code\n");
> > + goto end;
> > + }
> > + }
> > +end:
> > + av_free(source_str);
> > + av_free(device_ids);
> > + av_free(binary);
> > + av_free(buildlog);
> > + return ret;
> > +}
>
> we could probably delay this stuff to when we know how to deal with it
> security-wise
>
yes, agree with you.
>
> > +
> > +static int get_kernel_env_and_func(const char *kernel_name,
> > + AVOpenCLKernelEnv *env,
> > + av_opencl_kernel_function
> *function)
> > +{
> > + for (int i = 0; i < gpu_env.kernel_count; i++) {
> > + if (av_strcasecmp(kernel_name, gpu_env.kernel_names[i]) == 0) {
> > + env->context = gpu_env.context;
> > + env->command_queue = gpu_env.command_queue;
> > + env->program = gpu_env.program;
> > + env->kernel = gpu_env.kernels[i];
> > + *function = gpu_env.kernel_functions[i];
> > + return 0;
> > + }
> > + }
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not find kernel: %s\n",
> kernel_name);
> > + return AVERROR(EINVAL);
> > +}
> > +
> > +int av_opencl_run_kernel(const char *kernel_name, void **userdata)
> > +{
> > + AVOpenCLKernelEnv env = {0};
> > + av_opencl_kernel_function function;
> > + int status;
> > + av_strlcpy(env.kernel_name, kernel_name,
> AV_OPENCL_MAX_KERNEL_NAME_SIZE);
>
> > + status = get_kernel_env_and_func(kernel_name, &env, &function);
>
> you could avoid a level of indirection, and just include the code here
>
>
> > + if (!status) {
> > + return(function(userdata, &env));
> > + }
> > + return status;
> > +}
> > +
>
> > +int av_opencl_init(const char *build_option, AVOpenCLExternalInfo
> *ext_opencl_info)
> > +{
> > + int ret;
> > + if (!isinited) {
>
> > + /*initialize devices, context, comand_queue*/
>
> command
>
> > + ret = init_opencl_env(&gpu_env, ext_opencl_info);
> > + if (ret) {
> > + return ret;
> > + }
> > + /*initialize program, kernel_name, kernel_count*/
> > + ret = compile_kernel_file("ffmpeg-kernels", &gpu_env,
> build_option);
> > + if (ret) {
> > + return ret;
> > + }
> > + av_assert1(gpu_env.kernel_count > 0);
> > + isinited = 1;
> > + }
> > + return 0;
> > +}
>
> I suggest to merge this with init_opencl_env(), in order to avoid a
> level of indirection and save some lines.
>
The initial operation include two steps, first init opencl env, second
compile kernel. I think it merge the two functions in this function, the
function will be too large.
>
> > +
> > +void av_opencl_uninit(void)
> > +{
> > + int status;
> > + if (!isinited)
> > + return;
> > + av_freep(&(gpu_env.temp_buffer));
> > + if (gpu_env.is_user_created)
> > + return;
>
> > + gpu_env.runtime_kernel_count--;
> > + if (!gpu_env.runtime_kernel_count) {
>
> this is a bit weird, because the code is executed also in case
> runtime_kernel_count is 0.
>
it is because each filter sill call this, but the opencl env should not be
released before all the kernels are release. So I set a counter, when all
kernels are released, then release the opencl env.
>
> > + if (gpu_env.program) {
> > + status = clReleaseProgram(gpu_env.program);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not release
> OpenCL program: %s\n", opencl_errstr(status));
> > + }
> > + gpu_env.program = NULL;
> > + }
> > + if (gpu_env.command_queue) {
> > + status = clReleaseCommandQueue(gpu_env.command_queue);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not release
> OpenCL command queue: %s\n", opencl_errstr(status));
> > + }
> > + gpu_env.command_queue = NULL;
> > + }
> > + if (gpu_env.context) {
> > + status = clReleaseContext(gpu_env.context);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not release
> OpenCL context: %s\n", opencl_errstr(status));
> > + }
> > + gpu_env.context = NULL;
> > + }
> > + av_freep(&(gpu_env.device_ids));
> > + isinited = 0;
> > + }
> > +}
>
> Note: this and init are not thread safe, so we should specify that
> they should be called by a single thread in the docs.
>
> > +
> > +int av_opencl_is_inited(void)
> > +{
> > + return isinited;
> > +}
> > +
> > +void av_opencl_get_kernel_env(AVOpenCLKernelEnv *env)
> > +{
> > + env->context = gpu_env.context;
> > + env->command_queue = gpu_env.command_queue;
> > + env->program = gpu_env.program;
>
> nit: vertical align:
> env->context = gpu_env.context;
> env->command_queue = gpu_env.command_queue;
> env->program = gpu_env.program;
>
> > +}
> > +
> > +int av_opencl_buffer_create(void **cl_buf, size_t cl_buf_size, int
> flags, void *host_ptr)
> > +{
> > + int status;
> > + *cl_buf = clCreateBuffer(gpu_env.context, flags, cl_buf_size,
> host_ptr, &status);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not create OpenCL
> buffer: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + return 0;
> > +}
> > +
> > +void av_opencl_buffer_release(void *cl_buf)
> > +{
> > + int status = 0;
> > + if (!cl_buf)
> > + return;
> > + status = clReleaseMemObject(cl_buf);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not release OpenCL
> buffer: %s\n", opencl_errstr(status));
> > + }
> > +}
> > +
>
> > +int av_opencl_buffer_write(uint8_t *src_buf, void *dst_cl_buf, size_t
> buf_size)
>
> int av_opencl_buffer_write(void *dst_cl_buf, const uint8_t *src_buf,
> size_t buf_size)
>
> the usual C convention is:
> DST, SRC
>
> like in memcpy
> (also note the const)
>
> > +{
> > + int status;
> > + void *mapped = clEnqueueMapBuffer(gpu_env.command_queue, dst_cl_buf,
> > + CL_TRUE,CL_MAP_WRITE, 0,
> sizeof(uint8_t) * buf_size,
>
> you could probably omit sizeof(uint8_t) here and below since it is
> always trivially 1
>
> > + 0, NULL, NULL, &status);
> > +
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not map OpenCL
> buffer: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + memcpy(mapped, src_buf, sizeof(uint8_t) * buf_size);
> > +
> > + status = clEnqueueUnmapMemObject(gpu_env.command_queue, dst_cl_buf,
> mapped, 0, NULL, NULL);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not unmap OpenCL
> buffer: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + return 0;
> > +}
> > +
> > +int av_opencl_buffer_read(uint8_t *dst_buf, void *src_cl_buf, size_t
> buf_size)
> > +{
> > + int status;
> > + void *mapped = clEnqueueMapBuffer(gpu_env.command_queue, src_cl_buf,
> > + CL_TRUE,CL_MAP_READ, 0,
> sizeof(uint8_t) * buf_size,
> > + 0, NULL, NULL, &status);
> > +
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not map OpenCL
> buffer: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + memcpy(dst_buf, mapped, sizeof(uint8_t) * buf_size);
> > +
> > + status = clEnqueueUnmapMemObject(gpu_env.command_queue, src_cl_buf,
> mapped, 0, NULL, NULL);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not unmap OpenCL
> buffer: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + return 0;
> > +}
> > +
> > +int av_opencl_buffer_write_image(void *dst_cl_buf, size_t
> cl_buffer_size,
> > + uint8_t **data, int *plane_size,
> int plane_num, int offset)
> > +{
> > + int buffersize = 0;
> > + uint8_t *temp;
> > + int status;
> > + void *mapped;
> > + int i;
>
> > + if ((plane_num > 8) || (plane_num <= 0)) {
>
> nit: () can be avoided, also you could write
>
> if ((unsigned int)plane_num > 8) {
> ...
>
> > + return AVERROR(EIO);
>
> return AVERROR(EINVAL)
>
> Also how can you have more than 4 planes?
>
this is because the AVFrame has data[0]-data[7], and I thinks we should
follow it.
>
> > + }
>
> > + for (i = 0;i < plane_num;i++) {
> > + buffersize += plane_size[i];
> > + }
>
> > + if (buffersize > cl_buffer_size) {
> > + av_log(&openclutils, AV_LOG_ERROR, "CPU memory buffer size is
> too large\n");
>
> av_log(..., "Cannot write image to OpenCL buffer: buffer too small\n");
>
> seems more direct to me
>
> > + return AVERROR_EXTERNAL;
>
> return AVERROR(EINVAL)?
>
> > + }
> > + mapped = clEnqueueMapBuffer(gpu_env.command_queue, dst_cl_buf,
> > + CL_TRUE,CL_MAP_WRITE, 0,
> buffersize + offset,
> > + 0, NULL, NULL, &status);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not map OpenCL
> buffer: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + temp = mapped;
> > + temp += offset;
> > + for (i = 0; i < plane_num; i++) {
> > + memcpy(temp, data[i], plane_size[i]);
> > + temp += plane_size[i];
> > + }
> > + status = clEnqueueUnmapMemObject(gpu_env.command_queue, dst_cl_buf,
> mapped, 0, NULL, NULL);
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not unmap OpenCL
> buffer: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + return 0;
> > +}
> > +
> > +int av_opencl_buffer_read_image(void *src_cl_inbuf, size_t
> cl_buffer_size, uint8_t **data,
> > + int *plane_size, int
> plane_num)
> > +{
> > + int size = 0;
> > + int ret = 0;
> > + uint8_t *temp;
> > + int i;
> > + if ((plane_num > 8) || (plane_num <= 0)) {
> > + return AVERROR(EIO);
> > + }
> > + for (i = 0;i < plane_num;i++) {
> > + size += plane_size[i];
> > + }
> > + if (size > cl_buffer_size) {
> > + av_log(&openclutils, AV_LOG_ERROR, "CPU memory buffer size is
> too large\n");
> > + return AVERROR_EXTERNAL;
> > + }
> > + if (!(gpu_env.temp_buffer)) {
> > + gpu_env.temp_buffer = av_malloc(size);
> > + if (!gpu_env.temp_buffer)
> > + return AVERROR(ENOMEM);
> > + gpu_env.temp_buffer_size = size;
> > + }
> > +
> > + if(size > gpu_env.temp_buffer_size) {
> > + av_free(gpu_env.temp_buffer);
> > + gpu_env.temp_buffer = av_malloc(size);
> > + if (!gpu_env.temp_buffer)
> > + return AVERROR(ENOMEM);
> > + gpu_env.temp_buffer_size = size;
> > + }
> > + temp = gpu_env.temp_buffer;
> > + ret = av_opencl_buffer_read(gpu_env.temp_buffer, src_cl_inbuf,
> size);
> > + if (!ret) {
> > + for (i = 0;i < plane_num;i++) {
> > + memcpy(data[i], temp, plane_size[i]);
> > + temp += plane_size[i];
> > + }
> > + }
> > + return ret;
> > +}
> > +
> > +cl_device_id av_opencl_get_device_id(void)
> > +{
> > + if (!gpu_env.is_user_created) {
> > + return *(gpu_env.device_ids);
> > + } else
> > + return gpu_env.device_id;
> > +}
> > +
> > +cl_context av_opencl_get_context(void)
> > +{
> > + return gpu_env.context;
> > +}
> > +
> > +cl_command_queue av_opencl_get_command_queue(void)
> > +{
> > + return gpu_env.command_queue;
> > +}
> > +
> > diff --git a/libavutil/opencl.h b/libavutil/opencl.h
> > new file mode 100644
> > index 0000000..ee2f3a4
> > --- /dev/null
> > +++ b/libavutil/opencl.h
> > @@ -0,0 +1,219 @@
> > +/*
> > + * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
> > + * Copyright (C) 2012 Li Cao <li at multicorewareinc.com>
> > + * Copyright (C) 2012 Wei Gao <weigao at multicorewareinc.com>
> > + *
> > + * This file is part of FFmpeg.
> > + *
> > + * FFmpeg is free software; you can redistribute it and/or
> > + * modify it under the terms of the GNU Lesser General Public
> > + * License as published by the Free Software Foundation; either
> > + * version 2.1 of the License, or (at your option) any later version.
> > + *
> > + * FFmpeg is distributed in the hope that it will be useful,
> > + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
> > + * Lesser General Public License for more details.
> > + *
> > + * You should have received a copy of the GNU Lesser General Public
> > + * License along with FFmpeg; if not, write to the Free Software
> > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
> 02110-1301 USA
> > + */
> > +
> > +#include "config.h"
> > +
> > +#ifndef LIBAVUTIL_OPENCLWRAPPER_H
> > +#define LIBAVUTIL_OPENCLWRAPPER_H
> > +
> > +#include <CL/cl.h>
> > +
> > +#define AV_OPENCL_KERNEL( ... )# __VA_ARGS__
> > +
> > +#define AV_OPENCL_MAX_KERNEL_NAME_SIZE 150
> > +
> > +typedef struct AVOpenCLKernelEnv {
> > + cl_context context;
> > + cl_command_queue command_queue;
> > + cl_program program;
> > + cl_kernel kernel;
> > + char kernel_name[AV_OPENCL_MAX_KERNEL_NAME_SIZE];
> > +} AVOpenCLKernelEnv;
> > +
> > +typedef struct AVOpenCLExternalInfo {
> > + cl_platform_id platform;
> > + cl_device_type device_type;
> > + cl_context context;
> > + cl_device_id *device_ids;
> > + cl_device_id device_id;
> > + cl_command_queue command_queue;
> > + char *platform_name;
> > +} AVOpenCLExternalInfo;
> > +
> > +/**
> > + * User defined function, used to set the input parameter in the kernel
> > + *environment. This function launches kernel and copies data from GPU to
> > + *CPU, or from CPU to GPU.
> > + */
> > +typedef int (* av_opencl_kernel_function) (void **userdata,
> AVOpenCLKernelEnv *kenv);
> > +
> > +/**
> > + * Register a function for running the kernel specified by the kernel
> name.
> > + *@param kernel_name this kernel name is used to find the kernel in
> OpenCL runtime environment.
> > + *@param function user defined function, used to set the input
> parameter in the kernel environment
> > + *@return >=0 on success, a negative error value on failure
> > + */
> > +int av_opencl_register_kernel_function(const char *kernel_name,
> av_opencl_kernel_function function);
> > +
> > +/**
> > + *Load OpenCL kernel.
> > + *
> > + *@param kernel_name this kernel name is used to find the kernel in
> OpenCL runtime environment.
> > + *@param userdata this userdata is the all parameters for
> running the kernel specified by kernel name
> > + *@return >=0 on success, a negative error value on failure
> > + */
> > +int av_opencl_run_kernel(const char *kernel_name, void **userdata);
> > +
> > +/**
> > + * Init the run time OpenCL environment.
> > + *
> > + *This function must be called befor calling any function related to
> OpenCL.
> > + *
> > + *
> > + *@param build_option option of compile the kernel in OpenCL
> runtime environment,reference "OpenCL Specification Version: 1.2 chapter
> 5.6.4"
> > + *@param ext_opencl_info this is the extern OpenCL environment which
> the application program has created
> > + *@return >=0 on success, a negative error value on failure
> > + */
> > +int av_opencl_init(const char *build_option, AVOpenCLExternalInfo
> *ext_opencl_info);
> > +
> > +/**
> > + * Release OpenCL resources , this function must be called after
> calling any functions related to OpenCL.
> > + */
> > +void av_opencl_uninit(void);
> > +
> > +/**
> > + * Get the OpenCL status, this function is used the check whether or
> not the OpenCL environment has been created.
> > + *
> > + *@return 0 not inited, 1, inited;
> > + */
> > +int av_opencl_is_inited(void);
> > +
>
> > +/**
> > + * Create kernel object by a kernel name on the specified OpenCL run
> time
> > + * indicated by the env parameter.
>
> Create kernel object on the specified OpenCL run in env.
>
> > + *
> > + * @param kernel_name kernel name
> > + * @param env the kernel environment which has been
> created by av_opencl_init_run_env
>
> av_opencl_init()?
>
> Also, what if the kernel is already created (leak in this case)?
>
> > + * @return >=0 on success, a negative error code on failure
> > + */
> > +int av_opencl_create_kernel(const char *kernel_name, AVOpenCLKernelEnv
> *env);
> > +
> > +/**
> > + * Release kernel object.
> > + *
>
> > + *@param env the kernel environment which has been created by
> av_opencl_init_run_env.
>
> init()
>
> [...]
> --
> FFmpeg = Fabulous and Foolish Mean Philosofic Epic Geisha
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel at ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
More information about the ffmpeg-devel
mailing list