deshake_opencl_kernel.h 8.1 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225
  1. /*
  2. * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
  3. * Copyright (C) 2013 Lenny Wang
  4. *
  5. *
  6. * This file is part of FFmpeg.
  7. *
  8. * FFmpeg is free software; you can redistribute it and/or
  9. * modify it under the terms of the GNU Lesser General Public
  10. * License as published by the Free Software Foundation; either
  11. * version 2.1 of the License, or (at your option) any later version.
  12. *
  13. * FFmpeg is distributed in the hope that it will be useful,
  14. * but WITHOUT ANY WARRANTY; without even the implied warranty of
  15. * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
  16. * Lesser General Public License for more details.
  17. *
  18. * You should have received a copy of the GNU Lesser General Public
  19. * License along with FFmpeg; if not, write to the Free Software
  20. * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  21. */
  22. #ifndef AVFILTER_DESHAKE_OPENCL_KERNEL_H
  23. #define AVFILTER_DESHAKE_OPENCL_KERNEL_H
  24. #include "libavutil/opencl.h"
  25. const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL(
  26. inline unsigned char pixel(global const unsigned char *src, int x, int y,
  27. int w, int h,int stride, unsigned char def)
  28. {
  29. return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[x + y * stride];
  30. }
  31. unsigned char interpolate_nearest(float x, float y, global const unsigned char *src,
  32. int width, int height, int stride, unsigned char def)
  33. {
  34. return pixel(src, (int)(x + 0.5f), (int)(y + 0.5f), width, height, stride, def);
  35. }
  36. unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src,
  37. int width, int height, int stride, unsigned char def)
  38. {
  39. int x_c, x_f, y_c, y_f;
  40. int v1, v2, v3, v4;
  41. x_f = (int)x;
  42. y_f = (int)y;
  43. x_c = x_f + 1;
  44. y_c = y_f + 1;
  45. if (x_f < -1 || x_f > width || y_f < -1 || y_f > height) {
  46. return def;
  47. } else {
  48. v4 = pixel(src, x_f, y_f, width, height, stride, def);
  49. v2 = pixel(src, x_c, y_f, width, height, stride, def);
  50. v3 = pixel(src, x_f, y_c, width, height, stride, def);
  51. v1 = pixel(src, x_c, y_c, width, height, stride, def);
  52. return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) +
  53. v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y)));
  54. }
  55. }
  56. unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src,
  57. int width, int height, int stride, unsigned char def)
  58. {
  59. int x_c, x_f, y_c, y_f;
  60. unsigned char v1, v2, v3, v4;
  61. float f1, f2, f3, f4;
  62. x_f = (int)x;
  63. y_f = (int)y;
  64. x_c = x_f + 1;
  65. y_c = y_f + 1;
  66. if (x_f < - 1 || x_f > width || y_f < -1 || y_f > height)
  67. return def;
  68. else {
  69. v4 = pixel(src, x_f, y_f, width, height, stride, def);
  70. v2 = pixel(src, x_c, y_f, width, height, stride, def);
  71. v3 = pixel(src, x_f, y_c, width, height, stride, def);
  72. v1 = pixel(src, x_c, y_c, width, height, stride, def);
  73. f1 = 1 - sqrt((x_c - x) * (y_c - y));
  74. f2 = 1 - sqrt((x_c - x) * (y - y_f));
  75. f3 = 1 - sqrt((x - x_f) * (y_c - y));
  76. f4 = 1 - sqrt((x - x_f) * (y - y_f));
  77. return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4);
  78. }
  79. }
  80. inline const float clipf(float a, float amin, float amax)
  81. {
  82. if (a < amin) return amin;
  83. else if (a > amax) return amax;
  84. else return a;
  85. }
  86. inline int mirror(int v, int m)
  87. {
  88. while ((unsigned)v > (unsigned)m) {
  89. v = -v;
  90. if (v < 0)
  91. v += 2 * m;
  92. }
  93. return v;
  94. }
  95. kernel void avfilter_transform_luma(global unsigned char *src,
  96. global unsigned char *dst,
  97. float4 matrix,
  98. int interpolate,
  99. int fill,
  100. int src_stride_lu,
  101. int dst_stride_lu,
  102. int height,
  103. int width)
  104. {
  105. int x = get_global_id(0);
  106. int y = get_global_id(1);
  107. int idx_dst = y * dst_stride_lu + x;
  108. unsigned char def = 0;
  109. float x_s = x * matrix.x + y * matrix.y + matrix.z;
  110. float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
  111. if (x < width && y < height) {
  112. switch (fill) {
  113. case 0: //FILL_BLANK
  114. def = 0;
  115. break;
  116. case 1: //FILL_ORIGINAL
  117. def = src[y*src_stride_lu + x];
  118. break;
  119. case 2: //FILL_CLAMP
  120. y_s = clipf(y_s, 0, height - 1);
  121. x_s = clipf(x_s, 0, width - 1);
  122. def = src[(int)y_s * src_stride_lu + (int)x_s];
  123. break;
  124. case 3: //FILL_MIRROR
  125. y_s = mirror(y_s, height - 1);
  126. x_s = mirror(x_s, width - 1);
  127. def = src[(int)y_s * src_stride_lu + (int)x_s];
  128. break;
  129. }
  130. switch (interpolate) {
  131. case 0: //INTERPOLATE_NEAREST
  132. dst[idx_dst] = interpolate_nearest(x_s, y_s, src, width, height, src_stride_lu, def);
  133. break;
  134. case 1: //INTERPOLATE_BILINEAR
  135. dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, width, height, src_stride_lu, def);
  136. break;
  137. case 2: //INTERPOLATE_BIQUADRATIC
  138. dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, width, height, src_stride_lu, def);
  139. break;
  140. default:
  141. return;
  142. }
  143. }
  144. }
  145. kernel void avfilter_transform_chroma(global unsigned char *src,
  146. global unsigned char *dst,
  147. float4 matrix,
  148. int interpolate,
  149. int fill,
  150. int src_stride_lu,
  151. int dst_stride_lu,
  152. int src_stride_ch,
  153. int dst_stride_ch,
  154. int height,
  155. int width,
  156. int ch,
  157. int cw)
  158. {
  159. int x = get_global_id(0);
  160. int y = get_global_id(1);
  161. int pad_ch = get_global_size(1)>>1;
  162. global unsigned char *dst_u = dst + height * dst_stride_lu;
  163. global unsigned char *src_u = src + height * src_stride_lu;
  164. global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
  165. global unsigned char *src_v = src_u + ch * src_stride_ch;
  166. src = y < pad_ch ? src_u : src_v;
  167. dst = y < pad_ch ? dst_u : dst_v;
  168. y = select(y - pad_ch, y, y < pad_ch);
  169. float x_s = x * matrix.x + y * matrix.y + matrix.z;
  170. float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
  171. int idx_dst = y * dst_stride_ch + x;
  172. unsigned char def;
  173. if (x < cw && y < ch) {
  174. switch (fill) {
  175. case 0: //FILL_BLANK
  176. def = 0;
  177. break;
  178. case 1: //FILL_ORIGINAL
  179. def = src[y*src_stride_ch + x];
  180. break;
  181. case 2: //FILL_CLAMP
  182. y_s = clipf(y_s, 0, ch - 1);
  183. x_s = clipf(x_s, 0, cw - 1);
  184. def = src[(int)y_s * src_stride_ch + (int)x_s];
  185. break;
  186. case 3: //FILL_MIRROR
  187. y_s = mirror(y_s, ch - 1);
  188. x_s = mirror(x_s, cw - 1);
  189. def = src[(int)y_s * src_stride_ch + (int)x_s];
  190. break;
  191. }
  192. switch (interpolate) {
  193. case 0: //INTERPOLATE_NEAREST
  194. dst[idx_dst] = interpolate_nearest(x_s, y_s, src, cw, ch, src_stride_ch, def);
  195. break;
  196. case 1: //INTERPOLATE_BILINEAR
  197. dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, cw, ch, src_stride_ch, def);
  198. break;
  199. case 2: //INTERPOLATE_BIQUADRATIC
  200. dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, cw, ch, src_stride_ch, def);
  201. break;
  202. default:
  203. return;
  204. }
  205. }
  206. }
  207. );
  208. #endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */