1 /*M/////////////////////////////////////////////////////////////////////////////////////// 2 // 3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 4 // 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. 8 // 9 // 10 // License Agreement 11 // For Open Source Computer Vision Library 12 // 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. 17 // 18 // Redistribution and use in source and binary forms, with or without modification, 19 // are permitted provided that the following conditions are met: 20 // 21 // * Redistribution's of source code must retain the above copyright notice, 22 // this list of conditions and the following disclaimer. 23 // 24 // * Redistribution's in binary form must reproduce the above copyright notice, 25 // this list of conditions and the following disclaimer in the documentation 26 // and/or other materials provided with the distribution. 27 // 28 // * The name of the copyright holders may not be used to endorse or promote products 29 // derived from this software without specific prior written permission. 30 // 31 // This software is provided by the copyright holders and contributors as is and 32 // any express or implied warranties, including, but not limited to, the implied 33 // warranties of merchantability and fitness for a particular purpose are disclaimed. 34 // In no event shall the copyright holders or contributors be liable for any direct, 35 // indirect, incidental, special, exemplary, or consequential damages 36 // (including, but not limited to, procurement of substitute goods or services; 37 // loss of use, data, or profits; or business interruption) however caused 38 // and on any theory of liability, whether in contract, strict liability, 39 // or tort (including negligence or otherwise) arising in any way out of 40 // the use of this software, even if advised of the possibility of such damage. 41 // 42 //M*/ 43 44 #ifdef DOUBLE_SUPPORT 45 #ifdef cl_amd_fp64 46 #pragma OPENCL EXTENSION cl_amd_fp64:enable 47 #elif defined (cl_khr_fp64) 48 #pragma OPENCL EXTENSION cl_khr_fp64:enable 49 #endif 50 #endif 51 52 __kernel void inrange(__global const uchar * src1ptr, int src1_step, int src1_offset, 53 __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, 54 #ifdef HAVE_SCALAR 55 __global const srcT1 * src2, __global const srcT1 * src3, 56 #else 57 __global const uchar * src2ptr, int src2_step, int src2_offset, 58 __global const uchar * src3ptr, int src3_step, int src3_offset, 59 #endif 60 int rowsPerWI) 61 { 62 int x = get_global_id(0); 63 int y0 = get_global_id(1) * rowsPerWI; 64 65 if (x < dst_cols) 66 { 67 int src1_index = mad24(y0, src1_step, mad24(x, (int)sizeof(srcT1) * kercn, src1_offset)); 68 int dst_index = mad24(y0, dst_step, mad24(x, colsPerWI, dst_offset)); 69 #ifndef HAVE_SCALAR 70 int src2_index = mad24(y0, src2_step, mad24(x, (int)sizeof(srcT1) * kercn, src2_offset)); 71 int src3_index = mad24(y0, src3_step, mad24(x, (int)sizeof(srcT1) * kercn, src3_offset)); 72 #endif 73 74 for (int y = y0, y1 = min(dst_rows, y0 + rowsPerWI); y < y1; ++y, src1_index += src1_step, dst_index += dst_step) 75 { 76 #if kercn >= cn && kercn == 4 && depth <= 4 && !defined HAVE_SCALAR 77 srcT src1 = *(__global const srcT *)(src1ptr + src1_index); 78 srcT src2 = *(__global const srcT *)(src2ptr + src2_index); 79 srcT src3 = *(__global const srcT *)(src3ptr + src3_index); 80 __global dstT * dst = (__global dstT *)(dstptr + dst_index); 81 #if cn == 1 82 dst[0] = src2 > src1 || src3 < src1 ? (dstT)(0) : (dstT)(255); 83 #elif cn == 2 84 dst[0] = (dstT)(src2.xy > src1.xy || src3.xy < src1.xy || 85 src2.zw > src1.zw || src3.zw < src1.zw ? (dstT)(0) : (dstT)(255); 86 #elif cn == 4 87 dst[0] = (dstT)(src2.x > src1.x || src3.x < src1.x || 88 src2.y > src1.y || src3.y < src1.y || 89 src2.z > src1.z || src3.z < src1.z || 90 src2.w > src1.w || src3.w < src1.w ? 0 : 255); 91 #endif 92 #else 93 __global const srcT1 * src1 = (__global const srcT1 *)(src1ptr + src1_index); 94 __global uchar * dst = dstptr + dst_index; 95 #ifndef HAVE_SCALAR 96 __global const srcT1 * src2 = (__global const srcT1 *)(src2ptr + src2_index); 97 __global const srcT1 * src3 = (__global const srcT1 *)(src3ptr + src3_index); 98 #endif 99 100 #pragma unroll 101 for (int px = 0; px < colsPerWI; ++px, src1 += cn 102 #ifndef HAVE_SCALAR 103 , src2 += cn, src3 += cn 104 #endif 105 ) 106 { 107 dst[px] = 255; 108 109 for (int c = 0; c < cn; ++c) 110 if (src2[c] > src1[c] || src3[c] < src1[c]) 111 { 112 dst[px] = 0; 113 break; 114 } 115 } 116 #endif // kercn >= cn 117 #ifndef HAVE_SCALAR 118 src2_index += src2_step; 119 src3_index += src3_step; 120 #endif 121 } 122 } 123 } 124
