inference-engine/thirdparty/fluid/modules/gapi/test/opencl_kernels_test_gapi.hpp

   1 // This file is part of OpenCV project.
   2 // It is subject to the license terms in the LICENSE file found in the top-level directory
   3 // of this distribution and at http://opencv.org/license.html.
   4 //
   5 // Copyright (C) 2018-2019 Intel Corporation
   6
   7 #include "opencv2/core/ocl.hpp"
   8 #include "opencv2/core/ocl_genbase.hpp"
   9 #include "opencv2/core/opencl/ocl_defs.hpp"
  10
  11 #ifdef HAVE_OPENCL
  12 const char* opencl_symm7x7_src =
  13 "#if BORDER_REPLICATE\n"
  14 "#define GET_BORDER(elem) (elem)\n"
  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 "#else\n"
  17 "#define GET_BORDER(elem) (BORDER_CONSTANT_VALUE)\n"
  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 "#endif\n"
  20 "#define GET_A0(id, x, l_edge, a1) ((x) <= (l_edge + 2) ? GET_BORDER(a1) : (((const __global uchar*)(id))[-3]))\n"
  21 "#define GET_A1(id, x, l_edge, a2) ((x) <= (l_edge + 1) ? GET_BORDER(a2) : (((const __global uchar*)(id))[-2]))\n"
  22 "#define GET_A2(id, x, l_edge, b) ((x) <= (l_edge) ? GET_BORDER(b[0]) : (((const __global uchar*)(id))[-1]))\n"
  23 "#define GET_C0(id, x, r_edge, b) ((x) >= (r_edge) ? GET_BORDER(b[8 - 1]) : (((const __global uchar*)(id))[8]))\n"
  24 "#define GET_C1(id, x, r_edge, c0) ((x) >= (r_edge - 1) ? GET_BORDER(c0) : (((const __global uchar*)(id))[8 + 1]))\n"
  25 "#define GET_C2(id, x, r_edge, c1) ((x) >= (r_edge - 2) ? GET_BORDER(c1) : (((const __global uchar*)(id))[8 + 2]))\n"
  26 "__kernel void symm_7x7_test(\n"
  27 "__global const uchar * srcptr,\n"
  28 "int srcStep, int srcEndX, int srcEndY,\n"
  29 "__global uchar * dstptr, int dstStep,\n"
  30 "int rows, int cols,\n"
  31 "int tile_y_coord,\n"
  32 "__constant int * coeff)\n"
  33 "{\n"
  34 "int lEdge = 0, rEdge = cols - 8;\n"
  35 "int x = (get_global_id(0) < cols/8) ? get_global_id(0) * 8: cols - 8;\n"
  36 "int y = get_global_id(1);\n"
  37 "int yd = min(3, tile_y_coord);\n"
  38 "int dst_id = mad24(y, dstStep, x);\n"
  39 "y+=yd;\n"
  40 "int src_id = mad24(y, srcStep, x);\n"
  41 "int y_limit = y + tile_y_coord;\n"
  42 "y_limit-=yd;\n"
  43 "const __global uchar* psrc = (const __global uchar*)(srcptr + src_id);\n"
  44 "__global uchar* pdst = (__global uchar*)(dstptr + dst_id);\n"
  45 "#define BSIZE (7)\n"
  46 "float a0[BSIZE]; float a1[BSIZE]; float a2[BSIZE];\n"
  47 "float8 b[BSIZE];\n"
  48 "float c0[BSIZE]; float c1[BSIZE]; float c2[BSIZE];\n"
  49 "b[3] = convert_float8(vload8(0, (const __global uchar*)psrc));\n"
  50 "if( (y_limit <=2 ) || (y_limit >= srcEndY - 3) || (x >= rEdge-2) || (x <= lEdge + 2) )\n"
  51 "{\n"
  52 "a2[3] = GET_A2(psrc, x, lEdge, b[3]);\n"
  53 "a1[3] = GET_A1(psrc, x, lEdge, a2[3]);\n"
  54 "a0[3] = GET_A0(psrc, x, lEdge, a1[3]);\n"
  55 "c0[3] = GET_C0(psrc, x, rEdge, b[3]);\n"
  56 "c1[3] = GET_C1(psrc, x, rEdge, c0[3]);\n"
  57 "c2[3] = GET_C2(psrc, x, rEdge, c1[3]);\n"
  58 "if(y_limit > 0)\n"
  59 "{\n"
  60 "b[2] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep)));\n"
  61 "a2[2] = GET_A2(psrc - srcStep, x, lEdge, b[2]);\n"
  62 "a1[2] = GET_A1(psrc - srcStep, x, lEdge, a2[2]);\n"
  63 "a0[2] = GET_A0(psrc - srcStep, x, lEdge, a1[2]);\n"
  64 "c0[2] = GET_C0(psrc - srcStep, x, rEdge, b[2]);\n"
  65 "c1[2] = GET_C1(psrc - srcStep, x, rEdge, c0[2]);\n"
  66 "c2[2] = GET_C2(psrc - srcStep, x, rEdge, c1[2]);\n"
  67 "}\n"
  68 "else\n"
  69 "{\n"
  70 "SET_ALL(2, 3);\n"
  71 "}\n"
  72 "if( y_limit > 1 )\n"
  73 "{\n"
  74 "b[1] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*2)));\n"
  75 "a2[1] = GET_A2(psrc - srcStep*2, x, lEdge, b[1]);\n"
  76 "a1[1] = GET_A1(psrc - srcStep*2, x, lEdge, a2[1]);\n"
  77 "a0[1] = GET_A0(psrc - srcStep*2, x, lEdge, a1[1]);\n"
  78 "c0[1] = GET_C0(psrc - srcStep*2, x, rEdge, b[1]);\n"
  79 "c1[1] = GET_C1(psrc - srcStep*2, x, rEdge, c0[1]);\n"
  80 "c2[1] = GET_C2(psrc - srcStep*2, x, rEdge, c1[1]);\n"
  81 "}\n"
  82 "else\n"
  83 "{\n"
  84 "SET_ALL(1, 2);\n"
  85 "}\n"
  86 "if( y_limit > 2 )\n"
  87 "{\n"
  88 "b[0] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*3)));\n"
  89 "a2[0] = GET_A2(psrc - srcStep*3, x, lEdge, b[0]);\n"
  90 "a1[0] = GET_A1(psrc - srcStep*3, x, lEdge, a2[0]);\n"
  91 "a0[0] = GET_A0(psrc - srcStep*3, x, lEdge, a1[0]);\n"
  92 "c0[0] = GET_C0(psrc - srcStep*3, x, rEdge, b[0]);\n"
  93 "c1[0] = GET_C1(psrc - srcStep*3, x, rEdge, c0[0]);\n"
  94 "c2[0] = GET_C2(psrc - srcStep*3, x, rEdge, c1[0]);\n"
  95 "}\n"
  96 "else\n"
  97 "{\n"
  98 "SET_ALL(0, 1);\n"
  99 "}\n"
 100 "if( y_limit < srcEndY - 1 )\n"
 101 "{\n"
 102 "b[4] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep)));\n"
 103 "a2[4] = GET_A2(psrc + srcStep, x, lEdge, b[4]);\n"
 104 "a1[4] = GET_A1(psrc + srcStep, x, lEdge, a2[4]);\n"
 105 "a0[4] = GET_A0(psrc + srcStep, x, lEdge, a1[4]);\n"
 106 "c0[4] = GET_C0(psrc + srcStep, x, rEdge, b[4]);\n"
 107 "c1[4] = GET_C1(psrc + srcStep, x, rEdge, c0[4]);\n"
 108 "c2[4] = GET_C2(psrc + srcStep, x, rEdge, c1[4]);\n"
 109 "}\n"
 110 "else\n"
 111 "{\n"
 112 "SET_ALL(4, 3);\n"
 113 "}\n"
 114 "if( y_limit < srcEndY - 2 )\n"
 115 "{\n"
 116 "b[5] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*2)));\n"
 117 "a2[5] = GET_A2(psrc + srcStep*2, x, lEdge, b[5]);\n"
 118 "a1[5] = GET_A1(psrc + srcStep*2, x, lEdge, a2[5]);\n"
 119 "a0[5] = GET_A0(psrc + srcStep*2, x, lEdge, a1[5]);\n"
 120 "c0[5] = GET_C0(psrc + srcStep*2, x, rEdge, b[5]);\n"
 121 "c1[5] = GET_C1(psrc + srcStep*2, x, rEdge, c0[5]);\n"
 122 "c2[5] = GET_C2(psrc + srcStep*2, x, rEdge, c1[5]);\n"
 123 "}\n"
 124 "else\n"
 125 "{\n"
 126 "SET_ALL(5, 4);\n"
 127 "}\n"
 128 "if( y_limit < srcEndY - 3 )\n"
 129 "{\n"
 130 "b[6] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*3)));\n"
 131 "a2[6] = GET_A2(psrc + srcStep*3, x, lEdge, b[6]);\n"
 132 "a1[6] = GET_A1(psrc + srcStep*3, x, lEdge, a2[6]);\n"
 133 "a0[6] = GET_A0(psrc + srcStep*3, x, lEdge, a1[6]);\n"
 134 "c0[6] = GET_C0(psrc + srcStep*3, x, rEdge, b[6]);\n"
 135 "c1[6] = GET_C1(psrc + srcStep*3, x, rEdge, c0[6]);\n"
 136 "c2[6] = GET_C2(psrc + srcStep*3, x, rEdge, c1[6]);\n"
 137 "}\n"
 138 "else\n"
 139 "{\n"
 140 "SET_ALL(6, 5);\n"
 141 "}\n"
 142 "}\n"
 143 "else\n"
 144 "{\n"
 145 "a2[3] = (((const __global uchar*)(psrc))[-1]);\n"
 146 "a1[3] = (((const __global uchar*)(psrc))[-2]);\n"
 147 "a0[3] = (((const __global uchar*)(psrc))[-3]);\n"
 148 "c0[3] = (((const __global uchar*)(psrc))[8]);\n"
 149 "c1[3] = (((const __global uchar*)(psrc))[8 + 1]);\n"
 150 "c2[3] = (((const __global uchar*)(psrc))[8 + 2]);\n"
 151 "b[2] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep)));\n"
 152 "a2[2] = (((const __global uchar*)(psrc - srcStep))[-1]);\n"
 153 "a1[2] = (((const __global uchar*)(psrc - srcStep))[-2]);\n"
 154 "a0[2] = (((const __global uchar*)(psrc - srcStep))[-3]);\n"
 155 "c0[2] = (((const __global uchar*)(psrc - srcStep))[8]);\n"
 156 "c1[2] = (((const __global uchar*)(psrc - srcStep))[8 + 1]);\n"
 157 "c2[2] = (((const __global uchar*)(psrc - srcStep))[8 + 2]);\n"
 158 "b[1] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*2)));\n"
 159 "a2[1] = (((const __global uchar*)(psrc - srcStep*2))[-1]);\n"
 160 "a1[1] = (((const __global uchar*)(psrc - srcStep*2))[-2]);\n"
 161 "a0[1] = (((const __global uchar*)(psrc - srcStep*2))[-3]);\n"
 162 "c0[1] = (((const __global uchar*)(psrc - srcStep*2))[8]);\n"
 163 "c1[1] = (((const __global uchar*)(psrc - srcStep*2))[8 + 1]);\n"
 164 "c2[1] = (((const __global uchar*)(psrc - srcStep*2))[8 + 2]);\n"
 165 "b[0] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*3)));\n"
 166 "a2[0] = (((const __global uchar*)(psrc - srcStep*3))[-1]);\n"
 167 "a1[0] = (((const __global uchar*)(psrc - srcStep*3))[-2]);\n"
 168 "a0[0] = (((const __global uchar*)(psrc - srcStep*3))[-3]);\n"
 169 "c0[0] = (((const __global uchar*)(psrc - srcStep*3))[8]);\n"
 170 "c1[0] = (((const __global uchar*)(psrc - srcStep*3))[8 + 1]);\n"
 171 "c2[0] = (((const __global uchar*)(psrc - srcStep*3))[8 + 2]);\n"
 172 "b[4] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep)));\n"
 173 "a2[4] = (((const __global uchar*)(psrc + srcStep))[-1]);\n"
 174 "a1[4] = (((const __global uchar*)(psrc + srcStep))[-2]);\n"
 175 "a0[4] = (((const __global uchar*)(psrc + srcStep))[-3]);\n"
 176 "c0[4] = (((const __global uchar*)(psrc + srcStep))[8]);\n"
 177 "c1[4] = (((const __global uchar*)(psrc + srcStep))[8 + 1]);\n"
 178 "c2[4] = (((const __global uchar*)(psrc + srcStep))[8 + 2]);\n"
 179 "b[5] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*2)));\n"
 180 "a2[5] = (((const __global uchar*)(psrc + srcStep*2))[-1]);\n"
 181 "a1[5] = (((const __global uchar*)(psrc + srcStep*2))[-2]);\n"
 182 "a0[5] = (((const __global uchar*)(psrc + srcStep*2))[-3]);\n"
 183 "c0[5] = (((const __global uchar*)(psrc + srcStep*2))[8]);\n"
 184 "c1[5] = (((const __global uchar*)(psrc + srcStep*2))[8 + 1]);\n"
 185 "c2[5] = (((const __global uchar*)(psrc + srcStep*2))[8 + 2]);\n"
 186 "b[6] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*3)));\n"
 187 "a2[6] = (((const __global uchar*)(psrc + srcStep*3))[-1]);\n"
 188 "a1[6] = (((const __global uchar*)(psrc + srcStep*3))[-2]);\n"
 189 "a0[6] = (((const __global uchar*)(psrc + srcStep*3))[-3]);\n"
 190 "c0[6] = (((const __global uchar*)(psrc + srcStep*3))[8]);\n"
 191 "c1[6] = (((const __global uchar*)(psrc + srcStep*3))[8 + 1]);\n"
 192 "c2[6] = (((const __global uchar*)(psrc + srcStep*3))[8 + 2]);\n"
 193 "}\n"
 194 "float a0_sum[3]; float a1_sum[3]; float a2_sum[3];\n"
 195 "float8 b_sum[3];\n"
 196 "float c0_sum[3]; float c1_sum[3]; float c2_sum[3];\n"
 197 "a0_sum[0] = a0[0] + a0[6];\n"
 198 "a0_sum[1] = a0[1] + a0[5];\n"
 199 "a0_sum[2] = a0[2] + a0[4];\n"
 200 "a1_sum[0] = a1[0] + a1[6];\n"
 201 "a1_sum[1] = a1[1] + a1[5];\n"
 202 "a1_sum[2] = a1[2] + a1[4];\n"
 203 "a2_sum[0] = a2[0] + a2[6];\n"
 204 "a2_sum[1] = a2[1] + a2[5];\n"
 205 "a2_sum[2] = a2[2] + a2[4];\n"
 206 "c0_sum[0] = c0[0] + c0[6];\n"
 207 "c0_sum[1] = c0[1] + c0[5];\n"
 208 "c0_sum[2] = c0[2] + c0[4];\n"
 209 "c1_sum[0] = c1[0] + c1[6];\n"
 210 "c1_sum[1] = c1[1] + c1[5];\n"
 211 "c1_sum[2] = c1[2] + c1[4];\n"
 212 "c2_sum[0] = c2[0] + c2[6];\n"
 213 "c2_sum[1] = c2[1] + c2[5];\n"
 214 "c2_sum[2] = c2[2] + c2[4];\n"
 215 "b_sum[0] = b[0] + b[6];\n"
 216 "b_sum[1] = b[1] + b[5];\n"
 217 "b_sum[2] = b[2] + b[4];\n"
 218 "float8 A = b[3];\n"
 219 "float8 intermediate = A * (float)coeff[0];\n"
 220 "float8 B = b_sum[2] +\n"
 221 "(float8)(a2[3], b[3].s0123, b[3].s456) +\n"
 222 "(float8)(b[3].s123, b[3].s4567, c0[3]);\n"
 223 "intermediate += B * (float)coeff[1];\n"
 224 "float8 C = (float8)(a2_sum[2], b_sum[2].s0123, b_sum[2].s456) +\n"
 225 "(float8)(b_sum[2].s123, b_sum[2].s4567, c0_sum[2]);\n"
 226 "intermediate += C * (float)coeff[2];\n"
 227 "float8 D = b_sum[1] +\n"
 228 "(float8)(a1[3], a2[3], b[3].s0123, b[3].s45) +\n"
 229 "(float8)(b[3].s23, b[3].s4567, c0[3], c1[3]);\n"
 230 "intermediate += D * (float)coeff[3];\n"
 231 "float8 E = (float8)(a2_sum[1], b_sum[1].s0123, b_sum[1].s456) +\n"
 232 "(float8)( b_sum[1].s123, b_sum[1].s4567, c0_sum[1]) +\n"
 233 "(float8)( a1_sum[2], a2_sum[2], b_sum[2].s0123, b_sum[2].s45) +\n"
 234 "(float8)( b_sum[2].s23, b_sum[2].s4567, c0_sum[2], c1_sum[2]);\n"
 235 "intermediate += E * (float)coeff[4];\n"
 236 "float8 F = (float8)(a1_sum[1], a2_sum[1], b_sum[1].s0123, b_sum[1].s45) +\n"
 237 "(float8)(b_sum[1].s23, b_sum[1].s4567, c0_sum[1], c1_sum[1]);\n"
 238 "intermediate += F * (float)coeff[5];\n"
 239 "float8 G = b_sum[0] +\n"
 240 "(float8)(a0[3], a1[3], a2[3], b[3].s0123, b[3].s4) +\n"
 241 "(float8)(b[3].s3, b[3].s4567, c0[3], c1[3], c2[3]);\n"
 242 "intermediate += G * (float)coeff[6];\n"
 243 "float8 H = (float8)(a2_sum[0], b_sum[0].s0123, b_sum[0].s456) +\n"
 244 "(float8)(b_sum[0].s123, b_sum[0].s4567, c0_sum[0]) +\n"
 245 "(float8)(a0_sum[2], a1_sum[2], a2_sum[2], b_sum[2].s0123, b_sum[2].s4) +\n"
 246 "(float8)(b_sum[2].s3, b_sum[2].s4567, c0_sum[2], c1_sum[2], c2_sum[2]);\n"
 247 "intermediate += H * (float)coeff[7];\n"
 248 "float8 I = (float8)(a1_sum[0], a2_sum[0], b_sum[0].s0123, b_sum[0].s45) +\n"
 249 "(float8)(b_sum[0].s23, b_sum[0].s4567, c0_sum[0], c1_sum[0]) +\n"
 250 "(float8)(a0_sum[1], a1_sum[1], a2_sum[1], b_sum[1].s0123, b_sum[1].s4) +\n"
 251 "(float8)(b_sum[1].s3, b_sum[1].s4567, c0_sum[1], c1_sum[1], c2_sum[1]);\n"
 252 "intermediate += I * (float)coeff[8];\n"
 253 "float8 J = (float8)(a0_sum[0], a1_sum[0], a2_sum[0], b_sum[0].s0123, b_sum[0].s4) +\n"
 254 "(float8)(b_sum[0].s3, b_sum[0].s4567, c0_sum[0], c1_sum[0], c2_sum[0]);\n"
 255 "intermediate += J * (float)coeff[9];\n"
 256 "intermediate *= SCALE;\n"
 257 "vstore8(convert_uchar8_sat(intermediate), 0, (__global uchar*)(pdst));\n"
 258 "}\n"
 259 ;
 260 #endif