modules/ocl/src/opencl/imgproc_remap.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) 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 // Third party copyrights are property of their respective owners.
  16 //
  17 // @Authors
  18 //    Wu Zailong, bullet@yeah.net
  19 //
  20 // Redistribution and use in source and binary forms, with or without modification,
  21 // are permitted provided that the following conditions are met:
  22 //
  23 //   * Redistribution's of source code must retain the above copyright notice,
  24 //     this list of conditions and the following disclaimer.
  25 //
  26 //   * Redistribution's in binary form must reproduce the above copyright notice,
  27 //     this list of conditions and the following disclaimer in the documentation
  28 //     and/or other materials provided with the distribution.
  29 //
  30 //   * The name of the copyright holders may not be used to endorse or promote products
  31 //     derived from this software without specific prior written permission.
  32 //
  33 // This software is provided by the copyright holders and contributors as is and
  34 // any express or implied warranties, including, but not limited to, the implied
  35 // warranties of merchantability and fitness for a particular purpose are disclaimed.
  36 // In no event shall the Intel Corporation or contributors be liable for any direct,
  37 // indirect, incidental, special, exemplary, or consequential damages
  38 // (including, but not limited to, procurement of substitute goods or services;
  39 // loss of use, data, or profits; or business interruption) however caused
  40 // and on any theory of liability, whether in contract, strict liability,
  41 // or tort (including negligence or otherwise) arising in any way out of
  42 // the use of this software, even if advised of the possibility of such damage.
  43 //
  44 //M*/
  45
  46 #if defined (DOUBLE_SUPPORT)
  47 #ifdef cl_khr_fp64
  48 #pragma OPENCL EXTENSION cl_khr_fp64:enable
  49 #elif defined (cl_amd_fp64)
  50 #pragma OPENCL EXTENSION cl_amd_fp64:enable
  51 #endif
  52 #endif
  53
  54 enum
  55 {
  56     INTER_BITS = 5,
  57     INTER_TAB_SIZE = 1 << INTER_BITS,
  58     INTER_TAB_SIZE2 = INTER_TAB_SIZE * INTER_TAB_SIZE
  59 };
  60
  61 #ifdef INTER_NEAREST
  62 #define convertToWT
  63 #endif
  64
  65 #ifdef BORDER_CONSTANT
  66 #define EXTRAPOLATE(v2, v) v = scalar;
  67 #elif defined BORDER_REPLICATE
  68 #define EXTRAPOLATE(v2, v) \
  69     { \
  70         v2 = max(min(v2, (int2)(src_cols - 1, src_rows - 1)), (int2)(0)); \
  71         v = convertToWT(src[mad24(v2.y, src_step, v2.x + src_offset)]); \
  72     }
  73 #elif defined BORDER_WRAP
  74 #define EXTRAPOLATE(v2, v) \
  75     { \
  76         if (v2.x < 0) \
  77             v2.x -= ((v2.x - src_cols + 1) / src_cols) * src_cols; \
  78         if (v2.x >= src_cols) \
  79             v2.x %= src_cols; \
  80         \
  81         if (v2.y < 0) \
  82             v2.y -= ((v2.y - src_rows + 1) / src_rows) * src_rows; \
  83         if( v2.y >= src_rows ) \
  84             v2.y %= src_rows; \
  85         v = convertToWT(src[mad24(v2.y, src_step, v2.x + src_offset)]); \
  86     }
  87 #elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
  88 #ifdef BORDER_REFLECT
  89 #define DELTA int delta = 0
  90 #else
  91 #define DELTA int delta = 1
  92 #endif
  93 #define EXTRAPOLATE(v2, v) \
  94     { \
  95         DELTA; \
  96         if (src_cols == 1) \
  97             v2.x = 0; \
  98         else \
  99             do \
 100             { \
 101                 if( v2.x < 0 ) \
 102                     v2.x = -v2.x - 1 + delta; \
 103                 else \
 104                     v2.x = src_cols - 1 - (v2.x - src_cols) - delta; \
 105             } \
 106             while (v2.x >= src_cols || v2.x < 0); \
 107         \
 108         if (src_rows == 1) \
 109             v2.y = 0; \
 110         else \
 111             do \
 112             { \
 113                 if( v2.y < 0 ) \
 114                     v2.y = -v2.y - 1 + delta; \
 115                 else \
 116                     v2.y = src_rows - 1 - (v2.y - src_rows) - delta; \
 117             } \
 118             while (v2.y >= src_rows || v2.y < 0); \
 119         v = convertToWT(src[mad24(v2.y, src_step, v2.x + src_offset)]); \
 120     }
 121 #else
 122 #error No extrapolation method
 123 #endif
 124
 125 #define NEED_EXTRAPOLATION(gx, gy) (gx >= src_cols || gy >= src_rows || gx < 0 || gy < 0)
 126
 127 #ifdef INTER_NEAREST
 128
 129 __kernel void remap_2_32FC1(__global const T * restrict src, __global T * dst,
 130         __global float * map1, __global float * map2,
 131         int src_offset, int dst_offset, int map1_offset, int map2_offset,
 132         int src_step, int dst_step, int map1_step, int map2_step,
 133         int src_cols, int src_rows, int dst_cols, int dst_rows, T scalar)
 134 {
 135     int x = get_global_id(0);
 136     int y = get_global_id(1);
 137
 138     if (x < dst_cols && y < dst_rows)
 139     {
 140         int dstIdx = mad24(y, dst_step, x + dst_offset);
 141         int map1Idx = mad24(y, map1_step, x + map1_offset);
 142         int map2Idx = mad24(y, map2_step, x + map2_offset);
 143
 144         int gx = convert_int_sat_rte(map1[map1Idx]);
 145         int gy = convert_int_sat_rte(map2[map2Idx]);
 146
 147         if (NEED_EXTRAPOLATION(gx, gy))
 148         {
 149 #ifndef BORDER_CONSTANT
 150             int2 gxy = (int2)(gx, gy);
 151 #endif
 152             EXTRAPOLATE(gxy, dst[dstIdx]);
 153         }
 154         else
 155         {
 156             int srcIdx = mad24(gy, src_step, gx + src_offset);
 157             dst[dstIdx] = src[srcIdx];
 158         }
 159     }
 160 }
 161
 162 __kernel void remap_32FC2(__global const T * restrict src, __global T * dst, __global float2 * map1,
 163         int src_offset, int dst_offset, int map1_offset,
 164         int src_step, int dst_step, int map1_step,
 165         int src_cols, int src_rows, int dst_cols, int dst_rows, T scalar)
 166 {
 167     int x = get_global_id(0);
 168     int y = get_global_id(1);
 169
 170     if (x < dst_cols && y < dst_rows)
 171     {
 172         int dstIdx = mad24(y, dst_step, x + dst_offset);
 173         int map1Idx = mad24(y, map1_step, x + map1_offset);
 174
 175         int2 gxy = convert_int2_sat_rte(map1[map1Idx]);
 176         int gx = gxy.x, gy = gxy.y;
 177
 178         if (NEED_EXTRAPOLATION(gx, gy))
 179             EXTRAPOLATE(gxy, dst[dstIdx])
 180         else
 181         {
 182             int srcIdx = mad24(gy, src_step, gx + src_offset);
 183             dst[dstIdx] = src[srcIdx];
 184         }
 185     }
 186 }
 187
 188 __kernel void remap_16SC2(__global const T * restrict src, __global T * dst, __global short2 * map1,
 189         int src_offset, int dst_offset, int map1_offset,
 190         int src_step, int dst_step, int map1_step,
 191         int src_cols, int src_rows, int dst_cols, int dst_rows, T scalar)
 192 {
 193     int x = get_global_id(0);
 194     int y = get_global_id(1);
 195
 196     if (x < dst_cols && y < dst_rows)
 197     {
 198         int dstIdx = mad24(y, dst_step, x + dst_offset);
 199         int map1Idx = mad24(y, map1_step, x + map1_offset);
 200
 201         int2 gxy = convert_int2(map1[map1Idx]);
 202         int gx = gxy.x, gy = gxy.y;
 203
 204         if (NEED_EXTRAPOLATION(gx, gy))
 205             EXTRAPOLATE(gxy, dst[dstIdx])
 206         else
 207         {
 208             int srcIdx = mad24(gy, src_step, gx + src_offset);
 209             dst[dstIdx] = src[srcIdx];
 210         }
 211     }
 212 }
 213
 214 __kernel void remap_16SC2_16UC1(__global const T * restrict src, __global T * dst, __global short2 * map1, __global ushort * map2,
 215         int src_offset, int dst_offset, int map1_offset, int map2_offset,
 216         int src_step, int dst_step, int map1_step, int map2_step,
 217         int src_cols, int src_rows, int dst_cols, int dst_rows, T scalar)
 218 {
 219     int x = get_global_id(0);
 220     int y = get_global_id(1);
 221
 222     if (x < dst_cols && y < dst_rows)
 223     {
 224         int dstIdx = mad24(y, dst_step, x + dst_offset);
 225         int map1Idx = mad24(y, map1_step, x + map1_offset);
 226         int map2Idx = mad24(y, map2_step, x + map2_offset);
 227
 228         int map2Value = convert_int(map2[map2Idx]) & (INTER_TAB_SIZE2 - 1);
 229         int dx = (map2Value & (INTER_TAB_SIZE - 1)) < (INTER_TAB_SIZE >> 1) ? 1 : 0;
 230         int dy = (map2Value >> INTER_BITS) < (INTER_TAB_SIZE >> 1) ? 1 : 0;
 231         int2 gxy = convert_int2(map1[map1Idx]) + (int2)(dx, dy);
 232         int gx = gxy.x, gy = gxy.y;
 233
 234         if (NEED_EXTRAPOLATION(gx, gy))
 235             EXTRAPOLATE(gxy, dst[dstIdx])
 236         else
 237         {
 238             int srcIdx = mad24(gy, src_step, gx + src_offset);
 239             dst[dstIdx] = src[srcIdx];
 240         }
 241     }
 242 }
 243
 244 #elif INTER_LINEAR
 245
 246 __kernel void remap_2_32FC1(__global T const * restrict  src, __global T * dst,
 247         __global float * map1, __global float * map2,
 248         int src_offset, int dst_offset, int map1_offset, int map2_offset,
 249         int src_step, int dst_step, int map1_step, int map2_step,
 250         int src_cols, int src_rows, int dst_cols, int dst_rows, T nVal)
 251 {
 252     int x = get_global_id(0);
 253     int y = get_global_id(1);
 254
 255     if (x < dst_cols && y < dst_rows)
 256     {
 257         int dstIdx = mad24(y, dst_step, x + dst_offset);
 258         int map1Idx = mad24(y, map1_step, x + map1_offset);
 259         int map2Idx = mad24(y, map2_step, x + map2_offset);
 260
 261         float2 map_data = (float2)(map1[map1Idx], map2[map2Idx]);
 262
 263         int2 map_dataA = convert_int2_sat_rtn(map_data);
 264         int2 map_dataB = (int2)(map_dataA.x + 1, map_dataA.y);
 265         int2 map_dataC = (int2)(map_dataA.x, map_dataA.y + 1);
 266         int2 map_dataD = (int2)(map_dataA.x + 1, map_dataA.y +1);
 267
 268         float2 _u = map_data - convert_float2(map_dataA);
 269         WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)INTER_TAB_SIZE)) / (WT2)INTER_TAB_SIZE;
 270         WT scalar = convertToWT(nVal);
 271         WT a = scalar, b = scalar, c = scalar, d = scalar;
 272
 273         if (!NEED_EXTRAPOLATION(map_dataA.x, map_dataA.y))
 274             a = convertToWT(src[mad24(map_dataA.y, src_step, map_dataA.x + src_offset)]);
 275         else
 276             EXTRAPOLATE(map_dataA, a);
 277
 278         if (!NEED_EXTRAPOLATION(map_dataB.x, map_dataB.y))
 279             b = convertToWT(src[mad24(map_dataB.y, src_step, map_dataB.x + src_offset)]);
 280         else
 281             EXTRAPOLATE(map_dataB, b);
 282
 283         if (!NEED_EXTRAPOLATION(map_dataC.x, map_dataC.y))
 284             c = convertToWT(src[mad24(map_dataC.y, src_step, map_dataC.x + src_offset)]);
 285         else
 286             EXTRAPOLATE(map_dataC, c);
 287
 288         if (!NEED_EXTRAPOLATION(map_dataD.x, map_dataD.y))
 289             d = convertToWT(src[mad24(map_dataD.y, src_step, map_dataD.x + src_offset)]);
 290         else
 291             EXTRAPOLATE(map_dataD, d);
 292
 293         WT dst_data = a * (WT)(1 - u.x) * (WT)(1 - u.y) +
 294                       b * (WT)(u.x)     * (WT)(1 - u.y) +
 295                       c * (WT)(1 - u.x) * (WT)(u.y) +
 296                       d * (WT)(u.x)     * (WT)(u.y);
 297         dst[dstIdx] = convertToT(dst_data);
 298     }
 299 }
 300
 301 __kernel void remap_32FC2(__global T const * restrict  src, __global T * dst,
 302         __global float2 * map1,
 303         int src_offset, int dst_offset, int map1_offset,
 304         int src_step, int dst_step, int map1_step,
 305         int src_cols, int src_rows, int dst_cols, int dst_rows, T nVal)
 306 {
 307     int x = get_global_id(0);
 308     int y = get_global_id(1);
 309
 310     if (x < dst_cols && y < dst_rows)
 311     {
 312         int dstIdx = mad24(y, dst_step, x + dst_offset);
 313         int map1Idx = mad24(y, map1_step, x + map1_offset);
 314
 315         float2 map_data = map1[map1Idx];
 316         int2 map_dataA = convert_int2_sat_rtn(map_data);
 317         int2 map_dataB = (int2)(map_dataA.x + 1, map_dataA.y);
 318         int2 map_dataC = (int2)(map_dataA.x, map_dataA.y + 1);
 319         int2 map_dataD = (int2)(map_dataA.x + 1, map_dataA.y + 1);
 320
 321         float2 _u = map_data - convert_float2(map_dataA);
 322         WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)INTER_TAB_SIZE)) / (WT2)INTER_TAB_SIZE;
 323         WT scalar = convertToWT(nVal);
 324         WT a = scalar, b = scalar, c = scalar, d = scalar;
 325
 326         if (!NEED_EXTRAPOLATION(map_dataA.x, map_dataA.y))
 327             a = convertToWT(src[mad24(map_dataA.y, src_step, map_dataA.x + src_offset)]);
 328         else
 329             EXTRAPOLATE(map_dataA, a);
 330
 331         if (!NEED_EXTRAPOLATION(map_dataB.x, map_dataB.y))
 332             b = convertToWT(src[mad24(map_dataB.y, src_step, map_dataB.x + src_offset)]);
 333         else
 334             EXTRAPOLATE(map_dataB, b);
 335
 336         if (!NEED_EXTRAPOLATION(map_dataC.x, map_dataC.y))
 337             c = convertToWT(src[mad24(map_dataC.y, src_step, map_dataC.x + src_offset)]);
 338         else
 339             EXTRAPOLATE(map_dataC, c);
 340
 341         if (!NEED_EXTRAPOLATION(map_dataD.x, map_dataD.y))
 342             d = convertToWT(src[mad24(map_dataD.y, src_step, map_dataD.x + src_offset)]);
 343         else
 344             EXTRAPOLATE(map_dataD, d);
 345
 346         WT dst_data = a * (WT)(1 - u.x) * (WT)(1 - u.y) +
 347                       b * (WT)(u.x)     * (WT)(1 - u.y) +
 348                       c * (WT)(1 - u.x) * (WT)(u.y) +
 349                       d * (WT)(u.x)     * (WT)(u.y);
 350         dst[dstIdx] = convertToT(dst_data);
 351     }
 352 }
 353
 354 #endif