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