[FFmpeg-devel] [PATCH]opencl: automatically select the fastest opencl device
Lenny Wang
lenny at multicorewareinc.com
Tue Dec 3 07:08:17 CET 2013
On Mon, Dec 2, 2013 at 11:03 PM, Michael Niedermayer <michaelni at gmx.at> wrote:
> On Mon, Dec 02, 2013 at 10:39:20PM -0600, Lenny Wang wrote:
>> On Mon, Dec 2, 2013 at 10:12 PM, Michael Niedermayer <michaelni at gmx.at> wrote:
>> > On Mon, Dec 02, 2013 at 04:48:41PM -0600, Lenny Wang wrote:
>> >> On Mon, Dec 2, 2013 at 4:40 PM, Stefano Sabatini <stefasab at gmail.com> wrote:
>> >> > On date Monday 2013-12-02 11:46:09 -0600, Lenny Wang encoded:
>> >> >> On Mon, Dec 2, 2013 at 5:46 AM, Stefano Sabatini <stefasab at gmail.com> wrote:
>> >> > [...]
>> >> >> From: Lenny Wang <lwanghpc at gmail.com>
>> >> >> Date: Mon, 2 Dec 2013 11:40:00 -0600
>> >> >> Subject: [PATCH] cmdutils & opencl: add -opencl_bench option to test and show available OpenCL devices
>> >> >>
>> >> >> ---
>> >> >> cmdutils.c | 166 +++++++++++++++++++++++++++++++++++++++++
>> >> >> cmdutils.h | 15 ++++
>> >> >> cmdutils_common_opts.h | 1 +
>> >> >> cmdutils_opencl_bench_kernel.h | 86 +++++++++++++++++++++
>> >> >> 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 +-
>> >> >> 10 files changed, 335 insertions(+), 3 deletions(-)
>> >> >>
>> >> >> diff --git a/cmdutils.c b/cmdutils.c
>> >> >> index 46ade3f..4f61a80 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,170 @@ 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");
>> >> >> + ret = AVERROR(ENOMEM);
>> >> >> + goto end;
>> >> >> + }
>> >> >> + 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);
>> >> >
>> >> > nit: still weird indent, it should be:
>> >> >
>> >> > 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);
>> >> >
>> >> > LGTM, but Wei should approve the patch, thanks.
>> >>
>> >> Indent fixed. Pending Wei's approval.
>> >
>> >> cmdutils.c | 166 +++++++++++++++++++++++++++++++++++++++++
>> >> cmdutils.h | 15 +++
>> >> cmdutils_common_opts.h | 1
>> >> cmdutils_opencl_bench_kernel.h | 86 +++++++++++++++++++++
>> >> 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
>> >> 10 files changed, 335 insertions(+), 3 deletions(-)
>> >> 6fccaf0e1c18908b2ce57bb460749dfe14824831 add-opencl-bench-option.patch
>> >> From: Lenny Wang <lwanghpc at gmail.com>
>> >> Date: Mon, 2 Dec 2013 11:40:00 -0600
>> >> Subject: [PATCH] cmdutils & opencl: add -opencl_bench option to test and show available OpenCL devices
>> >>
>> >> ---
>> >> cmdutils.c | 166 +++++++++++++++++++++++++++++++++++++++++
>> >> cmdutils.h | 15 ++++
>> >> cmdutils_common_opts.h | 1 +
>> >> cmdutils_opencl_bench_kernel.h | 86 +++++++++++++++++++++
>> >> 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 +-
>> >> 10 files changed, 335 insertions(+), 3 deletions(-)
>> >>
>> >> diff --git a/cmdutils.c b/cmdutils.c
>> >> index 46ade3f..4f61a80 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,170 @@ 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");
>> >> + ret = AVERROR(ENOMEM);
>> >> + goto end;
>> >> + }
>> >> + 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);
>> >> + 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 ((OpenCLDeviceBenchmark*)a)->runtime - ((OpenCLDeviceBenchmark*)b)->runtime;
>> >> +}
>> >> +
>> >> +int opt_opencl_bench(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;
>> >> + OpenCLDeviceBenchmark *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(OpenCLDeviceBenchmark) * 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(OpenCLDeviceBenchmark), compare_ocl_device_desc);
>> >> + fprintf(stderr, "platform_idx\tdevice_idx\tdevice_name\truntime\n");
>> >> + for (i = 0; i < count; i++)
>> >> + fprintf(stdout, "%d\t%d\t%s\t%"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..bb60a26 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;
>> >> +} OpenCLDeviceBenchmark;
>> >> +
>> >> +/**
>> >> + * Perform benchmark on all OpenCL devices and print results.
>> >> + * This option processing function does not utilize the arguments.
>> >> + */
>> >> +int opt_opencl_bench(void *optctx, const char *opt, const char *arg);
>> >> +#endif
>> >> +
>> >> /**
>> >> * Limit the execution time.
>> >> */
>> >
>> > why is this in cmdutils.c/h ?
>> >
>>
>> It's designed and used as a cmd utility option similar to
>> -codes,-filters etc, it lists available opencl devices sorted based on
>> benchmark scores, so users can pick the fastest device using
>> '-opencl_option'.
>
> -codecs, filters, ... list libavcodec and filter codecs and filters
> that are things registered with these libs.
> the code here lists hardware devices and benchmarks them.
> we dont put dshow device extraction and benchmarking code in cmdutils
> either. nor alsa device listing, and benchmarking code.
>
> cmdutils.c would turn in a total mess if everthing from alsa over
> oss, dshow, caca, v4l to vfw would have listing code under #ifs in
> it
>
How about I put the implementation in /tools, and with only minimal
code in cmdutils (cmdutils currently has "-opencl_options"
implemented). It's a very useful feature and it'd be helpful
integrated into ffmpeg as a command line option.
More information about the ffmpeg-devel
mailing list