Home Explore Help
Register Sign In
luoyc
/
OpenCV
1
0
Fork 0
Files Issues 0 Pull Requests 0 Wiki
Tree: a9c35a4807
Branches Tags
OpenCV
/
opencv
/
modules
/
gapi
/
test
/
opencl_kernels_test_gapi.hpp

opencl_kernels_test_gapi.hpp 11 KB
History Raw

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260 
  1. // This file is part of OpenCV project.
    2. 

  2. // It is subject to the license terms in the LICENSE file found in the top-level directory
    3. 

  3. // of this distribution and at http://opencv.org/license.html.
    4. 

  4. //
    5. 

  5. // Copyright (C) 2018 Intel Corporation
    6. 

  6. 

  7. #include <opencv2/core/ocl.hpp>
    8. 

  8. #include <opencv2/core/ocl_genbase.hpp>
    9. 

  9. #include <opencv2/core/opencl/ocl_defs.hpp>
    10. 

  10. 

  11. #ifdef HAVE_OPENCL
    12. 

  12. const char* opencl_symm7x7_src =
    13. 

  13. "#if BORDER_REPLICATE\n"
    14. 

  14. "#define GET_BORDER(elem) (elem)\n"
    15. 

  15. "#define SET_ALL(i, j) a0[i] = a0[j]; a1[i] = a1[j]; a2[i] = a2[j]; b[i] = b[j]; c0[i] = c0[j]; c1[i] = c1[j]; c2[i] = c2[j];\n"
    16. 

  16. "#else\n"
    17. 

  17. "#define GET_BORDER(elem) (BORDER_CONSTANT_VALUE)\n"
    18. 

  18. "#define SET_ALL(i, j) a0[i] = a1[i] = a2[i] = c0[i] = c1[i] = c2[i] = BORDER_CONSTANT_VALUE; b[i] = BORDER_CONSTANT_VALUE;\n"
    19. 

  19. "#endif\n"
    20. 

  20. "#define GET_A0(id, x, l_edge, a1) ((x) <= (l_edge + 2) ? GET_BORDER(a1) : (((const __global uchar*)(id))[-3]))\n"
    21. 

  21. "#define GET_A1(id, x, l_edge, a2) ((x) <= (l_edge + 1) ? GET_BORDER(a2) : (((const __global uchar*)(id))[-2]))\n"
    22. 

  22. "#define GET_A2(id, x, l_edge, b) ((x) <= (l_edge) ? GET_BORDER(b.s0) : (((const __global uchar*)(id))[-1]))\n"
    23. 

  23. "#define GET_C0(id, x, r_edge, b) ((x) >= (r_edge) ? GET_BORDER(b.s7) : (((const __global uchar*)(id))[8]))\n"
    24. 

  24. "#define GET_C1(id, x, r_edge, c0) ((x) >= (r_edge - 1) ? GET_BORDER(c0) : (((const __global uchar*)(id))[8 + 1]))\n"
    25. 

  25. "#define GET_C2(id, x, r_edge, c1) ((x) >= (r_edge - 2) ? GET_BORDER(c1) : (((const __global uchar*)(id))[8 + 2]))\n"
    26. 

  26. "__kernel void symm_7x7_test(\n"
    27. 

  27. "__global const uchar * srcptr,\n"
    28. 

  28. "int srcStep, int srcEndX, int srcEndY,\n"
    29. 

  29. "__global uchar * dstptr, int dstStep,\n"
    30. 

  30. "int rows, int cols,\n"
    31. 

  31. "int tile_y_coord,\n"
    32. 

  32. "__constant int * coeff)\n"
    33. 

  33. "{\n"
    34. 

  34. "int lEdge = 0, rEdge = cols - 8;\n"
    35. 

  35. "int x = (get_global_id(0) < cols/8) ? get_global_id(0) * 8: cols - 8;\n"
    36. 

  36. "int y = get_global_id(1);\n"
    37. 

  37. "int yd = min(3, tile_y_coord);\n"
    38. 

  38. "int dst_id = mad24(y, dstStep, x);\n"
    39. 

  39. "y+=yd;\n"
    40. 

  40. "int src_id = mad24(y, srcStep, x);\n"
    41. 

  41. "int y_limit = y + tile_y_coord;\n"
    42. 

  42. "y_limit-=yd;\n"
    43. 

  43. "const __global uchar* psrc = (const __global uchar*)(srcptr + src_id);\n"
    44. 

  44. "__global uchar* pdst = (__global uchar*)(dstptr + dst_id);\n"
    45. 

  45. "#define BSIZE (7)\n"
    46. 

  46. "float a0[BSIZE]; float a1[BSIZE]; float a2[BSIZE];\n"
    47. 

  47. "float8 b[BSIZE];\n"
    48. 

  48. "float c0[BSIZE]; float c1[BSIZE]; float c2[BSIZE];\n"
    49. 

  49. "b[3] = convert_float8(vload8(0, (const __global uchar*)psrc));\n"
    50. 

  50. "if( (y_limit <=2 ) || (y_limit >= srcEndY - 3) || (x >= rEdge-2) || (x <= lEdge + 2) )\n"
    51. 

  51. "{\n"
    52. 

  52. "a2[3] = GET_A2(psrc, x, lEdge, b[3]);\n"
    53. 

  53. "a1[3] = GET_A1(psrc, x, lEdge, a2[3]);\n"
    54. 

  54. "a0[3] = GET_A0(psrc, x, lEdge, a1[3]);\n"
    55. 

  55. "c0[3] = GET_C0(psrc, x, rEdge, b[3]);\n"
    56. 

  56. "c1[3] = GET_C1(psrc, x, rEdge, c0[3]);\n"
    57. 

  57. "c2[3] = GET_C2(psrc, x, rEdge, c1[3]);\n"
    58. 

  58. "if(y_limit > 0)\n"
    59. 

  59. "{\n"
    60. 

  60. "b[2] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep)));\n"
    61. 

  61. "a2[2] = GET_A2(psrc - srcStep, x, lEdge, b[2]);\n"
    62. 

  62. "a1[2] = GET_A1(psrc - srcStep, x, lEdge, a2[2]);\n"
    63. 

  63. "a0[2] = GET_A0(psrc - srcStep, x, lEdge, a1[2]);\n"
    64. 

  64. "c0[2] = GET_C0(psrc - srcStep, x, rEdge, b[2]);\n"
    65. 

  65. "c1[2] = GET_C1(psrc - srcStep, x, rEdge, c0[2]);\n"
    66. 

  66. "c2[2] = GET_C2(psrc - srcStep, x, rEdge, c1[2]);\n"
    67. 

  67. "}\n"
    68. 

  68. "else\n"
    69. 

  69. "{\n"
    70. 

  70. "SET_ALL(2, 3);\n"
    71. 

  71. "}\n"
    72. 

  72. "if( y_limit > 1 )\n"
    73. 

  73. "{\n"
    74. 

  74. "b[1] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*2)));\n"
    75. 

  75. "a2[1] = GET_A2(psrc - srcStep*2, x, lEdge, b[1]);\n"
    76. 

  76. "a1[1] = GET_A1(psrc - srcStep*2, x, lEdge, a2[1]);\n"
    77. 

  77. "a0[1] = GET_A0(psrc - srcStep*2, x, lEdge, a1[1]);\n"
    78. 

  78. "c0[1] = GET_C0(psrc - srcStep*2, x, rEdge, b[1]);\n"
    79. 

  79. "c1[1] = GET_C1(psrc - srcStep*2, x, rEdge, c0[1]);\n"
    80. 

  80. "c2[1] = GET_C2(psrc - srcStep*2, x, rEdge, c1[1]);\n"
    81. 

  81. "}\n"
    82. 

  82. "else\n"
    83. 

  83. "{\n"
    84. 

  84. "SET_ALL(1, 2);\n"
    85. 

  85. "}\n"
    86. 

  86. "if( y_limit > 2 )\n"
    87. 

  87. "{\n"
    88. 

  88. "b[0] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*3)));\n"
    89. 

  89. "a2[0] = GET_A2(psrc - srcStep*3, x, lEdge, b[0]);\n"
    90. 

  90. "a1[0] = GET_A1(psrc - srcStep*3, x, lEdge, a2[0]);\n"
    91. 

  91. "a0[0] = GET_A0(psrc - srcStep*3, x, lEdge, a1[0]);\n"
    92. 

  92. "c0[0] = GET_C0(psrc - srcStep*3, x, rEdge, b[0]);\n"
    93. 

  93. "c1[0] = GET_C1(psrc - srcStep*3, x, rEdge, c0[0]);\n"
    94. 

  94. "c2[0] = GET_C2(psrc - srcStep*3, x, rEdge, c1[0]);\n"
    95. 

  95. "}\n"
    96. 

  96. "else\n"
    97. 

  97. "{\n"
    98. 

  98. "SET_ALL(0, 1);\n"
    99. 

  99. "}\n"
    100. 

  100. "if( y_limit < srcEndY - 1 )\n"
    101. 

  101. "{\n"
    102. 

  102. "b[4] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep)));\n"
    103. 

  103. "a2[4] = GET_A2(psrc + srcStep, x, lEdge, b[4]);\n"
    104. 

  104. "a1[4] = GET_A1(psrc + srcStep, x, lEdge, a2[4]);\n"
    105. 

  105. "a0[4] = GET_A0(psrc + srcStep, x, lEdge, a1[4]);\n"
    106. 

  106. "c0[4] = GET_C0(psrc + srcStep, x, rEdge, b[4]);\n"
    107. 

  107. "c1[4] = GET_C1(psrc + srcStep, x, rEdge, c0[4]);\n"
    108. 

  108. "c2[4] = GET_C2(psrc + srcStep, x, rEdge, c1[4]);\n"
    109. 

  109. "}\n"
    110. 

  110. "else\n"
    111. 

  111. "{\n"
    112. 

  112. "SET_ALL(4, 3);\n"
    113. 

  113. "}\n"
    114. 

  114. "if( y_limit < srcEndY - 2 )\n"
    115. 

  115. "{\n"
    116. 

  116. "b[5] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*2)));\n"
    117. 

  117. "a2[5] = GET_A2(psrc + srcStep*2, x, lEdge, b[5]);\n"
    118. 

  118. "a1[5] = GET_A1(psrc + srcStep*2, x, lEdge, a2[5]);\n"
    119. 

  119. "a0[5] = GET_A0(psrc + srcStep*2, x, lEdge, a1[5]);\n"
    120. 

  120. "c0[5] = GET_C0(psrc + srcStep*2, x, rEdge, b[5]);\n"
    121. 

  121. "c1[5] = GET_C1(psrc + srcStep*2, x, rEdge, c0[5]);\n"
    122. 

  122. "c2[5] = GET_C2(psrc + srcStep*2, x, rEdge, c1[5]);\n"
    123. 

  123. "}\n"
    124. 

  124. "else\n"
    125. 

  125. "{\n"
    126. 

  126. "SET_ALL(5, 4);\n"
    127. 

  127. "}\n"
    128. 

  128. "if( y_limit < srcEndY - 3 )\n"
    129. 

  129. "{\n"
    130. 

  130. "b[6] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*3)));\n"
    131. 

  131. "a2[6] = GET_A2(psrc + srcStep*3, x, lEdge, b[6]);\n"
    132. 

  132. "a1[6] = GET_A1(psrc + srcStep*3, x, lEdge, a2[6]);\n"
    133. 

  133. "a0[6] = GET_A0(psrc + srcStep*3, x, lEdge, a1[6]);\n"
    134. 

  134. "c0[6] = GET_C0(psrc + srcStep*3, x, rEdge, b[6]);\n"
    135. 

  135. "c1[6] = GET_C1(psrc + srcStep*3, x, rEdge, c0[6]);\n"
    136. 

  136. "c2[6] = GET_C2(psrc + srcStep*3, x, rEdge, c1[6]);\n"
    137. 

  137. "}\n"
    138. 

  138. "else\n"
    139. 

  139. "{\n"
    140. 

  140. "SET_ALL(6, 5);\n"
    141. 

  141. "}\n"
    142. 

  142. "}\n"
    143. 

  143. "else\n"
    144. 

  144. "{\n"
    145. 

  145. "a2[3] = (((const __global uchar*)(psrc))[-1]);\n"
    146. 

  146. "a1[3] = (((const __global uchar*)(psrc))[-2]);\n"
    147. 

  147. "a0[3] = (((const __global uchar*)(psrc))[-3]);\n"
    148. 

  148. "c0[3] = (((const __global uchar*)(psrc))[8]);\n"
    149. 

  149. "c1[3] = (((const __global uchar*)(psrc))[8 + 1]);\n"
    150. 

  150. "c2[3] = (((const __global uchar*)(psrc))[8 + 2]);\n"
    151. 

  151. "b[2] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep)));\n"
    152. 

  152. "a2[2] = (((const __global uchar*)(psrc - srcStep))[-1]);\n"
    153. 

  153. "a1[2] = (((const __global uchar*)(psrc - srcStep))[-2]);\n"
    154. 

  154. "a0[2] = (((const __global uchar*)(psrc - srcStep))[-3]);\n"
    155. 

  155. "c0[2] = (((const __global uchar*)(psrc - srcStep))[8]);\n"
    156. 

  156. "c1[2] = (((const __global uchar*)(psrc - srcStep))[8 + 1]);\n"
    157. 

  157. "c2[2] = (((const __global uchar*)(psrc - srcStep))[8 + 2]);\n"
    158. 

  158. "b[1] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*2)));\n"
    159. 

  159. "a2[1] = (((const __global uchar*)(psrc - srcStep*2))[-1]);\n"
    160. 

  160. "a1[1] = (((const __global uchar*)(psrc - srcStep*2))[-2]);\n"
    161. 

  161. "a0[1] = (((const __global uchar*)(psrc - srcStep*2))[-3]);\n"
    162. 

  162. "c0[1] = (((const __global uchar*)(psrc - srcStep*2))[8]);\n"
    163. 

  163. "c1[1] = (((const __global uchar*)(psrc - srcStep*2))[8 + 1]);\n"
    164. 

  164. "c2[1] = (((const __global uchar*)(psrc - srcStep*2))[8 + 2]);\n"
    165. 

  165. "b[0] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*3)));\n"
    166. 

  166. "a2[0] = (((const __global uchar*)(psrc - srcStep*3))[-1]);\n"
    167. 

  167. "a1[0] = (((const __global uchar*)(psrc - srcStep*3))[-2]);\n"
    168. 

  168. "a0[0] = (((const __global uchar*)(psrc - srcStep*3))[-3]);\n"
    169. 

  169. "c0[0] = (((const __global uchar*)(psrc - srcStep*3))[8]);\n"
    170. 

  170. "c1[0] = (((const __global uchar*)(psrc - srcStep*3))[8 + 1]);\n"
    171. 

  171. "c2[0] = (((const __global uchar*)(psrc - srcStep*3))[8 + 2]);\n"
    172. 

  172. "b[4] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep)));\n"
    173. 

  173. "a2[4] = (((const __global uchar*)(psrc + srcStep))[-1]);\n"
    174. 

  174. "a1[4] = (((const __global uchar*)(psrc + srcStep))[-2]);\n"
    175. 

  175. "a0[4] = (((const __global uchar*)(psrc + srcStep))[-3]);\n"
    176. 

  176. "c0[4] = (((const __global uchar*)(psrc + srcStep))[8]);\n"
    177. 

  177. "c1[4] = (((const __global uchar*)(psrc + srcStep))[8 + 1]);\n"
    178. 

  178. "c2[4] = (((const __global uchar*)(psrc + srcStep))[8 + 2]);\n"
    179. 

  179. "b[5] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*2)));\n"
    180. 

  180. "a2[5] = (((const __global uchar*)(psrc + srcStep*2))[-1]);\n"
    181. 

  181. "a1[5] = (((const __global uchar*)(psrc + srcStep*2))[-2]);\n"
    182. 

  182. "a0[5] = (((const __global uchar*)(psrc + srcStep*2))[-3]);\n"
    183. 

  183. "c0[5] = (((const __global uchar*)(psrc + srcStep*2))[8]);\n"
    184. 

  184. "c1[5] = (((const __global uchar*)(psrc + srcStep*2))[8 + 1]);\n"
    185. 

  185. "c2[5] = (((const __global uchar*)(psrc + srcStep*2))[8 + 2]);\n"
    186. 

  186. "b[6] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*3)));\n"
    187. 

  187. "a2[6] = (((const __global uchar*)(psrc + srcStep*3))[-1]);\n"
    188. 

  188. "a1[6] = (((const __global uchar*)(psrc + srcStep*3))[-2]);\n"
    189. 

  189. "a0[6] = (((const __global uchar*)(psrc + srcStep*3))[-3]);\n"
    190. 

  190. "c0[6] = (((const __global uchar*)(psrc + srcStep*3))[8]);\n"
    191. 

  191. "c1[6] = (((const __global uchar*)(psrc + srcStep*3))[8 + 1]);\n"
    192. 

  192. "c2[6] = (((const __global uchar*)(psrc + srcStep*3))[8 + 2]);\n"
    193. 

  193. "}\n"
    194. 

  194. "float a0_sum[3]; float a1_sum[3]; float a2_sum[3];\n"
    195. 

  195. "float8 b_sum[3];\n"
    196. 

  196. "float c0_sum[3]; float c1_sum[3]; float c2_sum[3];\n"
    197. 

  197. "a0_sum[0] = a0[0] + a0[6];\n"
    198. 

  198. "a0_sum[1] = a0[1] + a0[5];\n"
    199. 

  199. "a0_sum[2] = a0[2] + a0[4];\n"
    200. 

  200. "a1_sum[0] = a1[0] + a1[6];\n"
    201. 

  201. "a1_sum[1] = a1[1] + a1[5];\n"
    202. 

  202. "a1_sum[2] = a1[2] + a1[4];\n"
    203. 

  203. "a2_sum[0] = a2[0] + a2[6];\n"
    204. 

  204. "a2_sum[1] = a2[1] + a2[5];\n"
    205. 

  205. "a2_sum[2] = a2[2] + a2[4];\n"
    206. 

  206. "c0_sum[0] = c0[0] + c0[6];\n"
    207. 

  207. "c0_sum[1] = c0[1] + c0[5];\n"
    208. 

  208. "c0_sum[2] = c0[2] + c0[4];\n"
    209. 

  209. "c1_sum[0] = c1[0] + c1[6];\n"
    210. 

  210. "c1_sum[1] = c1[1] + c1[5];\n"
    211. 

  211. "c1_sum[2] = c1[2] + c1[4];\n"
    212. 

  212. "c2_sum[0] = c2[0] + c2[6];\n"
    213. 

  213. "c2_sum[1] = c2[1] + c2[5];\n"
    214. 

  214. "c2_sum[2] = c2[2] + c2[4];\n"
    215. 

  215. "b_sum[0] = b[0] + b[6];\n"
    216. 

  216. "b_sum[1] = b[1] + b[5];\n"
    217. 

  217. "b_sum[2] = b[2] + b[4];\n"
    218. 

  218. "float8 A = b[3];\n"
    219. 

  219. "float8 intermediate = A * (float)coeff[0];\n"
    220. 

  220. "float8 B = b_sum[2] +\n"
    221. 

  221. "(float8)(a2[3], b[3].s0123, b[3].s456) +\n"
    222. 

  222. "(float8)(b[3].s123, b[3].s4567, c0[3]);\n"
    223. 

  223. "intermediate += B * (float)coeff[1];\n"
    224. 

  224. "float8 C = (float8)(a2_sum[2], b_sum[2].s0123, b_sum[2].s456) +\n"
    225. 

  225. "(float8)(b_sum[2].s123, b_sum[2].s4567, c0_sum[2]);\n"
    226. 

  226. "intermediate += C * (float)coeff[2];\n"
    227. 

  227. "float8 D = b_sum[1] +\n"
    228. 

  228. "(float8)(a1[3], a2[3], b[3].s0123, b[3].s45) +\n"
    229. 

  229. "(float8)(b[3].s23, b[3].s4567, c0[3], c1[3]);\n"
    230. 

  230. "intermediate += D * (float)coeff[3];\n"
    231. 

  231. "float8 E = (float8)(a2_sum[1], b_sum[1].s0123, b_sum[1].s456) +\n"
    232. 

  232. "(float8)( b_sum[1].s123, b_sum[1].s4567, c0_sum[1]) +\n"
    233. 

  233. "(float8)( a1_sum[2], a2_sum[2], b_sum[2].s0123, b_sum[2].s45) +\n"
    234. 

  234. "(float8)( b_sum[2].s23, b_sum[2].s4567, c0_sum[2], c1_sum[2]);\n"
    235. 

  235. "intermediate += E * (float)coeff[4];\n"
    236. 

  236. "float8 F = (float8)(a1_sum[1], a2_sum[1], b_sum[1].s0123, b_sum[1].s45) +\n"
    237. 

  237. "(float8)(b_sum[1].s23, b_sum[1].s4567, c0_sum[1], c1_sum[1]);\n"
    238. 

  238. "intermediate += F * (float)coeff[5];\n"
    239. 

  239. "float8 G = b_sum[0] +\n"
    240. 

  240. "(float8)(a0[3], a1[3], a2[3], b[3].s0123, b[3].s4) +\n"
    241. 

  241. "(float8)(b[3].s3, b[3].s4567, c0[3], c1[3], c2[3]);\n"
    242. 

  242. "intermediate += G * (float)coeff[6];\n"
    243. 

  243. "float8 H = (float8)(a2_sum[0], b_sum[0].s0123, b_sum[0].s456) +\n"
    244. 

  244. "(float8)(b_sum[0].s123, b_sum[0].s4567, c0_sum[0]) +\n"
    245. 

  245. "(float8)(a0_sum[2], a1_sum[2], a2_sum[2], b_sum[2].s0123, b_sum[2].s4) +\n"
    246. 

  246. "(float8)(b_sum[2].s3, b_sum[2].s4567, c0_sum[2], c1_sum[2], c2_sum[2]);\n"
    247. 

  247. "intermediate += H * (float)coeff[7];\n"
    248. 

  248. "float8 I = (float8)(a1_sum[0], a2_sum[0], b_sum[0].s0123, b_sum[0].s45) +\n"
    249. 

  249. "(float8)(b_sum[0].s23, b_sum[0].s4567, c0_sum[0], c1_sum[0]) +\n"
    250. 

  250. "(float8)(a0_sum[1], a1_sum[1], a2_sum[1], b_sum[1].s0123, b_sum[1].s4) +\n"
    251. 

  251. "(float8)(b_sum[1].s3, b_sum[1].s4567, c0_sum[1], c1_sum[1], c2_sum[1]);\n"
    252. 

  252. "intermediate += I * (float)coeff[8];\n"
    253. 

  253. "float8 J = (float8)(a0_sum[0], a1_sum[0], a2_sum[0], b_sum[0].s0123, b_sum[0].s4) +\n"
    254. 

  254. "(float8)(b_sum[0].s3, b_sum[0].s4567, c0_sum[0], c1_sum[0], c2_sum[0]);\n"
    255. 

  255. "intermediate += J * (float)coeff[9];\n"
    256. 

  256. "intermediate *= SCALE;\n"
    257. 

  257. "vstore8(convert_uchar8_sat(intermediate), 0, (__global uchar*)(pdst));\n"
    258. 

  258. "}\n"
    259. 

  259. ;
    260. 

  260. #endif
    261.