modules/core/src/opencl/flip.cl

   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) 2013, OpenCV Foundation, all rights reserved.
  14 // Third party copyrights are property of their respective owners.
  15 //
  16 // Redistribution and use in source and binary forms, with or without modification,
  17 // are permitted provided that the following conditions are met:
  18 //
  19 //   * Redistribution's of source code must retain the above copyright notice,
  20 //     this list of conditions and the following disclaimer.
  21 //
  22 //   * Redistribution's in binary form must reproduce the above copyright notice,
  23 //     this list of conditions and the following disclaimer in the documentation
  24 //     and/or other materials provided with the distribution.
  25 //
  26 //   * The name of the copyright holders may not be used to endorse or promote products
  27 //     derived from this software without specific prior written permission.
  28 //
  29 // This software is provided by the copyright holders and contributors "as is" and
  30 // any express or implied warranties, including, but not limited to, the implied
  31 // warranties of merchantability and fitness for a particular purpose are disclaimed.
  32 // In no event shall the OpenCV Foundation or contributors be liable for any direct,
  33 // indirect, incidental, special, exemplary, or consequential damages
  34 // (including, but not limited to, procurement of substitute goods or services;
  35 // loss of use, data, or profits; or business interruption) however caused
  36 // and on any theory of liability, whether in contract, strict liability,
  37 // or tort (including negligence or otherwise) arising in any way out of
  38 // the use of this software, even if advised of the possibility of such damage.
  39 //
  40 //M*/
  41
  42 #if cn != 3
  43 #define loadpix(addr) *(__global const T *)(addr)
  44 #define storepix(val, addr)  *(__global T *)(addr) = val
  45 #define TSIZE (int)sizeof(T)
  46 #else
  47 #define loadpix(addr) vload3(0, (__global const T1 *)(addr))
  48 #define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr))
  49 #define TSIZE ((int)sizeof(T1)*3)
  50 #endif
  51
  52 __kernel void arithm_flip_rows(__global const uchar * srcptr, int src_step, int src_offset,
  53                                __global uchar * dstptr, int dst_step, int dst_offset,
  54                                int rows, int cols, int thread_rows, int thread_cols)
  55 {
  56     int x = get_global_id(0);
  57     int y = get_global_id(1)*PIX_PER_WI_Y;
  58
  59     if (x < cols)
  60     {
  61         #pragma unroll
  62         for (int cy = 0; cy < PIX_PER_WI_Y && y < thread_rows; ++cy, ++y)
  63         {
  64             T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
  65             T src1 = loadpix(srcptr + mad24(rows - y - 1, src_step, mad24(x, TSIZE, src_offset)));
  66
  67             storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
  68             storepix(src0, dstptr + mad24(rows - y - 1, dst_step, mad24(x, TSIZE, dst_offset)));
  69         }
  70     }
  71 }
  72
  73 __kernel void arithm_flip_rows_cols(__global const uchar * srcptr, int src_step, int src_offset,
  74                                     __global uchar * dstptr, int dst_step, int dst_offset,
  75                                     int rows, int cols, int thread_rows, int thread_cols)
  76 {
  77     int x = get_global_id(0);
  78     int y = get_global_id(1)*PIX_PER_WI_Y;
  79
  80     if (x < cols)
  81     {
  82         int x1 = cols - x - 1;
  83         #pragma unroll
  84         for (int cy = 0; cy < PIX_PER_WI_Y && y < thread_rows; ++cy, ++y)
  85         {
  86             T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
  87             T src1 = loadpix(srcptr + mad24(rows - y - 1, src_step, mad24(x1, TSIZE, src_offset)));
  88
  89             storepix(src0, dstptr + mad24(rows - y - 1, dst_step, mad24(x1, TSIZE, dst_offset)));
  90             storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
  91         }
  92     }
  93 }
  94
  95 __kernel void arithm_flip_cols(__global const uchar * srcptr, int src_step, int src_offset,
  96                                __global uchar * dstptr, int dst_step, int dst_offset,
  97                                int rows, int cols, int thread_rows, int thread_cols)
  98 {
  99     int x = get_global_id(0);
 100     int y = get_global_id(1)*PIX_PER_WI_Y;
 101
 102     if (x < thread_cols)
 103     {
 104         int x1 = cols - x - 1;
 105         #pragma unroll
 106         for (int cy = 0; cy < PIX_PER_WI_Y && y < rows; ++cy, ++y)
 107         {
 108             T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
 109             T src1 = loadpix(srcptr + mad24(y, src_step, mad24(x1, TSIZE, src_offset)));
 110
 111             storepix(src0, dstptr + mad24(y, dst_step, mad24(x1, TSIZE, dst_offset)));
 112             storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
 113         }
 114     }
 115 }