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

[Lenny Wang]

On Sun, Dec 8, 2013 at 8:13 PM, Michael Niedermayer <michaelni at gmx.at> wrote:
> On Fri, Dec 06, 2013 at 02:35:19PM -0600, Lenny Wang wrote:
>> On Fri, Dec 6, 2013 at 4:57 AM, Stefano Sabatini <stefasab at gmail.com> wrote:
>> > On date Friday 2013-12-06 03:00:09 -0600, Lenny Wang encoded:
>> >> >
>> >> > please move the opencl code from cmdutils.c into a seperate file
>> >> > its up to you and wei if you want the benchmarking and or listing
>> >> > code available from the ffmpeg tool or in a seperate tool. but
>> >> > opencl code should not be put in random files under ifdef, put it
>> >> > in seperate opencl specific files. build and link them only when
>> >> > opencl is available
>> >> >
>> >>
>> >> Sorry about the late reply.  Please review attached patch modified
>> >> based on Michael's suggestion, thanks.
>> >
>> >> From: Lenny Wang <lwanghpc at gmail.com>
>> >> Date: Fri, 6 Dec 2013 02:40:00 -0600
>> >> Subject: [PATCH] cmdutils & opencl: add -opencl_bench option to test and show available OpenCL devices
>> >>
>> >> ---
>> >>  Makefile                     |   1 +
>> >>  cmdutils.c                   |  54 +++++++++++-
>> >>  cmdutils.h                   |   6 ++
>> >>  cmdutils_common_opts.h       |   1 +
>> >>  doc/APIchanges               |   2 +
>> >>  doc/fftools-common-opts.texi |   4 +
>> >>  doc/utils.texi               |   4 +-
>> >>  libavutil/opencl.c           |  42 ++++++++++
>> >>  libavutil/opencl.h           |  16 ++++
>> >>  libavutil/version.h          |   2 +-
>> >>  tools/opencl_bench.c         | 194 +++++++++++++++++++++++++++++++++++++++++++
>> >>  tools/opencl_bench.h         |  37 +++++++++
>> >>  12 files changed, 359 insertions(+), 4 deletions(-)
>> >
>> > Uh, not yet.
>> >
>> > So, Michael suggested to move all cmdutils OpenCL code to a separate
>> > file. The tools directory is meant for self-contained tools, you
>> > should not depend on ff*/cmdutils code or even worse make ff*/cmdutils
>> > code depend on it.
>> >
>> > So I suggest this: you move all cmdutils OpenCL code to
>> > cmdutils_opencl.c/h. cmdutils_opencl.o will be linked only if OpenCL
>> > is detected. The opencl_options code should go in the same file as
>> > well (in the same or maybe in a separate follow-up patch).
>> >
>> > Function prototypes can go to a dedicated cmdutils_opencl.h file or
>> > stay in cmdutils.h, under #ifdef, as Michael prefers.
>>
>> Fixed based on Stefano's suggestion.
>
>>  Makefile                     |    1
>>  cmdutils.c                   |   24 ---
>>  cmdutils.h                   |    4
>>  cmdutils_common_opts.h       |    1
>>  cmdutils_opencl.c            |  274 +++++++++++++++++++++++++++++++++++++++++++
>>  doc/APIchanges               |    2
>>  doc/fftools-common-opts.texi |    4
>>  doc/utils.texi               |    4
>>  libavutil/opencl.c           |   42 ++++++
>>  libavutil/opencl.h           |   16 ++
>>  10 files changed, 346 insertions(+), 26 deletions(-)
>> 79e5ef131086be7343651bede013cea722dcaaf2  cmdutils_opencl.patch
>> From: Lenny Wang <lwanghpc at gmail.com>
>> Date: Fri, 6 Dec 2013 14:20:00 -0600
>> Subject: [PATCH] cmdutils & opencl: add cmdutils_opencl.c with new -opencl_bench option
> [...]
>
>> +    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);
>> +    ret = (av_gettime() - start)/OPENCL_NB_ITER;
>
> tabs are forbidden in ffmpeg git
>
>
> [...]
>> diff --git a/doc/APIchanges b/doc/APIchanges
>> index 08ba47f..5b0de25 100644
>> --- a/doc/APIchanges
>> +++ b/doc/APIchanges
>> @@ -14,6 +14,8 @@ libavutil:     2012-10-22
>>
>>
>>  API changes, most recent first:
>> +2013-12-xx - xxxxxxx - lavu 52.57.100 - opencl.h
>> +  Add av_opencl_benchmark() function.
>
> version.h bump missing
>
>
> [...]
>> diff --git a/libavutil/opencl.h b/libavutil/opencl.h
>> index e4ecbf8..cf0abd7 100644
>> --- a/libavutil/opencl.h
>> +++ b/libavutil/opencl.h
>> @@ -310,4 +310,20 @@ 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 error code in case
>> + * of failure
>> + */
>> +int64_t av_opencl_benchmark(AVOpenCLDeviceNode *device, cl_platform_id platform,
>> +                            int64_t (*benchmark)(AVOpenCLExternalEnv *ext_opencl_env));
>
> why is this in libavutil ?
> that is why is the benchmark code split between cmdutils_opencl
> and libavutil ?
>

We decided to make the benchmark API public so other library
developers can also take advantage of it.  Stefano and Wei agreed with
the design in previous discussions.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: add-opencl-bench-option.patch
Type: application/octet-stream
Size: 19228 bytes
Desc: not available
URL: <http://ffmpeg.org/pipermail/ffmpeg-devel/attachments/20131208/c8adfbb8/attachment.obj>