Home | History | Annotate | Download | only in opencl 
      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