[FFmpeg-devel] [PATCH]opencl: automatically select the fastest opencl device
Stefano Sabatini
stefasab at gmail.com
Mon Dec 2 09:50:19 CET 2013
On date Monday 2013-12-02 01:05:57 -0600, Lenny Wang encoded:
> On Mon, Dec 2, 2013 at 12:19 AM, Wei Gao <highgod0401 at gmail.com> wrote:
[...]
> 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(-)
Missing minor bump and doc/APIchanges addition.
>
> 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);
> + }
nit:
av_free(inbuf);
av_log(...);
return ...
or
av_log(...);
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);
weird 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;
nit: OpenCLDeviceBenchmark
> +
> +/**
> + * 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);
nit: opt_opencl_bench()
> +#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
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));
> +
> #endif /* LIBAVUTIL_OPENCL_H */
No more comments from me.
--
FFmpeg = Fascinating and Fast Minimal Puristic Epic Gadget
More information about the ffmpeg-devel
mailing list