2 * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
4 * This file is part of FFmpeg.
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.
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.
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
21 #ifndef AVFILTER_UNSHARP_OPENCL_KERNEL_H
22 #define AVFILTER_UNSHARP_OPENCL_KERNEL_H
24 #include "libavutil/opencl.h"
26 const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL(
27 inline unsigned char clip_uint8(int a)
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,
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;
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;
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;
77 temp_src_stride = src_stride_lu;
78 temp_dst_stride = dst_stride_lu;
81 temp_steps_x = step_x_lu;
82 temp_steps_y = step_y_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;
92 temp_src_stride = src_stride_ch;
93 temp_dst_stride = dst_stride_ch;
96 temp_steps_x = step_x_ch;
97 temp_steps_y = step_y_ch;
99 temp_amount = amount_ch;
100 temp_scalebits = scalebits_ch;
101 temp_halfscale = halfscale_ch;
103 y = (global_id - width * height - ch * cw) / cw;
104 x = (global_id - width * height - ch * cw) % cw;
107 temp_src_stride = src_stride_ch;
108 temp_dst_stride = dst_stride_ch;
111 temp_steps_x = step_x_ch;
112 temp_steps_y = step_y_ch;
114 temp_amount = amount_ch;
115 temp_scalebits = scalebits_ch;
116 temp_halfscale = halfscale_ch;
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];
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);
131 temp_dst[x + y * temp_dst_stride] = temp_src[x + y * temp_src_stride];
137 #endif /* AVFILTER_UNSHARP_OPENCL_KERNEL_H */