FFmpeg
 All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups Pages
unsharp_opencl_kernel.h
Go to the documentation of this file.
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 */
Generated on Wed Jul 10 2013 23:48:10 for FFmpeg by   doxygen 1.8.2