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) 2014, Itseez, Inc., all rights reserved. 6 // Third party copyrights are property of their respective owners. 7 8 9 #define noconvert 10 11 #ifdef ONLY_SUM_CONVERT 12 13 __kernel void sumConvert(__global const uchar * src1ptr, int src1_step, int src1_offset, 14 __global const uchar * src2ptr, int src2_step, int src2_offset, 15 __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, 16 coeffT scale, coeffT delta) 17 { 18 int x = get_global_id(0); 19 int y = get_global_id(1); 20 21 if (y < dst_rows && x < dst_cols) 22 { 23 int src1_index = mad24(y, src1_step, mad24(x, (int)sizeof(srcT), src1_offset)); 24 int src2_index = mad24(y, src2_step, mad24(x, (int)sizeof(srcT), src2_offset)); 25 int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT), dst_offset)); 26 27 __global const srcT * src1 = (__global const srcT *)(src1ptr + src1_index); 28 __global const srcT * src2 = (__global const srcT *)(src2ptr + src2_index); 29 __global dstT * dst = (__global dstT *)(dstptr + dst_index); 30 31 #if wdepth <= 4 32 dst[0] = convertToDT( mad24((WT)(scale), convertToWT(src1[0]) + convertToWT(src2[0]), (WT)(delta)) ); 33 #else 34 dst[0] = convertToDT( mad((WT)(scale), convertToWT(src1[0]) + convertToWT(src2[0]), (WT)(delta)) ); 35 #endif 36 } 37 } 38 39 #else 40 41 /////////////////////////////////////////////////////////////////////////////////////////////////// 42 /////////////////////////////////Macro for border type//////////////////////////////////////////// 43 ///////////////////////////////////////////////////////////////////////////////////////////////// 44 45 #ifdef BORDER_CONSTANT 46 // CCCCCC|abcdefgh|CCCCCCC 47 #define EXTRAPOLATE(x, maxV) 48 #elif defined BORDER_REPLICATE 49 // aaaaaa|abcdefgh|hhhhhhh 50 #define EXTRAPOLATE(x, maxV) \ 51 { \ 52 (x) = clamp((x), 0, (maxV)-1); \ 53 } 54 #elif defined BORDER_WRAP 55 // cdefgh|abcdefgh|abcdefg 56 #define EXTRAPOLATE(x, maxV) \ 57 { \ 58 (x) = ( (x) + (maxV) ) % (maxV); \ 59 } 60 #elif defined BORDER_REFLECT 61 // fedcba|abcdefgh|hgfedcb 62 #define EXTRAPOLATE(x, maxV) \ 63 { \ 64 (x) = min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) ); \ 65 } 66 #elif defined BORDER_REFLECT_101 67 // gfedcb|abcdefgh|gfedcba 68 #define EXTRAPOLATE(x, maxV) \ 69 { \ 70 (x) = min(((maxV)-1)*2-(x), max((x),-(x)) ); \ 71 } 72 #else 73 #error No extrapolation method 74 #endif 75 76 #if CN != 3 77 #define loadpix(addr) *(__global const srcT *)(addr) 78 #define storepix(val, addr) *(__global dstT *)(addr) = val 79 #define SRCSIZE (int)sizeof(srcT) 80 #define DSTSIZE (int)sizeof(dstT) 81 #else 82 #define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) 83 #define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr)) 84 #define SRCSIZE (int)sizeof(srcT1)*3 85 #define DSTSIZE (int)sizeof(dstT1)*3 86 #endif 87 88 #define SRC(_x,_y) convertToWT(loadpix(Src + mad24(_y, src_step, SRCSIZE * _x))) 89 90 #ifdef BORDER_CONSTANT 91 // CCCCCC|abcdefgh|CCCCCCC 92 #define ELEM(_x,_y,r_edge,t_edge,const_v) (_x)<0 | (_x) >= (r_edge) | (_y)<0 | (_y) >= (t_edge) ? (const_v) : SRC((_x),(_y)) 93 #else 94 #define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y)) 95 #endif 96 97 // horizontal and vertical filter kernels 98 // should be defined on host during compile time to avoid overhead 99 #define DIG(a) a, 100 __constant WT1 mat_kernelX[] = { KERNEL_MATRIX_X }; 101 __constant WT1 mat_kernelY[] = { KERNEL_MATRIX_Y }; 102 103 __kernel void laplacian(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width, 104 __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, 105 WT1 scale, WT1 delta) 106 { 107 __local WT lsmem[BLK_Y + 2 * RADIUS][BLK_X + 2 * RADIUS]; 108 __local WT lsmemDy1[BLK_Y][BLK_X + 2 * RADIUS]; 109 __local WT lsmemDy2[BLK_Y][BLK_X + 2 * RADIUS]; 110 111 int lix = get_local_id(0); 112 int liy = get_local_id(1); 113 114 int x = get_global_id(0); 115 116 int srcX = x + srcOffsetX - RADIUS; 117 118 int clocY = liy; 119 do 120 { 121 int yb = clocY + srcOffsetY - RADIUS; 122 EXTRAPOLATE(yb, (height)); 123 124 int clocX = lix; 125 int cSrcX = srcX; 126 do 127 { 128 int xb = cSrcX; 129 EXTRAPOLATE(xb,(width)); 130 lsmem[clocY][clocX] = ELEM(xb, yb, (width), (height), 0 ); 131 132 clocX += BLK_X; 133 cSrcX += BLK_X; 134 } 135 while(clocX < BLK_X+(RADIUS*2)); 136 137 clocY += BLK_Y; 138 } 139 while (clocY < BLK_Y+(RADIUS*2)); 140 barrier(CLK_LOCAL_MEM_FENCE); 141 142 WT scale_v = (WT)scale; 143 WT delta_v = (WT)delta; 144 for (int y = 0; y < dst_rows; y+=BLK_Y) 145 { 146 int i, clocX = lix; 147 WT sum1 = (WT) 0; 148 WT sum2 = (WT) 0; 149 do 150 { 151 sum1 = (WT) 0; 152 sum2 = (WT) 0; 153 for (i=0; i<=2*RADIUS; i++) 154 { 155 sum1 = mad(lsmem[liy + i][clocX], mat_kernelY[i], sum1); 156 sum2 = mad(lsmem[liy + i][clocX], mat_kernelX[i], sum2); 157 } 158 lsmemDy1[liy][clocX] = sum1; 159 lsmemDy2[liy][clocX] = sum2; 160 clocX += BLK_X; 161 } 162 while(clocX < BLK_X+(RADIUS*2)); 163 barrier(CLK_LOCAL_MEM_FENCE); 164 165 if ((x < dst_cols) && (y + liy < dst_rows)) 166 { 167 sum1 = (WT) 0; 168 sum2 = (WT) 0; 169 for (i=0; i<=2*RADIUS; i++) 170 { 171 sum1 = mad(lsmemDy1[liy][lix+i], mat_kernelX[i], sum1); 172 sum2 = mad(lsmemDy2[liy][lix+i], mat_kernelY[i], sum2); 173 } 174 175 WT sum = mad(scale_v, (sum1 + sum2), delta_v); 176 storepix(convertToDT(sum), Dst + mad24(y + liy, dst_step, mad24(x, DSTSIZE, dst_offset))); 177 } 178 179 for (int i = liy * BLK_X + lix; i < (RADIUS*2) * (BLK_X+(RADIUS*2)); i += BLK_X * BLK_Y) 180 { 181 int clocX = i % (BLK_X+(RADIUS*2)); 182 int clocY = i / (BLK_X+(RADIUS*2)); 183 lsmem[clocY][clocX] = lsmem[clocY + BLK_Y][clocX]; 184 } 185 barrier(CLK_LOCAL_MEM_FENCE); 186 187 int yb = y + liy + BLK_Y + srcOffsetY + RADIUS; 188 EXTRAPOLATE(yb, (height)); 189 190 clocX = lix; 191 int cSrcX = x + srcOffsetX - RADIUS; 192 do 193 { 194 int xb = cSrcX; 195 EXTRAPOLATE(xb,(width)); 196 lsmem[liy + 2*RADIUS][clocX] = ELEM(xb, yb, (width), (height), 0 ); 197 198 clocX += BLK_X; 199 cSrcX += BLK_X; 200 } 201 while(clocX < BLK_X+(RADIUS*2)); 202 barrier(CLK_LOCAL_MEM_FENCE); 203 } 204 } 205 206 #endif
