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 #if cn != 3
      9 #define loadpix(addr) *(__global const T *)(addr)
     10 #define storepix(val, addr)  *(__global T *)(addr) = val
     11 #define TSIZE (int)sizeof(T)
     12 #else
     13 #define loadpix(addr) vload3(0, (__global const T1 *)(addr))
     14 #define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr))
     15 #define TSIZE ((int)sizeof(T1)*3)
     16 #endif
     17 
     18 __kernel void repeat(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
     19                      __global uchar * dstptr, int dst_step, int dst_offset)
     20 {
     21     int x = get_global_id(0);
     22     int y0 = get_global_id(1) * rowsPerWI;
     23 
     24     if (x < src_cols)
     25     {
     26         int src_index = mad24(y0, src_step, mad24(x, (int)sizeof(T), src_offset));
     27         int dst_index0 = mad24(y0, dst_step, mad24(x, (int)sizeof(T), dst_offset));
     28 
     29         for (int y = y0, y1 = min(src_rows, y0 + rowsPerWI); y < y1; ++y, src_index += src_step, dst_index0 += dst_step)
     30         {
     31             T srcelem = loadpix(srcptr + src_index);
     32 
     33             #pragma unroll
     34             for (int ey = 0; ey < ny; ++ey)
     35             {
     36                 int dst_index = mad24(ey * src_rows, dst_step, dst_index0);
     37 
     38                 #pragma unroll
     39                 for (int ex = 0; ex < nx; ++ex)
     40                 {
     41                     storepix(srcelem, dstptr + dst_index);
     42                     dst_index = mad24(src_cols, (int)sizeof(T), dst_index);
     43                 }
     44             }
     45         }
     46     }
     47 }
     48