Subversion Repositories Kolibri OS

Rev

Blame | Last modification | View Log | RSS feed

  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;
  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. }
  277.