Showing the changes to libavfilter/unsharp_opencl_kernel.h (1 of 4 files). Show all
Unified · Split
| @@ -0,0 +1,137 @@ | ||
| 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 */ | |