cmdutils_opencl.c

   1 /*
   2  * Copyright (C) 2013 Lenny Wang
   3  *
   4  * This file is part of FFmpeg.
   5  *
   6  * FFmpeg is free software; you can redistribute it and/or
   7  * modify it under the terms of the GNU Lesser General Public
   8  * License as published by the Free Software Foundation; either
   9  * version 2.1 of the License, or (at your option) any later version.
  10  *
  11  * FFmpeg is distributed in the hope that it will be useful,
  12  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  13  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  14  * Lesser General Public License for more details.
  15  *
  16  * You should have received a copy of the GNU Lesser General Public
  17  * License along with FFmpeg; if not, write to the Free Software
  18  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  19  */
  20
  21 #include "libavutil/opt.h"
  22 #include "libavutil/time.h"
  23 #include "libavutil/log.h"
  24 #include "libavutil/opencl.h"
  25 #include "libavutil/avstring.h"
  26 #include "cmdutils.h"
  27
  28 typedef struct {
  29     int platform_idx;
  30     int device_idx;
  31     char device_name[64];
  32     int64_t runtime;
  33 } OpenCLDeviceBenchmark;
  34
  35 const char *ocl_bench_source = AV_OPENCL_KERNEL(
  36 inline unsigned char clip_uint8(int a)
  37 {
  38     if (a & (~0xFF))
  39         return (-a)>>31;
  40     else
  41         return a;
  42 }
  43
  44 kernel void unsharp_bench(
  45                     global unsigned char *src,
  46                     global unsigned char *dst,
  47                     global int *mask,
  48                     int width,
  49                     int height)
  50 {
  51     int i, j, local_idx, lc_idx, sum = 0;
  52     int2 thread_idx, block_idx, global_idx, lm_idx;
  53     thread_idx.x = get_local_id(0);
  54     thread_idx.y = get_local_id(1);
  55     block_idx.x = get_group_id(0);
  56     block_idx.y = get_group_id(1);
  57     global_idx.x = get_global_id(0);
  58     global_idx.y = get_global_id(1);
  59     local uchar data[32][32];
  60     local int lc[128];
  61
  62     for (i = 0; i <= 1; i++) {
  63         lm_idx.y = -8 + (block_idx.y + i) * 16 + thread_idx.y;
  64         lm_idx.y = lm_idx.y < 0 ? 0 : lm_idx.y;
  65         lm_idx.y = lm_idx.y >= height ? height - 1: lm_idx.y;
  66         for (j = 0; j <= 1; j++) {
  67             lm_idx.x = -8 + (block_idx.x + j) * 16 + thread_idx.x;
  68             lm_idx.x = lm_idx.x < 0 ? 0 : lm_idx.x;
  69             lm_idx.x = lm_idx.x >= width ? width - 1: lm_idx.x;
  70             data[i*16 + thread_idx.y][j*16 + thread_idx.x] = src[lm_idx.y*width + lm_idx.x];
  71         }
  72     }
  73     local_idx = thread_idx.y*16 + thread_idx.x;
  74     if (local_idx < 128)
  75         lc[local_idx] = mask[local_idx];
  76     barrier(CLK_LOCAL_MEM_FENCE);
  77
  78     \n#pragma unroll\n
  79     for (i = -4; i <= 4; i++) {
  80         lm_idx.y = 8 + i + thread_idx.y;
  81         \n#pragma unroll\n
  82         for (j = -4; j <= 4; j++) {
  83             lm_idx.x = 8 + j + thread_idx.x;
  84             lc_idx = (i + 4)*8 + j + 4;
  85             sum += (int)data[lm_idx.y][lm_idx.x] * lc[lc_idx];
  86         }
  87     }
  88     int temp = (int)data[thread_idx.y + 8][thread_idx.x + 8];
  89     int res = temp + (((temp - (int)((sum + 1<<15) >> 16))) >> 16);
  90     if (global_idx.x < width && global_idx.y < height)
  91         dst[global_idx.x + global_idx.y*width] = clip_uint8(res);
  92 }
  93 );
  94
  95 #define OCLCHECK(method, ... )                                                 \
  96 do {                                                                           \
  97     status = method(__VA_ARGS__);                                              \
  98     if (status != CL_SUCCESS) {                                                \
  99         av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n",                   \
 100                av_opencl_errstr(status));                                      \
 101         ret = AVERROR_EXTERNAL;                                                \
 102         goto end;                                                              \
 103     }                                                                          \
 104 } while (0)
 105
 106 #define CREATEBUF(out, flags, size)                                            \
 107 do {                                                                           \
 108     out = clCreateBuffer(ext_opencl_env->context, flags, size, NULL, &status); \
 109     if (status != CL_SUCCESS) {                                                \
 110         av_log(NULL, AV_LOG_ERROR, "Could not create OpenCL buffer\n");        \
 111         ret = AVERROR_EXTERNAL;                                                \
 112         goto end;                                                              \
 113     }                                                                          \
 114 } while (0)
 115
 116 static void fill_rand_int(int *data, int n)
 117 {
 118     int i;
 119     srand(av_gettime());
 120     for (i = 0; i < n; i++)
 121         data[i] = rand();
 122 }
 123
 124 #define OPENCL_NB_ITER 5
 125 static int64_t run_opencl_bench(AVOpenCLExternalEnv *ext_opencl_env)
 126 {
 127     int i, arg = 0, width = 1920, height = 1088;
 128     int64_t start, ret = 0;
 129     cl_int status;
 130     size_t kernel_len;
 131     char *inbuf;
 132     int *mask = NULL;
 133     int buf_size = width * height * sizeof(char);
 134     int mask_size = sizeof(uint32_t) * 128;
 135
 136     cl_mem cl_mask = NULL, cl_inbuf = NULL, cl_outbuf = NULL;
 137     cl_kernel kernel = NULL;
 138     cl_program program = NULL;
 139     size_t local_work_size_2d[2] = {16, 16};
 140     size_t global_work_size_2d[2] = {(size_t)width, (size_t)height};
 141
 142     if (!(inbuf = av_malloc(buf_size)) || !(mask = av_malloc(mask_size))) {
 143         av_log(NULL, AV_LOG_ERROR, "Out of memory\n");
 144         ret = AVERROR(ENOMEM);
 145         goto end;
 146     }
 147     fill_rand_int((int*)inbuf, buf_size/4);
 148     fill_rand_int(mask, mask_size/4);
 149
 150     CREATEBUF(cl_mask, CL_MEM_READ_ONLY, mask_size);
 151     CREATEBUF(cl_inbuf, CL_MEM_READ_ONLY, buf_size);
 152     CREATEBUF(cl_outbuf, CL_MEM_READ_WRITE, buf_size);
 153
 154     kernel_len = strlen(ocl_bench_source);
 155     program = clCreateProgramWithSource(ext_opencl_env->context, 1, &ocl_bench_source,
 156                                         &kernel_len, &status);
 157     if (status != CL_SUCCESS || !program) {
 158         av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark program\n");
 159         ret = AVERROR_EXTERNAL;
 160         goto end;
 161     }
 162     status = clBuildProgram(program, 1, &(ext_opencl_env->device_id), NULL, NULL, NULL);
 163     if (status != CL_SUCCESS) {
 164         av_log(NULL, AV_LOG_ERROR, "OpenCL unable to build benchmark program\n");
 165         ret = AVERROR_EXTERNAL;
 166         goto end;
 167     }
 168     kernel = clCreateKernel(program, "unsharp_bench", &status);
 169     if (status != CL_SUCCESS) {
 170         av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark kernel\n");
 171         ret = AVERROR_EXTERNAL;
 172         goto end;
 173     }
 174
 175     OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_inbuf, CL_TRUE, 0,
 176              buf_size, inbuf, 0, NULL, NULL);
 177     OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_mask, CL_TRUE, 0,
 178              mask_size, mask, 0, NULL, NULL);
 179     OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_inbuf);
 180     OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_outbuf);
 181     OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_mask);
 182     OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &width);
 183     OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &height);
 184
 185     start = av_gettime_relative();
 186     for (i = 0; i < OPENCL_NB_ITER; i++)
 187         OCLCHECK(clEnqueueNDRangeKernel, ext_opencl_env->command_queue, kernel, 2, NULL,
 188                  global_work_size_2d, local_work_size_2d, 0, NULL, NULL);
 189     clFinish(ext_opencl_env->command_queue);
 190     ret = (av_gettime_relative() - start)/OPENCL_NB_ITER;
 191 end:
 192     if (kernel)
 193         clReleaseKernel(kernel);
 194     if (program)
 195         clReleaseProgram(program);
 196     if (cl_inbuf)
 197         clReleaseMemObject(cl_inbuf);
 198     if (cl_outbuf)
 199         clReleaseMemObject(cl_outbuf);
 200     if (cl_mask)
 201         clReleaseMemObject(cl_mask);
 202     av_free(inbuf);
 203     av_free(mask);
 204     return ret;
 205 }
 206
 207 static int compare_ocl_device_desc(const void *a, const void *b)
 208 {
 209     const OpenCLDeviceBenchmark* va = (const OpenCLDeviceBenchmark*)a;
 210     const OpenCLDeviceBenchmark* vb = (const OpenCLDeviceBenchmark*)b;
 211     return FFDIFFSIGN(va->runtime , vb->runtime);
 212 }
 213
 214 int opt_opencl_bench(void *optctx, const char *opt, const char *arg)
 215 {
 216     int i, j, nb_devices = 0, count = 0, ret = 0;
 217     int64_t score = 0;
 218     AVOpenCLDeviceList *device_list;
 219     AVOpenCLDeviceNode *device_node = NULL;
 220     OpenCLDeviceBenchmark *devices = NULL;
 221     cl_platform_id platform;
 222
 223     ret = av_opencl_get_device_list(&device_list);
 224     if (ret < 0) {
 225         return ret;
 226     }
 227     for (i = 0; i < device_list->platform_num; i++)
 228         nb_devices += device_list->platform_node[i]->device_num;
 229     if (!nb_devices) {
 230         av_log(NULL, AV_LOG_ERROR, "No OpenCL device detected!\n");
 231         av_opencl_free_device_list(&device_list);
 232         return AVERROR(EINVAL);
 233     }
 234     if (!(devices = av_malloc_array(nb_devices, sizeof(OpenCLDeviceBenchmark)))) {
 235         av_log(NULL, AV_LOG_ERROR, "Could not allocate buffer\n");
 236         av_opencl_free_device_list(&device_list);
 237         return AVERROR(ENOMEM);
 238     }
 239
 240     for (i = 0; i < device_list->platform_num; i++) {
 241         for (j = 0; j < device_list->platform_node[i]->device_num; j++) {
 242             device_node = device_list->platform_node[i]->device_node[j];
 243             platform = device_list->platform_node[i]->platform_id;
 244             score = av_opencl_benchmark(device_node, platform, run_opencl_bench);
 245             if (score > 0) {
 246                 devices[count].platform_idx = i;
 247                 devices[count].device_idx = j;
 248                 devices[count].runtime = score;
 249                 av_strlcpy(devices[count].device_name, device_node->device_name,
 250                            sizeof(devices[count].device_name));
 251                 count++;
 252             }
 253         }
 254     }
 255     qsort(devices, count, sizeof(OpenCLDeviceBenchmark), compare_ocl_device_desc);
 256     fprintf(stderr, "platform_idx\tdevice_idx\tdevice_name\truntime\n");
 257     for (i = 0; i < count; i++)
 258         fprintf(stdout, "%d\t%d\t%s\t%"PRId64"\n",
 259                 devices[i].platform_idx, devices[i].device_idx,
 260                 devices[i].device_name, devices[i].runtime);
 261
 262     av_opencl_free_device_list(&device_list);
 263     av_free(devices);
 264     return 0;
 265 }
 266
 267 int opt_opencl(void *optctx, const char *opt, const char *arg)
 268 {
 269     char *key, *value;
 270     const char *opts = arg;
 271     int ret = 0;
 272     while (*opts) {
 273         ret = av_opt_get_key_value(&opts, "=", ":", 0, &key, &value);
 274         if (ret < 0)
 275             return ret;
 276         ret = av_opencl_set_option(key, value);
 277         if (ret < 0)
 278             return ret;
 279         if (*opts)
 280             opts++;
 281     }
 282     return ret;
 283 }