[FFmpeg-devel] [PATCH]opencl: automatically select the fastest opencl device
Wei Gao
highgod0401 at gmail.com
Tue Dec 3 02:15:23 CET 2013
2013/12/3 Lenny Wang <lenny at multicorewareinc.com>
> 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.
>
LGTM, Thanks.
>
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel at ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
>
More information about the ffmpeg-devel
mailing list