36 inline unsigned char clip_uint8(
int a)
44 kernel
void unsharp_bench(
45 global
unsigned char *
src,
46 global
unsigned char *dst,
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];
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];
73 local_idx = thread_idx.y*16 + thread_idx.x;
75 lc[local_idx] = mask[local_idx];
76 barrier(CLK_LOCAL_MEM_FENCE);
79 for (i = -4; i <= 4; i++) {
80 lm_idx.y = 8 + i + thread_idx.y;
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];
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);
95 #define OCLCHECK(method, ... ) \
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; \
106 #define CREATEBUF(out, flags, size) \
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; \
120 for (i = 0; i <
n; i++)
124 #define OPENCL_NB_ITER 5
127 int i,
arg = 0, width = 1920, height = 1088;
128 int64_t
start, ret = 0;
133 int buf_size = width * height *
sizeof(char);
134 int mask_size =
sizeof(uint32_t) * 128;
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};
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);
156 &kernel_len, &status);
157 if (status != CL_SUCCESS || !program) {
163 if (status != CL_SUCCESS) {
168 kernel = clCreateKernel(program,
"unsharp_bench", &status);
169 if (status != CL_SUCCESS) {
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);
188 global_work_size_2d, local_work_size_2d, 0,
NULL,
NULL);
193 clReleaseKernel(kernel);
195 clReleaseProgram(program);
197 clReleaseMemObject(cl_inbuf);
199 clReleaseMemObject(cl_outbuf);
201 clReleaseMemObject(cl_mask);
216 int i, j, nb_devices = 0,
count = 0, ret = 0;
221 cl_platform_id platform;
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);
#define AV_OPENCL_KERNEL(...)
ptrdiff_t const GLvoid * data
#define OCLCHECK(method,...)
#define CREATEBUF(out, flags, size)
static void fill_rand_int(int *data, int n)
static av_cold int end(AVCodecContext *avctx)
int av_opencl_get_device_list(AVOpenCLDeviceList **device_list)
Get OpenCL device list.
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
static const uint16_t mask[17]
static int compare_ocl_device_desc(const void *a, const void *b)
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.
#define FFDIFFSIGN(x, y)
Comparator.
int opt_opencl(void *optctx, const char *opt, const char *arg)
static int64_t run_opencl_bench(AVOpenCLExternalEnv *ext_opencl_env)
GLsizei GLboolean const GLfloat * value
int av_opencl_set_option(const char *key, const char *val)
Set option in the global OpenCL context.
int64_t av_gettime(void)
Get the current time in microseconds.
AVOpenCLPlatformNode ** platform_node
cl_command_queue command_queue
int64_t av_gettime_relative(void)
Get the current time in microseconds since some unspecified starting point.
void av_opencl_free_device_list(AVOpenCLDeviceList **device_list)
Free OpenCL device list.
const char * ocl_bench_source
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.
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.
#define av_malloc_array(a, b)
#define AVERROR_EXTERNAL
Generic error in an external library.
int opt_opencl_bench(void *optctx, const char *opt, const char *arg)