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