1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2014, Itseez, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
#ifdef DOUBLE_SUPPORT
#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64:enable
#elif defined (cl_khr_fp64)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#endif
#ifdef DEPTH_0
#define MIN_VAL 0
#define MAX_VAL UCHAR_MAX
#elif defined DEPTH_1
#define MIN_VAL SCHAR_MIN
#define MAX_VAL SCHAR_MAX
#elif defined DEPTH_2
#define MIN_VAL 0
#define MAX_VAL USHRT_MAX
#elif defined DEPTH_3
#define MIN_VAL SHRT_MIN
#define MAX_VAL SHRT_MAX
#elif defined DEPTH_4
#define MIN_VAL INT_MIN
#define MAX_VAL INT_MAX
#elif defined DEPTH_5
#define MIN_VAL (-FLT_MAX)
#define MAX_VAL FLT_MAX
#elif defined DEPTH_6
#define MIN_VAL (-DBL_MAX)
#define MAX_VAL DBL_MAX
#endif
#define noconvert
#define INDEX_MAX UINT_MAX
#if wdepth <= 4
#define MIN_ABS(a) convertFromU(abs(a))
#define MIN_ABS2(a, b) convertFromU(abs_diff(a, b))
#define MIN(a, b) min(a, b)
#define MAX(a, b) max(a, b)
#else
#define MIN_ABS(a) fabs(a)
#define MIN_ABS2(a, b) fabs(a - b)
#define MIN(a, b) fmin(a, b)
#define MAX(a, b) fmax(a, b)
#endif
#if kercn != 3
#define loadpix(addr) *(__global const srcT *)(addr)
#define srcTSIZE (int)sizeof(srcT)
#else
#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
#define srcTSIZE ((int)sizeof(srcT1) * 3)
#endif
#ifndef HAVE_MASK
#undef srcTSIZE
#define srcTSIZE (int)sizeof(srcT1)
#endif
#ifdef NEED_MINVAL
#ifdef NEED_MINLOC
#define CALC_MIN(p, inc) \
if (minval > temp.p) \
{ \
minval = temp.p; \
minloc = id + inc; \
}
#else
#define CALC_MIN(p, inc) \
minval = MIN(minval, temp.p);
#endif
#else
#define CALC_MIN(p, inc)
#endif
#ifdef NEED_MAXVAL
#ifdef NEED_MAXLOC
#define CALC_MAX(p, inc) \
if (maxval < temp.p) \
{ \
maxval = temp.p; \
maxloc = id + inc; \
}
#else
#define CALC_MAX(p, inc) \
maxval = MAX(maxval, temp.p);
#endif
#else
#define CALC_MAX(p, inc)
#endif
#ifdef OP_CALC2
#define CALC_MAX2(p) \
maxval2 = MAX(maxval2, temp2.p);
#else
#define CALC_MAX2(p)
#endif
#define CALC_P(p, inc) \
CALC_MIN(p, inc) \
CALC_MAX(p, inc) \
CALC_MAX2(p)
__kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_offset, int cols,
int total, int groupnum, __global uchar * dstptr
#ifdef HAVE_MASK
, __global const uchar * mask, int mask_step, int mask_offset
#endif
#ifdef HAVE_SRC2
, __global const uchar * src2ptr, int src2_step, int src2_offset
#endif
)
{
int lid = get_local_id(0);
int gid = get_group_id(0);
int id = get_global_id(0)
#ifndef HAVE_MASK
* kercn;
#else
;
#endif
srcptr += src_offset;
#ifdef HAVE_MASK
mask += mask_offset;
#endif
#ifdef HAVE_SRC2
src2ptr += src2_offset;
#endif
#ifdef NEED_MINVAL
__local dstT1 localmem_min[WGS2_ALIGNED];
dstT1 minval = MAX_VAL;
#ifdef NEED_MINLOC
__local uint localmem_minloc[WGS2_ALIGNED];
uint minloc = INDEX_MAX;
#endif
#endif
#ifdef NEED_MAXVAL
dstT1 maxval = MIN_VAL;
__local dstT1 localmem_max[WGS2_ALIGNED];
#ifdef NEED_MAXLOC
__local uint localmem_maxloc[WGS2_ALIGNED];
uint maxloc = INDEX_MAX;
#endif
#endif
#ifdef OP_CALC2
__local dstT1 localmem_max2[WGS2_ALIGNED];
dstT1 maxval2 = MIN_VAL;
#endif
int src_index;
#ifdef HAVE_MASK
int mask_index;
#endif
#ifdef HAVE_SRC2
int src2_index;
#endif
dstT temp;
#ifdef HAVE_SRC2
dstT temp2;
#endif
for (int grain = groupnum * WGS
#ifndef HAVE_MASK
* kercn
#endif
; id < total; id += grain)
{
#ifdef HAVE_MASK
#ifdef HAVE_MASK_CONT
mask_index = id;
#else
mask_index = mad24(id / cols, mask_step, id % cols);
#endif
if (mask[mask_index])
#endif
{
#ifdef HAVE_SRC_CONT
src_index = id * srcTSIZE;//mul24(id, srcTSIZE);
#else
src_index = mad24(id / cols, src_step, mul24(id % cols, srcTSIZE));
#endif
temp = convertToDT(loadpix(srcptr + src_index));
#ifdef OP_ABS
temp = MIN_ABS(temp);
#endif
#ifdef HAVE_SRC2
#ifdef HAVE_SRC2_CONT
src2_index = id * srcTSIZE; //mul24(id, srcTSIZE);
#else
src2_index = mad24(id / cols, src2_step, mul24(id % cols, srcTSIZE));
#endif
temp2 = convertToDT(loadpix(src2ptr + src2_index));
temp = MIN_ABS2(temp, temp2);
#ifdef OP_CALC2
temp2 = MIN_ABS(temp2);
#endif
#endif
#if kercn == 1
#ifdef NEED_MINVAL
#ifdef NEED_MINLOC
if (minval > temp)
{
minval = temp;
minloc = id;
}
#else
minval = MIN(minval, temp);
#endif
#endif
#ifdef NEED_MAXVAL
#ifdef NEED_MAXLOC
if (maxval < temp)
{
maxval = temp;
maxloc = id;
}
#else
maxval = MAX(maxval, temp);
#endif
#ifdef OP_CALC2
maxval2 = MAX(maxval2, temp2);
#endif
#endif
#elif kercn >= 2
CALC_P(s0, 0)
CALC_P(s1, 1)
#if kercn >= 3
CALC_P(s2, 2)
#if kercn >= 4
CALC_P(s3, 3)
#if kercn >= 8
CALC_P(s4, 4)
CALC_P(s5, 5)
CALC_P(s6, 6)
CALC_P(s7, 7)
#if kercn == 16
CALC_P(s8, 8)
CALC_P(s9, 9)
CALC_P(sA, 10)
CALC_P(sB, 11)
CALC_P(sC, 12)
CALC_P(sD, 13)
CALC_P(sE, 14)
CALC_P(sF, 15)
#endif
#endif
#endif
#endif
#endif
}
}
if (lid < WGS2_ALIGNED)
{
#ifdef NEED_MINVAL
localmem_min[lid] = minval;
#endif
#ifdef NEED_MAXVAL
localmem_max[lid] = maxval;
#endif
#ifdef NEED_MINLOC
localmem_minloc[lid] = minloc;
#endif
#ifdef NEED_MAXLOC
localmem_maxloc[lid] = maxloc;
#endif
#ifdef OP_CALC2
localmem_max2[lid] = maxval2;
#endif
}
barrier(CLK_LOCAL_MEM_FENCE);
if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
{
int lid3 = lid - WGS2_ALIGNED;
#ifdef NEED_MINVAL
#ifdef NEED_MINLOC
if (localmem_min[lid3] >= minval)
{
if (localmem_min[lid3] == minval)
localmem_minloc[lid3] = min(localmem_minloc[lid3], minloc);
else
localmem_minloc[lid3] = minloc,
localmem_min[lid3] = minval;
}
#else
localmem_min[lid3] = MIN(localmem_min[lid3], minval);
#endif
#endif
#ifdef NEED_MAXVAL
#ifdef NEED_MAXLOC
if (localmem_max[lid3] <= maxval)
{
if (localmem_max[lid3] == maxval)
localmem_maxloc[lid3] = min(localmem_maxloc[lid3], maxloc);
else
localmem_maxloc[lid3] = maxloc,
localmem_max[lid3] = maxval;
}
#else
localmem_max[lid3] = MAX(localmem_max[lid3], maxval);
#endif
#endif
#ifdef OP_CALC2
localmem_max2[lid3] = MAX(localmem_max2[lid3], maxval2);
#endif
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
{
if (lid < lsize)
{
int lid2 = lsize + lid;
#ifdef NEED_MINVAL
#ifdef NEED_MINLOC
if (localmem_min[lid] >= localmem_min[lid2])
{
if (localmem_min[lid] == localmem_min[lid2])
localmem_minloc[lid] = min(localmem_minloc[lid2], localmem_minloc[lid]);
else
localmem_minloc[lid] = localmem_minloc[lid2],
localmem_min[lid] = localmem_min[lid2];
}
#else
localmem_min[lid] = MIN(localmem_min[lid], localmem_min[lid2]);
#endif
#endif
#ifdef NEED_MAXVAL
#ifdef NEED_MAXLOC
if (localmem_max[lid] <= localmem_max[lid2])
{
if (localmem_max[lid] == localmem_max[lid2])
localmem_maxloc[lid] = min(localmem_maxloc[lid2], localmem_maxloc[lid]);
else
localmem_maxloc[lid] = localmem_maxloc[lid2],
localmem_max[lid] = localmem_max[lid2];
}
#else
localmem_max[lid] = MAX(localmem_max[lid], localmem_max[lid2]);
#endif
#endif
#ifdef OP_CALC2
localmem_max2[lid] = MAX(localmem_max2[lid], localmem_max2[lid2]);
#endif
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if (lid == 0)
{
int pos = 0;
#ifdef NEED_MINVAL
*(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_min[0];
pos = mad24(groupnum, (int)sizeof(dstT1), pos);
#endif
#ifdef NEED_MAXVAL
*(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_max[0];
pos = mad24(groupnum, (int)sizeof(dstT1), pos);
#endif
#ifdef NEED_MINLOC
*(__global uint *)(dstptr + mad24(gid, (int)sizeof(uint), pos)) = localmem_minloc[0];
pos = mad24(groupnum, (int)sizeof(uint), pos);
#endif
#ifdef NEED_MAXLOC
*(__global uint *)(dstptr + mad24(gid, (int)sizeof(uint), pos)) = localmem_maxloc[0];
#ifdef OP_CALC2
pos = mad24(groupnum, (int)sizeof(uint), pos);
#endif
#endif
#ifdef OP_CALC2
*(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_max2[0];
#endif
}
}