1 /*M///////////////////////////////////////////////////////////////////////////////////////
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install,
7 // copy or use the software.
11 // For Open Source Computer Vision Library
13 // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
14 // Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
15 // Copyright (C) 2013, OpenCV Foundation, all rights reserved.
16 // Third party copyrights are property of their respective owners.
19 // Jia Haipeng, jiahaipeng95@gmail.com
22 // Redistribution and use in source and binary forms, with or without modification,
23 // are permitted provided that the following conditions are met:
25 // * Redistribution's of source code must retain the above copyright notice,
26 // this list of conditions and the following disclaimer.
28 // * Redistribution's in binary form must reproduce the above copyright notice,
29 // this list of conditions and the following disclaimer in the documentation
30 // and/or other materials provided with the distribution.
32 // * The name of the copyright holders may not be used to endorse or promote products
33 // derived from this software without specific prior written permission.
35 // This software is provided by the copyright holders and contributors as is and
36 // any express or implied warranties, including, but not limited to, the implied
37 // warranties of merchantability and fitness for a particular purpose are disclaimed.
38 // In no event shall the copyright holders or contributors be liable for any direct,
39 // indirect, incidental, special, exemplary, or consequential damages
40 // (including, but not limited to, procurement of substitute goods or services;
41 // loss of use, data, or profits; or business interruption) however caused
42 // and on any theory of liability, whether in contract, strict liability,
43 // or tort (including negligence or otherwise) arising in any way out of
44 // the use of this software, even if advised of the possibility of such damage.
50 after compiling this program user gets a single kernel called KF.
51 the following flags should be passed:
52 1) one of "-D BINARY_OP", "-D UNARY_OP", "-D MASK_BINARY_OP" or "-D MASK_UNARY_OP"
53 2) the actual operation performed, one of "-D OP_...", see below the list of operations.
54 2a) "-D dstDepth=<destination depth> [-D cn=<num channels]"
55 for some operations, like min/max/and/or/xor it's enough
56 2b) "-D srcDepth1=<source1 depth> -D srcDepth2=<source2 depth> -D dstDepth=<destination depth>
57 -D workDepth=<work depth> [-D cn=<num channels>]" - for mixed-type operations
62 #pragma OPENCL EXTENSION cl_amd_fp64:enable
63 #elif defined cl_khr_fp64
64 #pragma OPENCL EXTENSION cl_khr_fp64:enable
82 #define srcT1_C1 srcT1
83 #define srcT2_C1 srcT2
88 #define storedst(val) *(__global dstT *)(dstptr + dst_index) = val
89 #define storedst2(val) *(__global dstT *)(dstptr2 + dst_index2) = val
91 #define storedst(val) vstore3(val, 0, (__global dstT_C1 *)(dstptr + dst_index))
92 #define storedst2(val) vstore3(val, 0, (__global dstT_C1 *)(dstptr2 + dst_index2))
104 #define srcT1_C1 dstT_C1
112 #define srcT2_C1 dstT_C1
117 #define srcelem1 *(__global srcT1 *)(srcptr1 + src1_index)
118 #define srcelem2 *(__global srcT2 *)(srcptr2 + src2_index)
120 #define srcelem1 vload3(0, (__global srcT1_C1 *)(srcptr1 + src1_index))
121 #define srcelem2 vload3(0, (__global srcT2_C1 *)(srcptr2 + src2_index))
124 #define convertToDT noconvert
130 #define convertToWT2 convertToWT1
133 #define srcelem1 convertToWT1(*(__global srcT1 *)(srcptr1 + src1_index))
134 #define srcelem2 convertToWT2(*(__global srcT2 *)(srcptr2 + src2_index))
136 #define srcelem1 convertToWT1(vload3(0, (__global srcT1_C1 *)(srcptr1 + src1_index)))
137 #define srcelem2 convertToWT2(vload3(0, (__global srcT2_C1 *)(srcptr2 + src2_index)))
150 #define PROCESS_ELEM storedst(convertToDT(srcelem1 + srcelem2))
153 #define PROCESS_ELEM storedst(convertToDT(srcelem1 - srcelem2))
155 #elif defined OP_RSUB
156 #define PROCESS_ELEM storedst(convertToDT(srcelem2 - srcelem1))
158 #elif defined OP_ABSDIFF
159 #define PROCESS_ELEM \
160 workT v = srcelem1 - srcelem2; \
161 storedst(convertToDT(v >= (workT)(0) ? v : -v))
164 #define PROCESS_ELEM storedst(srcelem1 & srcelem2)
167 #define PROCESS_ELEM storedst(srcelem1 | srcelem2)
170 #define PROCESS_ELEM storedst(srcelem1 ^ srcelem2)
173 #define PROCESS_ELEM storedst(~srcelem1)
176 #define PROCESS_ELEM storedst(min(srcelem1, srcelem2))
179 #define PROCESS_ELEM storedst(max(srcelem1, srcelem2))
182 #define PROCESS_ELEM storedst(convertToDT(srcelem1 * srcelem2))
184 #elif defined OP_MUL_SCALE
187 #define EXTRA_PARAMS , workST srcelem2_, scaleT scale
189 #define srcelem2 srcelem2_
191 #define EXTRA_PARAMS , scaleT scale
193 #define PROCESS_ELEM storedst(convertToDT(srcelem1 * scale * srcelem2))
196 #define PROCESS_ELEM \
197 workT e2 = srcelem2, zero = (workT)(0); \
198 storedst(convertToDT(e2 != zero ? srcelem1 / e2 : zero))
200 #elif defined OP_DIV_SCALE
203 #define EXTRA_PARAMS , workST srcelem2_, scaleT scale
205 #define srcelem2 srcelem2_
207 #define EXTRA_PARAMS , scaleT scale
209 #define PROCESS_ELEM \
210 workT e2 = srcelem2, zero = (workT)(0); \
211 storedst(convertToDT(e2 == zero ? zero : (srcelem1 * (workT)(scale) / e2)))
213 #elif defined OP_RDIV_SCALE
216 #define EXTRA_PARAMS , workST srcelem2_, scaleT scale
218 #define srcelem2 srcelem2_
220 #define EXTRA_PARAMS , scaleT scale
222 #define PROCESS_ELEM \
223 workT e1 = srcelem1, zero = (workT)(0); \
224 storedst(convertToDT(e1 == zero ? zero : (srcelem2 * (workT)(scale) / e1)))
226 #elif defined OP_RECIP_SCALE
228 #define EXTRA_PARAMS , scaleT scale
229 #define PROCESS_ELEM \
230 workT e1 = srcelem1, zero = (workT)(0); \
231 storedst(convertToDT(e1 != zero ? scale / e1 : zero))
233 #elif defined OP_ADDW
235 #define EXTRA_PARAMS , scaleT alpha, scaleT beta, scaleT gamma
237 #define PROCESS_ELEM storedst(convertToDT(mad24(srcelem1, alpha, mad24(srcelem2, beta, gamma))))
239 #define PROCESS_ELEM storedst(convertToDT(mad(srcelem1, alpha, mad(srcelem2, beta, gamma))))
243 #define PROCESS_ELEM storedst(hypot(srcelem1, srcelem2))
245 #elif defined OP_ABS_NOSAT
246 #define PROCESS_ELEM \
247 dstT v = convertToDT(srcelem1); \
248 storedst(v >= 0 ? v : -v)
250 #elif defined OP_PHASE_RADIANS
251 #define PROCESS_ELEM \
252 workT tmp = atan2(srcelem2, srcelem1); \
253 if(tmp < 0) tmp += 6.283185307179586232f; \
256 #elif defined OP_PHASE_DEGREES
257 #define PROCESS_ELEM \
258 workT tmp = atan2(srcelem2, srcelem1)*57.29577951308232286465f; \
259 if(tmp < 0) tmp += 360; \
263 #define PROCESS_ELEM storedst(exp(srcelem1))
266 #define PROCESS_ELEM storedst(pow(srcelem1, srcelem2))
268 #elif defined OP_POWN
271 #define PROCESS_ELEM storedst(pown(srcelem1, srcelem2))
273 #elif defined OP_SQRT
274 #define PROCESS_ELEM storedst(sqrt(srcelem1))
277 #define PROCESS_ELEM \
278 dstT v = (dstT)(srcelem1);\
279 storedst(v > (dstT)(0) ? log(v) : log(-v))
284 #define PROCESS_ELEM storedst((dstT)(srcelem1 CMP_OPERATOR srcelem2 ? (dstT)(255) : (dstT)(0)))
286 #elif defined OP_CONVERT_SCALE_ABS
288 #define EXTRA_PARAMS , workT1 alpha, workT1 beta
290 #define PROCESS_ELEM \
291 workT value = mad24(srcelem1, (workT)(alpha), (workT)(beta)); \
292 storedst(convertToDT(value >= 0 ? value : -value))
294 #define PROCESS_ELEM \
295 workT value = mad(srcelem1, (workT)(alpha), (workT)(beta)); \
296 storedst(convertToDT(value >= 0 ? value : -value))
299 #elif defined OP_SCALE_ADD
301 #define EXTRA_PARAMS , workT1 alpha
303 #define PROCESS_ELEM storedst(convertToDT(mad24(srcelem1, (workT)(alpha), srcelem2)))
305 #define PROCESS_ELEM storedst(convertToDT(mad(srcelem1, (workT)(alpha), srcelem2)))
308 #elif defined OP_CTP_AD || defined OP_CTP_AR
310 #define CV_EPSILON FLT_EPSILON
312 #define CV_EPSILON DBL_EPSILON
315 #define TO_DEGREE cartToPolar *= (180 / CV_PI);
316 #elif defined OP_CTP_AR
319 #define PROCESS_ELEM \
320 dstT x = srcelem1, y = srcelem2; \
321 dstT x2 = x * x, y2 = y * y; \
322 dstT magnitude = sqrt(x2 + y2); \
323 dstT tmp = y >= 0 ? 0 : CV_PI * 2; \
324 tmp = x < 0 ? CV_PI : tmp; \
325 dstT tmp1 = y >= 0 ? CV_PI * 0.5f : CV_PI * 1.5f; \
326 dstT cartToPolar = y2 <= x2 ? x * y / mad((dstT)(0.28f), y2, x2 + CV_EPSILON) + tmp : (tmp1 - x * y / mad((dstT)(0.28f), x2, y2 + CV_EPSILON)); \
328 storedst(magnitude); \
329 storedst2(cartToPolar)
331 #elif defined OP_PTC_AD || defined OP_PTC_AR
333 #define FROM_DEGREE \
334 dstT ascale = CV_PI/180.0f; \
335 dstT alpha = y * ascale
337 #define FROM_DEGREE \
340 #define PROCESS_ELEM \
341 dstT x = srcelem1, y = srcelem2; \
343 storedst(cos(alpha) * x); \
344 storedst2(sin(alpha) * x)
346 #elif defined OP_PATCH_NANS
348 #define EXTRA_PARAMS , int val
349 #define PROCESS_ELEM \
350 if (( srcelem1 & 0x7fffffff) > 0x7f800000 ) \
354 #error "unknown op type"
357 #if defined OP_CTP_AD || defined OP_CTP_AR || defined OP_PTC_AD || defined OP_PTC_AR
359 #define EXTRA_PARAMS , __global uchar* dstptr2, int dststep2, int dstoffset2
361 #define EXTRA_INDEX int dst_index2 = mad24(y, dststep2, mad24(x, (int)sizeof(dstT_C1) * cn, dstoffset2))
364 #if defined UNARY_OP || defined MASK_UNARY_OP
366 #if defined OP_AND || defined OP_OR || defined OP_XOR || defined OP_ADD || defined OP_SAT_ADD || \
367 defined OP_SUB || defined OP_SAT_SUB || defined OP_RSUB || defined OP_SAT_RSUB || \
368 defined OP_ABSDIFF || defined OP_CMP || defined OP_MIN || defined OP_MAX || defined OP_POW || \
369 defined OP_MUL || defined OP_DIV || defined OP_POWN
371 #define EXTRA_PARAMS , workST srcelem2_
373 #define srcelem2 srcelem2_
378 #define srcelem2 (workT)(srcelem2_.x, srcelem2_.y, srcelem2_.z)
383 #if defined BINARY_OP
385 __kernel void KF(__global const uchar * srcptr1, int srcstep1, int srcoffset1,
386 __global const uchar * srcptr2, int srcstep2, int srcoffset2,
387 __global uchar * dstptr, int dststep, int dstoffset,
388 int rows, int cols EXTRA_PARAMS )
390 int x = get_global_id(0);
391 int y = get_global_id(1);
393 if (x < cols && y < rows)
395 int src1_index = mad24(y, srcstep1, mad24(x, (int)sizeof(srcT1_C1) * cn, srcoffset1));
396 #if !(defined(OP_RECIP_SCALE) || defined(OP_NOT))
397 int src2_index = mad24(y, srcstep2, mad24(x, (int)sizeof(srcT2_C1) * cn, srcoffset2));
399 int dst_index = mad24(y, dststep, mad24(x, (int)sizeof(dstT_C1) * cn, dstoffset));
406 #elif defined MASK_BINARY_OP
408 __kernel void KF(__global const uchar * srcptr1, int srcstep1, int srcoffset1,
409 __global const uchar * srcptr2, int srcstep2, int srcoffset2,
410 __global const uchar * mask, int maskstep, int maskoffset,
411 __global uchar * dstptr, int dststep, int dstoffset,
412 int rows, int cols EXTRA_PARAMS )
414 int x = get_global_id(0);
415 int y = get_global_id(1);
417 if (x < cols && y < rows)
419 int mask_index = mad24(y, maskstep, x + maskoffset);
420 if( mask[mask_index] )
422 int src1_index = mad24(y, srcstep1, mad24(x, (int)sizeof(srcT1_C1) * cn, srcoffset1));
423 int src2_index = mad24(y, srcstep2, mad24(x, (int)sizeof(srcT2_C1) * cn, srcoffset2));
424 int dst_index = mad24(y, dststep, mad24(x, (int)sizeof(dstT_C1) * cn, dstoffset));
431 #elif defined UNARY_OP
433 __kernel void KF(__global const uchar * srcptr1, int srcstep1, int srcoffset1,
434 __global uchar * dstptr, int dststep, int dstoffset,
435 int rows, int cols EXTRA_PARAMS )
437 int x = get_global_id(0);
438 int y = get_global_id(1);
440 if (x < cols && y < rows)
442 int src1_index = mad24(y, srcstep1, mad24(x, (int)sizeof(srcT1_C1) * cn, srcoffset1));
443 int dst_index = mad24(y, dststep, mad24(x, (int)sizeof(dstT_C1) * cn, dstoffset));
449 #elif defined MASK_UNARY_OP
451 __kernel void KF(__global const uchar * srcptr1, int srcstep1, int srcoffset1,
452 __global const uchar * mask, int maskstep, int maskoffset,
453 __global uchar * dstptr, int dststep, int dstoffset,
454 int rows, int cols EXTRA_PARAMS )
456 int x = get_global_id(0);
457 int y = get_global_id(1);
459 if (x < cols && y < rows)
461 int mask_index = mad24(y, maskstep, x + maskoffset);
462 if( mask[mask_index] )
464 int src1_index = mad24(y, srcstep1, mad24(x, (int)sizeof(srcT1_C1) * cn, srcoffset1));
465 int dst_index = mad24(y, dststep, mad24(x, (int)sizeof(dstT_C1) * cn, dstoffset));
474 #error "Unknown operation type"