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 "cmdutils.h"
26 
27 typedef struct {
28  int platform_idx;
29  int device_idx;
30  char device_name[64];
31  int64_t runtime;
32 } OpenCLDeviceBenchmark;
33 
34 const char *ocl_bench_source = AV_OPENCL_KERNEL(
35 inline unsigned char clip_uint8(int a)
36 {
37  if (a & (~0xFF))
38  return (-a)>>31;
39  else
40  return a;
41 }
42 
43 kernel void unsharp_bench(
44  global unsigned char *src,
45  global unsigned char *dst,
46  global int *mask,
47  int width,
48  int height)
49 {
50  int i, j, local_idx, lc_idx, sum = 0;
51  int2 thread_idx, block_idx, global_idx, lm_idx;
52  thread_idx.x = get_local_id(0);
53  thread_idx.y = get_local_id(1);
54  block_idx.x = get_group_id(0);
55  block_idx.y = get_group_id(1);
56  global_idx.x = get_global_id(0);
57  global_idx.y = get_global_id(1);
58  local uchar data[32][32];
59  local int lc[128];
60 
61  for (i = 0; i <= 1; i++) {
62  lm_idx.y = -8 + (block_idx.y + i) * 16 + thread_idx.y;
63  lm_idx.y = lm_idx.y < 0 ? 0 : lm_idx.y;
64  lm_idx.y = lm_idx.y >= height ? height - 1: lm_idx.y;
65  for (j = 0; j <= 1; j++) {
66  lm_idx.x = -8 + (block_idx.x + j) * 16 + thread_idx.x;
67  lm_idx.x = lm_idx.x < 0 ? 0 : lm_idx.x;
68  lm_idx.x = lm_idx.x >= width ? width - 1: lm_idx.x;
69  data[i*16 + thread_idx.y][j*16 + thread_idx.x] = src[lm_idx.y*width + lm_idx.x];
70  }
71  }
72  local_idx = thread_idx.y*16 + thread_idx.x;
73  if (local_idx < 128)
74  lc[local_idx] = mask[local_idx];
75  barrier(CLK_LOCAL_MEM_FENCE);
76 
77  \n#pragma unroll\n
78  for (i = -4; i <= 4; i++) {
79  lm_idx.y = 8 + i + thread_idx.y;
80  \n#pragma unroll\n
81  for (j = -4; j <= 4; j++) {
82  lm_idx.x = 8 + j + thread_idx.x;
83  lc_idx = (i + 4)*8 + j + 4;
84  sum += (int)data[lm_idx.y][lm_idx.x] * lc[lc_idx];
85  }
86  }
87  int temp = (int)data[thread_idx.y + 8][thread_idx.x + 8];
88  int res = temp + (((temp - (int)((sum + 1<<15) >> 16))) >> 16);
89  if (global_idx.x < width && global_idx.y < height)
90  dst[global_idx.x + global_idx.y*width] = clip_uint8(res);
91 }
92 );
93 
94 #define OCLCHECK(method, ... ) \
95 do { \
96  status = method(__VA_ARGS__); \
97  if (status != CL_SUCCESS) { \
98  av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n", \
99  av_opencl_errstr(status)); \
100  ret = AVERROR_EXTERNAL; \
101  goto end; \
102  } \
103 } while (0)
104 
105 #define CREATEBUF(out, flags, size) \
106 do { \
107  out = clCreateBuffer(ext_opencl_env->context, flags, size, NULL, &status); \
108  if (status != CL_SUCCESS) { \
109  av_log(NULL, AV_LOG_ERROR, "Could not create OpenCL buffer\n"); \
110  ret = AVERROR_EXTERNAL; \
111  goto end; \
112  } \
113 } while (0)
114 
115 static void fill_rand_int(int *data, int n)
116 {
117  int i;
118  srand(av_gettime());
119  for (i = 0; i < n; i++)
120  data[i] = rand();
121 }
122 
123 #define OPENCL_NB_ITER 5
124 static int64_t run_opencl_bench(AVOpenCLExternalEnv *ext_opencl_env)
125 {
126  int i, arg = 0, width = 1920, height = 1088;
127  int64_t start, ret = 0;
128  cl_int status;
129  size_t kernel_len;
130  char *inbuf;
131  int *mask;
132  int buf_size = width * height * sizeof(char);
133  int mask_size = sizeof(uint32_t) * 128;
134 
135  cl_mem cl_mask, cl_inbuf, cl_outbuf;
136  cl_kernel kernel = NULL;
137  cl_program program = NULL;
138  size_t local_work_size_2d[2] = {16, 16};
139  size_t global_work_size_2d[2] = {(size_t)width, (size_t)height};
140 
141  if (!(inbuf = av_malloc(buf_size)) || !(mask = av_malloc(mask_size))) {
142  av_log(NULL, AV_LOG_ERROR, "Out of memory\n");
143  ret = AVERROR(ENOMEM);
144  goto end;
145  }
146  fill_rand_int((int*)inbuf, buf_size/4);
147  fill_rand_int(mask, mask_size/4);
148 
149  CREATEBUF(cl_mask, CL_MEM_READ_ONLY, mask_size);
150  CREATEBUF(cl_inbuf, CL_MEM_READ_ONLY, buf_size);
151  CREATEBUF(cl_outbuf, CL_MEM_READ_WRITE, buf_size);
152 
153  kernel_len = strlen(ocl_bench_source);
154  program = clCreateProgramWithSource(ext_opencl_env->context, 1, &ocl_bench_source,
155  &kernel_len, &status);
156  if (status != CL_SUCCESS || !program) {
157  av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark program\n");
158  ret = AVERROR_EXTERNAL;
159  goto end;
160  }
161  status = clBuildProgram(program, 1, &(ext_opencl_env->device_id), NULL, NULL, NULL);
162  if (status != CL_SUCCESS) {
163  av_log(NULL, AV_LOG_ERROR, "OpenCL unable to build benchmark program\n");
164  ret = AVERROR_EXTERNAL;
165  goto end;
166  }
167  kernel = clCreateKernel(program, "unsharp_bench", &status);
168  if (status != CL_SUCCESS) {
169  av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark kernel\n");
170  ret = AVERROR_EXTERNAL;
171  goto end;
172  }
173 
174  OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_inbuf, CL_TRUE, 0,
175  buf_size, inbuf, 0, NULL, NULL);
176  OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_mask, CL_TRUE, 0,
177  mask_size, mask, 0, NULL, NULL);
178  OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_inbuf);
179  OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_outbuf);
180  OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_mask);
181  OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &width);
182  OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &height);
183 
184  start = av_gettime();
185  for (i = 0; i < OPENCL_NB_ITER; i++)
186  OCLCHECK(clEnqueueNDRangeKernel, ext_opencl_env->command_queue, kernel, 2, NULL,
187  global_work_size_2d, local_work_size_2d, 0, NULL, NULL);
188  clFinish(ext_opencl_env->command_queue);
189  ret = (av_gettime() - start)/OPENCL_NB_ITER;
190 end:
191  if (kernel)
192  clReleaseKernel(kernel);
193  if (program)
194  clReleaseProgram(program);
195  if (cl_inbuf)
196  clReleaseMemObject(cl_inbuf);
197  if (cl_outbuf)
198  clReleaseMemObject(cl_outbuf);
199  if (cl_mask)
200  clReleaseMemObject(cl_mask);
201  av_free(inbuf);
202  av_free(mask);
203  return ret;
204 }
205 
206 static int compare_ocl_device_desc(const void *a, const void *b)
207 {
208  return ((OpenCLDeviceBenchmark*)a)->runtime - ((OpenCLDeviceBenchmark*)b)->runtime;
209 }
210 
211 int opt_opencl_bench(void *optctx, const char *opt, const char *arg)
212 {
213  int i, j, nb_devices = 0, count = 0;
214  int64_t score = 0;
215  AVOpenCLDeviceList *device_list;
216  AVOpenCLDeviceNode *device_node = NULL;
217  OpenCLDeviceBenchmark *devices = NULL;
218  cl_platform_id platform;
219 
220  av_opencl_get_device_list(&device_list);
221  for (i = 0; i < device_list->platform_num; i++)
222  nb_devices += device_list->platform_node[i]->device_num;
223  if (!nb_devices) {
224  av_log(NULL, AV_LOG_ERROR, "No OpenCL device detected!\n");
225  return AVERROR(EINVAL);
226  }
227  if (!(devices = av_malloc(sizeof(OpenCLDeviceBenchmark) * nb_devices))) {
228  av_log(NULL, AV_LOG_ERROR, "Could not allocate buffer\n");
229  return AVERROR(ENOMEM);
230  }
231 
232  for (i = 0; i < device_list->platform_num; i++) {
233  for (j = 0; j < device_list->platform_node[i]->device_num; j++) {
234  device_node = device_list->platform_node[i]->device_node[j];
235  platform = device_list->platform_node[i]->platform_id;
236  score = av_opencl_benchmark(device_node, platform, run_opencl_bench);
237  if (score > 0) {
238  devices[count].platform_idx = i;
239  devices[count].device_idx = j;
240  devices[count].runtime = score;
241  strcpy(devices[count].device_name, device_node->device_name);
242  count++;
243  }
244  }
245  }
246  qsort(devices, count, sizeof(OpenCLDeviceBenchmark), compare_ocl_device_desc);
247  fprintf(stderr, "platform_idx\tdevice_idx\tdevice_name\truntime\n");
248  for (i = 0; i < count; i++)
249  fprintf(stdout, "%d\t%d\t%s\t%"PRId64"\n",
250  devices[i].platform_idx, devices[i].device_idx,
251  devices[i].device_name, devices[i].runtime);
252 
253  av_opencl_free_device_list(&device_list);
254  av_free(devices);
255  return 0;
256 }
257 
258 int opt_opencl(void *optctx, const char *opt, const char *arg)
259 {
260  char *key, *value;
261  const char *opts = arg;
262  int ret = 0;
263  while (*opts) {
264  ret = av_opt_get_key_value(&opts, "=", ":", 0, &key, &value);
265  if (ret < 0)
266  return ret;
267  ret = av_opencl_set_option(key, value);
268  if (ret < 0)
269  return ret;
270  if (*opts)
271  opts++;
272  }
273  return ret;
274 }
Generated on Sun Mar 23 2014 23:49:48 for FFmpeg by   doxygen 1.8.2