summaryrefslogtreecommitdiff
path: root/cmdutils_opencl.c (plain)
blob: 655d1c9546c1efebd552a26ff79b49fa564df4b6
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
28typedef struct {
29 int platform_idx;
30 int device_idx;
31 char device_name[64];
32 int64_t runtime;
33} OpenCLDeviceBenchmark;
34
35const char *ocl_bench_source = AV_OPENCL_KERNEL(
36inline unsigned char clip_uint8(int a)
37{
38 if (a & (~0xFF))
39 return (-a)>>31;
40 else
41 return a;
42}
43
44kernel 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, ... ) \
96do { \
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) \
107do { \
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
116static 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
125static 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;
191end:
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
207static 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
214int 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
267int 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}
284
generated by cgit at 2024-05-02 10:33:31 (GMT)