[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