Subversion Repositories Kolibri OS

Rev

Go to most recent revision | Blame | Last modification | View Log | RSS feed

  1. /*
  2.  * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
  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. #ifndef AVFILTER_UNSHARP_OPENCL_KERNEL_H
  22. #define AVFILTER_UNSHARP_OPENCL_KERNEL_H
  23.  
  24. #include "libavutil/opencl.h"
  25.  
  26. const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL(
  27. inline unsigned char clip_uint8(int a)
  28. {
  29.     if (a & (~0xFF))
  30.         return (-a)>>31;
  31.     else
  32.         return a;
  33. }
  34.  
  35. kernel void unsharp(global  unsigned char *src,
  36.                     global  unsigned char *dst,
  37.                     const global  unsigned int *mask_lu,
  38.                     const global  unsigned int *mask_ch,
  39.                     int amount_lu,
  40.                     int amount_ch,
  41.                     int step_x_lu,
  42.                     int step_y_lu,
  43.                     int step_x_ch,
  44.                     int step_y_ch,
  45.                     int scalebits_lu,
  46.                     int scalebits_ch,
  47.                     int halfscale_lu,
  48.                     int halfscale_ch,
  49.                     int src_stride_lu,
  50.                     int src_stride_ch,
  51.                     int dst_stride_lu,
  52.                     int dst_stride_ch,
  53.                     int height,
  54.                     int width,
  55.                     int ch,
  56.                     int cw)
  57. {
  58.     global unsigned char *dst_y = dst;
  59.     global unsigned char *dst_u = dst_y + height * dst_stride_lu;
  60.     global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
  61.  
  62.     global unsigned char *src_y = src;
  63.     global unsigned char *src_u = src_y + height * src_stride_lu;
  64.     global unsigned char *src_v = src_u + ch * src_stride_ch;
  65.  
  66.     global unsigned char *temp_dst;
  67.     global unsigned char *temp_src;
  68.     const global unsigned int *temp_mask;
  69.     int global_id = get_global_id(0);
  70.     int i, j, x, y, temp_src_stride, temp_dst_stride, temp_height, temp_width, temp_steps_x, temp_steps_y,
  71.         temp_amount, temp_scalebits, temp_halfscale, sum, idx_x, idx_y, temp, res;
  72.     if (global_id < width * height) {
  73.         y = global_id / width;
  74.         x = global_id % width;
  75.         temp_dst = dst_y;
  76.         temp_src = src_y;
  77.         temp_src_stride = src_stride_lu;
  78.         temp_dst_stride = dst_stride_lu;
  79.         temp_height = height;
  80.         temp_width = width;
  81.         temp_steps_x = step_x_lu;
  82.         temp_steps_y = step_y_lu;
  83.         temp_mask = mask_lu;
  84.         temp_amount = amount_lu;
  85.         temp_scalebits = scalebits_lu;
  86.         temp_halfscale = halfscale_lu;
  87.     } else if ((global_id >= width * height) && (global_id < width * height + ch * cw)) {
  88.         y = (global_id - width * height) / cw;
  89.         x = (global_id - width * height) % cw;
  90.         temp_dst = dst_u;
  91.         temp_src = src_u;
  92.         temp_src_stride = src_stride_ch;
  93.         temp_dst_stride = dst_stride_ch;
  94.         temp_height = ch;
  95.         temp_width = cw;
  96.         temp_steps_x = step_x_ch;
  97.         temp_steps_y = step_y_ch;
  98.         temp_mask = mask_ch;
  99.         temp_amount = amount_ch;
  100.         temp_scalebits = scalebits_ch;
  101.         temp_halfscale = halfscale_ch;
  102.     } else {
  103.         y = (global_id - width * height - ch * cw) / cw;
  104.         x = (global_id - width * height - ch * cw) % cw;
  105.         temp_dst = dst_v;
  106.         temp_src = src_v;
  107.         temp_src_stride = src_stride_ch;
  108.         temp_dst_stride = dst_stride_ch;
  109.         temp_height = ch;
  110.         temp_width = cw;
  111.         temp_steps_x = step_x_ch;
  112.         temp_steps_y = step_y_ch;
  113.         temp_mask = mask_ch;
  114.         temp_amount = amount_ch;
  115.         temp_scalebits = scalebits_ch;
  116.         temp_halfscale = halfscale_ch;
  117.     }
  118.     if (temp_amount) {
  119.         sum = 0;
  120.         for (j = 0; j <= 2 * temp_steps_y; j++) {
  121.             idx_y = (y - temp_steps_y + j) <= 0 ? 0 : (y - temp_steps_y + j) >= temp_height ? temp_height-1 : y - temp_steps_y + j;
  122.             for (i = 0; i <= 2 * temp_steps_x; i++) {
  123.                 idx_x = (x - temp_steps_x + i) <= 0 ? 0 : (x - temp_steps_x + i) >= temp_width ? temp_width-1 : x - temp_steps_x + i;
  124.                 sum += temp_mask[i + j * (2 * temp_steps_x + 1)] * temp_src[idx_x + idx_y * temp_src_stride];
  125.             }
  126.         }
  127.         temp = (int)temp_src[x + y * temp_src_stride];
  128.         res = temp + (((temp - (int)((sum + temp_halfscale) >> temp_scalebits)) * temp_amount) >> 16);
  129.         temp_dst[x + y * temp_dst_stride] = clip_uint8(res);
  130.     } else {
  131.         temp_dst[x + y * temp_dst_stride] = temp_src[x + y * temp_src_stride];
  132.     }
  133. }
  134.  
  135. );
  136.  
  137. #endif /* AVFILTER_UNSHARP_OPENCL_KERNEL_H */
  138.