FFmpeg
 All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups Pages
cmdutils_opencl.c
Go to the documentation of this file.
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 {
31  char device_name[64];
32  int64_t runtime;
34 
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;
133  int buf_size = width * height * sizeof(char);
134  int mask_size = sizeof(uint32_t) * 128;
135 
136  cl_mem cl_mask, cl_inbuf, cl_outbuf;
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  return ((OpenCLDeviceBenchmark*)a)->runtime - ((OpenCLDeviceBenchmark*)b)->runtime;
210 }
211 
212 int opt_opencl_bench(void *optctx, const char *opt, const char *arg)
213 {
214  int i, j, nb_devices = 0, count = 0;
215  int64_t score = 0;
216  AVOpenCLDeviceList *device_list;
217  AVOpenCLDeviceNode *device_node = NULL;
218  OpenCLDeviceBenchmark *devices = NULL;
219  cl_platform_id platform;
220 
221  av_opencl_get_device_list(&device_list);
222  for (i = 0; i < device_list->platform_num; i++)
223  nb_devices += device_list->platform_node[i]->device_num;
224  if (!nb_devices) {
225  av_log(NULL, AV_LOG_ERROR, "No OpenCL device detected!\n");
226  return AVERROR(EINVAL);
227  }
228  if (!(devices = av_malloc_array(nb_devices, sizeof(OpenCLDeviceBenchmark)))) {
229  av_log(NULL, AV_LOG_ERROR, "Could not allocate buffer\n");
230  return AVERROR(ENOMEM);
231  }
232 
233  for (i = 0; i < device_list->platform_num; i++) {
234  for (j = 0; j < device_list->platform_node[i]->device_num; j++) {
235  device_node = device_list->platform_node[i]->device_node[j];
236  platform = device_list->platform_node[i]->platform_id;
237  score = av_opencl_benchmark(device_node, platform, run_opencl_bench);
238  if (score > 0) {
239  devices[count].platform_idx = i;
240  devices[count].device_idx = j;
241  devices[count].runtime = score;
242  av_strlcpy(devices[count].device_name, device_node->device_name,
243  sizeof(devices[count].device_name));
244  count++;
245  }
246  }
247  }
248  qsort(devices, count, sizeof(OpenCLDeviceBenchmark), compare_ocl_device_desc);
249  fprintf(stderr, "platform_idx\tdevice_idx\tdevice_name\truntime\n");
250  for (i = 0; i < count; i++)
251  fprintf(stdout, "%d\t%d\t%s\t%"PRId64"\n",
252  devices[i].platform_idx, devices[i].device_idx,
253  devices[i].device_name, devices[i].runtime);
254 
255  av_opencl_free_device_list(&device_list);
256  av_free(devices);
257  return 0;
258 }
259 
260 int opt_opencl(void *optctx, const char *opt, const char *arg)
261 {
262  char *key, *value;
263  const char *opts = arg;
264  int ret = 0;
265  while (*opts) {
266  ret = av_opt_get_key_value(&opts, "=", ":", 0, &key, &value);
267  if (ret < 0)
268  return ret;
269  ret = av_opencl_set_option(key, value);
270  if (ret < 0)
271  return ret;
272  if (*opts)
273  opts++;
274  }
275  return ret;
276 }
#define NULL
Definition: coverity.c:32
#define AV_OPENCL_KERNEL(...)
Definition: opencl.h:45
ptrdiff_t const GLvoid * data
Definition: opengl_enc.c:101
else temp
Definition: vf_mcdeint.c:257
const char * b
Definition: vf_curves.c:109
#define OCLCHECK(method,...)
#define CREATEBUF(out, flags, size)
if()
Definition: avfilter.c:975
#define av_malloc(s)
AVOptions.
static void fill_rand_int(int *data, int n)
static av_cold int end(AVCodecContext *avctx)
Definition: avrndec.c:67
cl_platform_id platform_id
Definition: opencl.h:54
#define av_log(a,...)
int av_opencl_get_device_list(AVOpenCLDeviceList **device_list)
Get OpenCL device list.
Definition: opencl.c:330
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
static const uint16_t mask[17]
Definition: lzw.c:38
#define AVERROR(e)
Definition: error.h:43
static int compare_ocl_device_desc(const void *a, const void *b)
const char * arg
Definition: jacosubdec.c:66
cl_device_id device_id
Definition: opencl.h:69
GLsizei count
Definition: opengl_enc.c:109
size_t av_strlcpy(char *dst, const char *src, size_t size)
Copy the string src to dst, but no more than size - 1 bytes, and null-terminate dst.
Definition: avstring.c:83
AVOpenCLDeviceNode ** device_node
Definition: opencl.h:57
int opt_opencl(void *optctx, const char *opt, const char *arg)
ret
Definition: avfilter.c:974
static int64_t run_opencl_bench(AVOpenCLExternalEnv *ext_opencl_env)
GLsizei GLboolean const GLfloat * value
Definition: opengl_enc.c:109
int n
Definition: avisynth_c.h:547
int av_opencl_set_option(const char *key, const char *val)
Set option in the global OpenCL context.
Definition: opencl.c:375
int64_t av_gettime(void)
Get the current time in microseconds.
Definition: time.c:39
AVOpenCLPlatformNode ** platform_node
Definition: opencl.h:62
AVS_Value src
Definition: avisynth_c.h:482
#define OPENCL_NB_ITER
BYTE int const BYTE int int int height
Definition: avisynth_c.h:676
cl_command_queue command_queue
Definition: opencl.h:70
char * device_name
Definition: opencl.h:49
int64_t av_gettime_relative(void)
Get the current time in microseconds since some unspecified starting point.
Definition: time.c:56
void av_opencl_free_device_list(AVOpenCLDeviceList **device_list)
Free OpenCL device list.
Definition: opencl.c:348
const char * ocl_bench_source
cl_context context
Definition: opencl.h:68
#define av_free(p)
int av_opt_get_key_value(const char **ropts, const char *key_val_sep, const char *pairs_sep, unsigned flags, char **rkey, char **rval)
Extract a key-value pair from the beginning of a string.
Definition: opt.c:1361
int64_t av_opencl_benchmark(AVOpenCLDeviceNode *device_node, cl_platform_id platform, int64_t(*benchmark)(AVOpenCLExternalEnv *ext_opencl_env))
Benchmark an OpenCL device with a user defined callback function.
Definition: opencl.c:813
OpenCL wrapper.
void INT64 start
Definition: avisynth_c.h:553
#define av_malloc_array(a, b)
#define AVERROR_EXTERNAL
Generic error in an external library.
Definition: error.h:57
for(j=16;j >0;--j)
int opt_opencl_bench(void *optctx, const char *opt, const char *arg)
static int width