git.sesse.net Git - ffmpeg/blob - libavfilter/unsharp_opencl_kernel.h

   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 */