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#ifdef DOUBLE_SUPPORT 9#ifdef cl_amd_fp64 10#pragma OPENCL EXTENSION cl_amd_fp64:enable 11#elif defined (cl_khr_fp64) 12#pragma OPENCL EXTENSION cl_khr_fp64:enable 13#endif 14#endif 15 16#define noconvert 17 18#if cn != 3 19#define loadpix(addr) *(__global const srcT *)(addr) 20#define storepix(val, addr) *(__global dstT *)(addr) = val 21#define srcTSIZE (int)sizeof(srcT) 22#define dstTSIZE (int)sizeof(dstT) 23#else 24#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) 25#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr)) 26#define srcTSIZE ((int)sizeof(srcT1)*3) 27#define dstTSIZE ((int)sizeof(dstT1)*3) 28#endif 29 30__kernel void normalizek(__global const uchar * srcptr, int src_step, int src_offset, 31 __global const uchar * mask, int mask_step, int mask_offset, 32 __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols 33#ifdef HAVE_SCALE 34 , float scale 35#endif 36#ifdef HAVE_DELTA 37 , float delta 38#endif 39 ) 40{ 41 int x = get_global_id(0); 42 int y0 = get_global_id(1) * rowsPerWI; 43 44 if (x < dst_cols) 45 { 46 int src_index = mad24(y0, src_step, mad24(x, srcTSIZE, src_offset)); 47 int mask_index = mad24(y0, mask_step, x + mask_offset); 48 int dst_index = mad24(y0, dst_step, mad24(x, dstTSIZE, dst_offset)); 49 50 for (int y = y0, y1 = min(y0 + rowsPerWI, dst_rows); y < y1; 51 ++y, src_index += src_step, dst_index += dst_step, mask_index += mask_step) 52 { 53 if (mask[mask_index]) 54 { 55 workT value = convertToWT(loadpix(srcptr + src_index)); 56#ifdef HAVE_SCALE 57#ifdef HAVE_DELTA 58 value = fma(value, (workT)(scale), (workT)(delta)); 59#else 60 value *= (workT)(scale); 61#endif 62#else // not scale 63#ifdef HAVE_DELTA 64 value += (workT)(delta); 65#endif 66#endif 67 68 storepix(convertToDT(value), dstptr + dst_index); 69 } 70 } 71 } 72} 73