modules/imgproc/src/opencl/laplacian5.cl

   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) = max(min((x), (maxV) - 1), 0); \
  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