[FFmpeg-devel] [PATCH]opencl: automatically select the fastest opencl device

Lenny Wang lenny at multicorewareinc.com
Mon Dec 2 08:05:57 CET 2013


On Mon, Dec 2, 2013 at 12:19 AM, Wei Gao <highgod0401 at gmail.com> wrote:
> 2013/12/2 Lenny Wang <lenny at multicorewareinc.com>
>
>>
>>
>> Fixed all the nits, new patch attached, thanks.
>>
>>> From: Lenny Wang <lwanghpc at gmail.com>
>>
>> Date: Sun, 1 Dec 2013 18:15:00 -0600
>>
>> Subject: [PATCH] cmdutils & opencl: add -opencl_bench option to test and
>>> show available OpenCL devices
>>
>>
>>> ---
>>
>>  cmdutils.c                     | 164
>>> +++++++++++++++++++++++++++++++++++++++++
>>
>>  cmdutils.h                     |  15 ++++
>>
>>  cmdutils_common_opts.h         |   1 +
>>
>>  cmdutils_opencl_bench_kernel.h |  86 +++++++++++++++++++++
>>
>>  doc/fftools-common-opts.texi   |   4 +
>>
>>  doc/utils.texi                 |   4 +-
>>
>>  libavutil/opencl.c             |  42 +++++++++++
>>
>>  libavutil/opencl.h             |  15 ++++
>>
>>  8 files changed, 329 insertions(+), 2 deletions(-)
>>
>>
>>> diff --git a/cmdutils.c b/cmdutils.c
>>
>> index 46ade3f..872b099 100644
>>
>> --- a/cmdutils.c
>>
>> +++ b/cmdutils.c
>>
>> @@ -60,6 +60,8 @@
>>
>>  #endif
>>
>>  #if CONFIG_OPENCL
>>
>>  #include "libavutil/opencl.h"
>>
>> +#include "libavutil/time.h"
>>
>> +#include "cmdutils_opencl_bench_kernel.h"
>>
>>  #endif
>>
>>
>>
>>
>>
>> @@ -986,6 +988,168 @@ int opt_timelimit(void *optctx, const char *opt,
>>> const char *arg)
>>
>>  }
>>
>>
>>
>>  #if CONFIG_OPENCL
>>
>> +#define OCLCHECK(method, ... )
>>>       \
>>
>> +do {
>>>       \
>>
>> +    status = method(__VA_ARGS__);
>>>      \
>>
>> +    if (status != CL_SUCCESS) {
>>>      \
>>
>> +        av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n",
>>>       \
>>
>> +               av_opencl_errstr(status));
>>>      \
>>
>> +        ret = AVERROR_EXTERNAL;
>>>      \
>>
>> +        goto end;
>>>      \
>>
>> +    }
>>>      \
>>
>> +} while (0)
>>
>> +
>>
>> +#define CREATEBUF(out, flags, size)
>>>      \
>>
>> +do {
>>>       \
>>
>> +    out = clCreateBuffer(ext_opencl_env->context, flags, size, NULL,
>>> &status); \
>>
>> +    if (status != CL_SUCCESS) {
>>>      \
>>
>> +        av_log(NULL, AV_LOG_ERROR, "Could not create OpenCL buffer\n");
>>>      \
>>
>> +        ret = AVERROR_EXTERNAL;
>>>      \
>>
>> +        goto end;
>>>      \
>>
>> +    }
>>>      \
>>
>> +} while (0)
>>
>> +
>>
>> +static void fill_rand_int(int *data, int n)
>>
>> +{
>>
>> +    int i;
>>
>> +    srand(av_gettime());
>>
>> +    for (i = 0; i < n; i++)
>>
>> +        data[i] = rand();
>>
>> +}
>>
>> +
>>
>> +#define OPENCL_NB_ITER 5
>>
>> +static int64_t run_opencl_bench(AVOpenCLExternalEnv *ext_opencl_env)
>>
>> +{
>>
>> +    int i, arg = 0, width = 1920, height = 1088;
>>
>> +    int64_t start, ret = 0;
>>
>> +    cl_int status;
>>
>> +    size_t kernel_len;
>>
>> +    char *inbuf;
>>
>> +    int *mask;
>>
>> +    int buf_size = width * height * sizeof(char);
>>
>> +    int mask_size = sizeof(uint32_t) * 128;
>>
>> +
>>
>> +    cl_mem cl_mask, cl_inbuf, cl_outbuf;
>>
>> +    cl_kernel kernel = NULL;
>>
>> +    cl_program program = NULL;
>>
>> +    size_t local_work_size_2d[2] = {16, 16};
>>
>> +    size_t global_work_size_2d[2] = {(size_t)width, (size_t)height};
>>
>> +
>>
>> +    if (!(inbuf = av_malloc(buf_size)) || !(mask = av_malloc(mask_size)))
>>> {
>>
>> +        av_log(NULL, AV_LOG_ERROR, "Out of memory\n");
>>
>> +        return AVERROR(ENOMEM);
>>
>> +    }
>>
>> +    fill_rand_int((int*)inbuf, buf_size/4);
>>
>> +    fill_rand_int(mask, mask_size/4);
>>
>> +
>>
>> +    CREATEBUF(cl_mask, CL_MEM_READ_ONLY, mask_size);
>>
>> +    CREATEBUF(cl_inbuf, CL_MEM_READ_ONLY, buf_size);
>>
>> +    CREATEBUF(cl_outbuf, CL_MEM_READ_WRITE, buf_size);
>>
>> +
>>
>> +    kernel_len = strlen(ocl_bench_source);
>>
>> +    program = clCreateProgramWithSource(ext_opencl_env->context, 1,
>>> &ocl_bench_source,
>>
>> +                                        &kernel_len, &status);
>>
>> +    if (status != CL_SUCCESS || !program) {
>>
>> +        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark
>>> program\n");
>>
>> +        ret = AVERROR_EXTERNAL;
>>
>> +        goto end;
>>
>> +    }
>>
>> +    status = clBuildProgram(program, 1, &(ext_opencl_env->device_id),
>>> NULL, NULL, NULL);
>>
>> +    if (status != CL_SUCCESS) {
>>
>> +        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to build benchmark
>>> program\n");
>>
>> +        ret = AVERROR_EXTERNAL;
>>
>> +        goto end;
>>
>> +    }
>>
>> +    kernel = clCreateKernel(program, "unsharp_bench", &status);
>>
>> +    if (status != CL_SUCCESS) {
>>
>> +        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark
>>> kernel\n");
>>
>> +        ret = AVERROR_EXTERNAL;
>>
>> +        goto end;
>>
>> +    }
>>
>> +
>>
>> +    OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue,
>>> cl_inbuf, CL_TRUE, 0,
>>
>> +             buf_size, inbuf, 0, NULL, NULL);
>>
>> +    OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue,
>>> cl_mask, CL_TRUE, 0,
>>
>> +             mask_size, mask, 0, NULL, NULL);
>>
>> +    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_inbuf);
>>
>> +    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_outbuf);
>>
>> +    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_mask);
>>
>> +    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &width);
>>
>> +    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &height);
>>
>> +
>>
>> +    start = av_gettime();
>>
>> +    for (i = 0; i < OPENCL_NB_ITER; i++)
>>
>> +    OCLCHECK(clEnqueueNDRangeKernel, ext_opencl_env->command_queue,
>>> kernel, 2, NULL,
>>
>> +                 global_work_size_2d, local_work_size_2d, 0, NULL, NULL);
>>
>> + clFinish(ext_opencl_env->command_queue);
>>
>> clFinish indent
>
> +    ret = (av_gettime() - start)/OPENCL_NB_ITER;
>>
>> +end:
>>
>> +    if (kernel)
>>
>> +        clReleaseKernel(kernel);
>>
>> +    if (program)
>>
>> +        clReleaseProgram(program);
>>
>> +    if (cl_inbuf)
>>
>> +        clReleaseMemObject(cl_inbuf);
>>
>> +    if (cl_outbuf)
>>
>> +        clReleaseMemObject(cl_outbuf);
>>
>> +    if (cl_mask)
>>
>> +        clReleaseMemObject(cl_mask);
>>
>> +    av_free(inbuf);
>>
>> +    av_free(mask);
>>
>> +    return ret;
>>
>> +}
>>
>> +
>>
>> +static int compare_ocl_device_desc(const void *a, const void *b)
>>
>> +{
>>
>> +    return ((oclDeviceBenchmark*)a)->runtime -
>>> ((oclDeviceBenchmark*)b)->runtime;
>>
>> +}
>>
>> +
>>
>> +int opt_test_opencl(void *optctx, const char *opt, const char *arg)
>>
>> +{
>>
>> +    int i, j, nb_devices = 0, count = 0;
>>
>> +    int64_t score = 0;
>>
>> +    AVOpenCLDeviceList *device_list;
>>
>> +    AVOpenCLDeviceNode *device_node = NULL;
>>
>> +    oclDeviceBenchmark *devices = NULL;
>>
>> +    cl_platform_id platform;
>>
>> +
>>
>> +    av_opencl_get_device_list(&device_list);
>>
>> +    for (i = 0; i < device_list->platform_num; i++)
>>
>> +        nb_devices += device_list->platform_node[i]->device_num;
>>
>> +    if (!nb_devices) {
>>
>> +        av_log(NULL, AV_LOG_ERROR, "No OpenCL device detected!\n");
>>
>> +        return AVERROR(EINVAL);
>>
>> +    }
>>
>> +    if (!(devices = av_malloc(sizeof(oclDeviceBenchmark) * nb_devices))) {
>>
>> +        av_log(NULL, AV_LOG_ERROR, "Could not allocate buffer\n");
>>
>> +        return AVERROR(ENOMEM);
>>
>> +    }
>>
>> +
>>
>> +    for (i = 0; i < device_list->platform_num; i++) {
>>
>> +        for (j = 0; j < device_list->platform_node[i]->device_num; j++) {
>>
>> +            device_node = device_list->platform_node[i]->device_node[j];
>>
>> +            platform = device_list->platform_node[i]->platform_id;
>>
>> +            score = av_opencl_benchmark(device_node, platform,
>>> run_opencl_bench);
>>
>> +            if (score > 0) {
>>
>> +                devices[count].platform_idx = i;
>>
>> +                devices[count].device_idx = j;
>>
>> +                devices[count].runtime = score;
>>
>> +                strcpy(devices[count].device_name,
>>> device_node->device_name);
>>
>> +                count++;
>>
>> +            }
>>
>> +        }
>>
>> +    }
>>
>> +    qsort(devices, count, sizeof(oclDeviceBenchmark),
>>> compare_ocl_device_desc);
>>
>> +    for (i = 0; i < count; i++)
>>
>> +
>>>  printf("platform_idx=%d\tdevice_idx=%d\tdevice_name=%s\truntime=%"PRId64"\n",
>>
>> +               devices[i].platform_idx, devices[i].device_idx,
>>
>> +               devices[i].device_name, devices[i].runtime);
>>
>> +
>>
>> +    av_opencl_free_device_list(&device_list);
>>
>> +    av_free(devices);
>>
>> +    return 0;
>>
>> +}
>>
>> +
>>
>>  int opt_opencl(void *optctx, const char *opt, const char *arg)
>>
>>  {
>>
>>      char *key, *value;
>>
>> diff --git a/cmdutils.h b/cmdutils.h
>>
>> index b814961..082d19c 100644
>>
>> --- a/cmdutils.h
>>
>> +++ b/cmdutils.h
>>
>> @@ -98,8 +98,23 @@ int opt_max_alloc(void *optctx, const char *opt, const
>>> char *arg);
>>
>>
>>
>>  int opt_codec_debug(void *optctx, const char *opt, const char *arg);
>>
>>
>>
>> +#if CONFIG_OPENCL
>>
>>  int opt_opencl(void *optctx, const char *opt, const char *arg);
>>
>>
>>
>> +typedef struct {
>>
>> +    int platform_idx;
>>
>> +    int device_idx;
>>
>> +    char device_name[64];
>>
>> +    int64_t runtime;
>>
>> +} oclDeviceBenchmark;
>>
>> +
>>
>> +/**
>>
>> + * Perform benchmark on all OpenCL devices and print results.
>>
>> + * This option processing function does not utilize the arguments.
>>
>> + */
>>
>> +int opt_test_opencl(void *optctx, const char *opt, const char *arg);
>>
>> +#endif
>>
>> +
>>
>>  /**
>>
>>   * Limit the execution time.
>>
>>   */
>>
>> diff --git a/cmdutils_common_opts.h b/cmdutils_common_opts.h
>>
>> index 3e3f0ac..8de75b9 100644
>>
>> --- a/cmdutils_common_opts.h
>>
>> +++ b/cmdutils_common_opts.h
>>
>> @@ -22,5 +22,6 @@
>>
>>      { "max_alloc"  , HAS_ARG,  {.func_arg = opt_max_alloc},     "set
>>> maximum size of a single allocated block", "bytes" },
>>
>>      { "cpuflags"   , HAS_ARG | OPT_EXPERT, { .func_arg = opt_cpuflags },
>>> "force specific cpu flags", "flags" },
>>
>>  #if CONFIG_OPENCL
>>
>> +    { "opencl_bench", OPT_EXIT, {.func_arg = opt_test_opencl},  "run
>>> benchmark across available OpenCL devices" },
>>
>>      { "opencl_options", HAS_ARG, {.func_arg = opt_opencl},      "set
>>> OpenCL environment options" },
>>
>>  #endif
>>
>> diff --git a/cmdutils_opencl_bench_kernel.h
>>> b/cmdutils_opencl_bench_kernel.h
>>
>> new file mode 100644
>>
>> index 0000000..ce6bad5
>>
>> --- /dev/null
>>
>> +++ b/cmdutils_opencl_bench_kernel.h
>>
>> @@ -0,0 +1,86 @@
>>
>> +/*
>>
>> + * Copyright (C) 2013 Lenny Wang
>>
>> + *
>>
>> + * 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
>>
>> + */
>>
>> +
>>
>> +#ifndef CMDUTILS_OPENCL_BENCH_KERNEL_H
>>
>> +#define CMDUTILS_OPENCL_BENCH_KERNEL_H
>>
>> +
>>
>> +#include "libavutil/opencl.h"
>>
>> +
>>
>> +const char *ocl_bench_source = AV_OPENCL_KERNEL(
>>
>> +inline unsigned char clip_uint8(int a)
>>
>> +{
>>
>> +    if (a & (~0xFF))
>>
>> +        return (-a)>>31;
>>
>> +    else
>>
>> +        return a;
>>
>> +}
>>
>> +
>>
>> +kernel void unsharp_bench(
>>
>> +                    global unsigned char *src,
>>
>> +                    global unsigned char *dst,
>>
>> +                    global int *mask,
>>
>> +                    int width,
>>
>> +                    int height)
>>
>> +{
>>
>> +    int i, j, local_idx, lc_idx, sum = 0;
>>
>> +    int2 thread_idx, block_idx, global_idx, lm_idx;
>>
>> +    thread_idx.x = get_local_id(0);
>>
>> +    thread_idx.y = get_local_id(1);
>>
>> +    block_idx.x = get_group_id(0);
>>
>> +    block_idx.y = get_group_id(1);
>>
>> +    global_idx.x = get_global_id(0);
>>
>> +    global_idx.y = get_global_id(1);
>>
>> +    local uchar data[32][32];
>>
>> +    local int lc[128];
>>
>> +
>>
>> +    for (i = 0; i <= 1; i++) {
>>
>> +        lm_idx.y = -8 + (block_idx.y + i) * 16 + thread_idx.y;
>>
>> +        lm_idx.y = lm_idx.y < 0 ? 0 : lm_idx.y;
>>
>> +        lm_idx.y = lm_idx.y >= height ? height - 1: lm_idx.y;
>>
>> +        for (j = 0; j <= 1; j++) {
>>
>> +            lm_idx.x = -8 + (block_idx.x + j) * 16 + thread_idx.x;
>>
>> +            lm_idx.x = lm_idx.x < 0 ? 0 : lm_idx.x;
>>
>> +            lm_idx.x = lm_idx.x >= width ? width - 1: lm_idx.x;
>>
>> +            data[i*16 + thread_idx.y][j*16 + thread_idx.x] =
>>> src[lm_idx.y*width + lm_idx.x];
>>
>> +        }
>>
>> +    }
>>
>> +    local_idx = thread_idx.y*16 + thread_idx.x;
>>
>> +    if (local_idx < 128)
>>
>> +        lc[local_idx] = mask[local_idx];
>>
>> +    barrier(CLK_LOCAL_MEM_FENCE);
>>
>> +
>>
>> +    \n#pragma unroll\n
>>
>> +    for (i = -4; i <= 4; i++) {
>>
>> +        lm_idx.y = 8 + i + thread_idx.y;
>>
>> +        \n#pragma unroll\n
>>
>> +        for (j = -4; j <= 4; j++) {
>>
>> +            lm_idx.x = 8 + j + thread_idx.x;
>>
>> +            lc_idx = (i + 4)*8 + j + 4;
>>
>> +            sum += (int)data[lm_idx.y][lm_idx.x] * lc[lc_idx];
>>
>> +        }
>>
>> +    }
>>
>> +    int temp = (int)data[thread_idx.y + 8][thread_idx.x + 8];
>>
>> +    int res = temp + (((temp - (int)((sum + 1<<15) >> 16))) >> 16);
>>
>> +    if (global_idx.x < width && global_idx.y < height)
>>
>> +        dst[global_idx.x + global_idx.y*width] = clip_uint8(res);
>>
>> +}
>>
>> +);
>>
>> +
>>
>> +#endif
>>
>> diff --git a/doc/fftools-common-opts.texi b/doc/fftools-common-opts.texi
>>
>> index 617af2f..d63b3d1 100644
>>
>> --- a/doc/fftools-common-opts.texi
>>
>> +++ b/doc/fftools-common-opts.texi
>>
>> @@ -250,6 +250,10 @@ Possible flags for this option are:
>>
>>  @end table
>>
>>  @end table
>>
>>
>>
>> + at item -opencl_bench
>>
>> +Benchmark available OpenCL devices and output the results. This option
>>
>> +is only available when FFmpeg has been compiled with
>>> @code{--enable-opencl}.
>>
>> +
>>
>>  @item -opencl_options options (@emph{global})
>>
>>  Set OpenCL environment options. This option is only available when
>>
>>  FFmpeg has been compiled with @code{--enable-opencl}.
>>
>> diff --git a/doc/utils.texi b/doc/utils.texi
>>
>> index 9377139..3055662 100644
>>
>> --- a/doc/utils.texi
>>
>> +++ b/doc/utils.texi
>>
>> @@ -1051,13 +1051,13 @@ See reference "OpenCL Specification Version: 1.2
>>> chapter 5.6.4".
>>
>>  Select the index of the platform to run OpenCL code.
>>
>>
>>
>>  The specified index must be one of the indexes in the device list
>>
>> -which can be obtained with @code{av_opencl_get_device_list()}.
>>
>> +which can be obtained with @code{ffmpeg -opencl_bench} or
>>> @code{av_opencl_get_device_list()}.
>>
>>
>>
>>  @item device_idx
>>
>>  Select the index of the device used to run OpenCL code.
>>
>>
>>
>>  The specifed index must be one of the indexes in the device list which
>>
>> -can be obtained with @code{av_opencl_get_device_list()}.
>>
>> +can be obtained with @code{ffmpeg -opencl_bench} or
>>> @code{av_opencl_get_device_list()}.
>>
>>
>>
>>  @end table
>>
>>
>>
>> diff --git a/libavutil/opencl.c b/libavutil/opencl.c
>>
>> index ae4c476..ed723f9 100644
>>
>> --- a/libavutil/opencl.c
>>
>> +++ b/libavutil/opencl.c
>>
>> @@ -761,3 +761,45 @@ int av_opencl_buffer_read_image(uint8_t **dst_data,
>>> int *plane_size, int plane_n
>>
>>      }
>>
>>      return 0;
>>
>>  }
>>
>> +
>>
>> +int64_t av_opencl_benchmark(AVOpenCLDeviceNode *device_node,
>>> cl_platform_id platform,
>>
>> +                            int64_t (*benchmark)(AVOpenCLExternalEnv
>>> *ext_opencl_env))
>>
>> +{
>>
>> +    int64_t ret = 0;
>>
>> +    cl_int status;
>>
>> +    cl_context_properties cps[3];
>>
>> +    AVOpenCLExternalEnv *ext_opencl_env = NULL;
>>
>> +
>>
>> +    ext_opencl_env = av_opencl_alloc_external_env();
>>
>> +    ext_opencl_env->device_id = device_node->device_id;
>>
>> +    ext_opencl_env->device_type = device_node->device_type;
>>
>> +    av_log(&opencl_ctx, AV_LOG_VERBOSE, "Performing test on OpenCL device
>>> %s\n",
>>
>> +           device_node->device_name);
>>
>> +
>>
>> +    cps[0] = CL_CONTEXT_PLATFORM;
>>
>> +    cps[1] = (cl_context_properties)platform;
>>
>> +    cps[2] = 0;
>>
>> +    ext_opencl_env->context = clCreateContextFromType(cps,
>>> ext_opencl_env->device_type,
>>
>> +                                                      NULL, NULL,
>>> &status);
>>
>> +    if (status != CL_SUCCESS || !ext_opencl_env->context) {
>>
>> +        ret = AVERROR_EXTERNAL;
>>
>> +        goto end;
>>
>> +    }
>>
>> +    ext_opencl_env->command_queue =
>>> clCreateCommandQueue(ext_opencl_env->context,
>>
>> +
>>> ext_opencl_env->device_id, 0, &status);
>>
>> +    if (status != CL_SUCCESS || !ext_opencl_env->command_queue) {
>>
>> +        ret = AVERROR_EXTERNAL;
>>
>> +        goto end;
>>
>> +    }
>>
>> +    ret = benchmark(ext_opencl_env);
>>
>> +    if (ret < 0)
>>
>> +        av_log(&opencl_ctx, AV_LOG_ERROR, "Benchmark failed with OpenCL
>>> device %s\n",
>>
>> +               device_node->device_name);
>>
>> +end:
>>
>> +    if (ext_opencl_env->command_queue)
>>
>> +        clReleaseCommandQueue(ext_opencl_env->command_queue);
>>
>> +    if (ext_opencl_env->context)
>>
>> +        clReleaseContext(ext_opencl_env->context);
>>
>> +    av_opencl_free_external_env(&ext_opencl_env);
>>
>> +    return ret;
>>
>> +}
>>
>> diff --git a/libavutil/opencl.h b/libavutil/opencl.h
>>
>> index e4ecbf8..fd31e5e 100644
>>
>> --- a/libavutil/opencl.h
>>
>> +++ b/libavutil/opencl.h
>>
>> @@ -310,4 +310,19 @@ void av_opencl_release_kernel(AVOpenCLKernelEnv *env);
>>
>>   */
>>
>>  void av_opencl_uninit(void);
>>
>>
>>
>> +/**
>>
>> + * Benchmark an OpenCL device with a user defined callback function.
>>>  This function
>>
>> + * sets up an external OpenCL environment including context and command
>>> queue on
>>
>> + * the device then tears it down in the end.  The callback function
>>> should perform
>>
>> + * the rest of the work.
>>
>> + *
>>
>> + * @param device            pointer to the OpenCL device to be used
>>
>> + * @param platform          cl_platform_id handle to which the device
>>> belongs to
>>
>> + * @param benchmark         callback function to perform the benchmark,
>>> return a
>>
>> + *                          negative value in case of failure
>>
>> + * @return the score passed from the callback function, a negative value
>>> in case of failure
>>
>> + */
>>
>> +int64_t av_opencl_benchmark(AVOpenCLDeviceNode *device, cl_platform_id
>>> platform,
>>
>> +                            int64_t (*benchmark)(AVOpenCLExternalEnv
>>> *ext_opencl_env));
>>
>> +
>>
>>  #endif /* LIBAVUTIL_OPENCL_H */
>>
>>
>> seems an indent problem of clFinish.
>
> Looks good to me, thanks.
>

Indent fixed.  Sorry about that.  New patch attached.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: add-opencl-bench-option.patch
Type: application/octet-stream
Size: 17026 bytes
Desc: not available
URL: <http://ffmpeg.org/pipermail/ffmpeg-devel/attachments/20131202/582c7672/attachment.obj>


More information about the ffmpeg-devel mailing list