modules/superres/src/opencl/superres_btvl1.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, Multicoreware, Inc., 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 //    Jin Ma jin@multicorewareinc.com
  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 __kernel void buildMotionMapsKernel(__global float* forwardMotionX,
  47                                     __global float* forwardMotionY,
  48                                     __global float* backwardMotionX,
  49                                     __global float* backwardMotionY,
  50                                     __global float* forwardMapX,
  51                                     __global float* forwardMapY,
  52                                     __global float* backwardMapX,
  53                                     __global float* backwardMapY,
  54                                     int forwardMotionX_row,
  55                                     int forwardMotionX_col,
  56                                     int forwardMotionX_step,
  57                                     int forwardMotionY_step,
  58                                     int backwardMotionX_step,
  59                                     int backwardMotionY_step,
  60                                     int forwardMapX_step,
  61                                     int forwardMapY_step,
  62                                     int backwardMapX_step,
  63                                     int backwardMapY_step
  64                                    )
  65 {
  66     int x = get_global_id(0);
  67     int y = get_global_id(1);
  68
  69     if(x < forwardMotionX_col && y < forwardMotionX_row)
  70     {
  71         float fx = forwardMotionX[y * forwardMotionX_step + x];
  72         float fy = forwardMotionY[y * forwardMotionY_step + x];
  73
  74         float bx = backwardMotionX[y * backwardMotionX_step + x];
  75         float by = backwardMotionY[y * backwardMotionY_step + x];
  76
  77         forwardMapX[y * forwardMapX_step + x] = x + bx;
  78         forwardMapY[y * forwardMapY_step + x] = y + by;
  79
  80         backwardMapX[y * backwardMapX_step + x] = x + fx;
  81         backwardMapY[y * backwardMapY_step + x] = y + fy;
  82     }
  83 }
  84
  85 __kernel void upscaleKernel(__global float* src,
  86                             __global float* dst,
  87                             int src_step,
  88                             int dst_step,
  89                             int src_row,
  90                             int src_col,
  91                             int scale,
  92                             int channels
  93                            )
  94 {
  95     int x = get_global_id(0);
  96     int y = get_global_id(1);
  97
  98     if(x < src_col && y < src_row)
  99     {
 100         if(channels == 1)
 101         {
 102             dst[y * scale * dst_step + x * scale] = src[y * src_step + x];
 103         }
 104         else
 105         {
 106             vstore4(vload4(0, src + y * channels * src_step + 4 * x), 0, dst + y * channels * scale * dst_step + 4 * x * scale);
 107         }
 108     }
 109 }
 110
 111
 112 float diffSign(float a, float b)
 113 {
 114     return a > b ? 1.0f : a < b ? -1.0f : 0.0f;
 115 }
 116
 117 float4 diffSign4(float4 a, float4 b)
 118 {
 119     float4 pos;
 120     pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;
 121     pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;
 122     pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;
 123     pos.w = 0.0f;
 124     return pos;
 125 }
 126
 127 __kernel void diffSignKernel(__global float* src1,
 128                              __global float* src2,
 129                              __global float* dst,
 130                              int src1_row,
 131                              int src1_col,
 132                              int dst_step,
 133                              int src1_step,
 134                              int src2_step)
 135 {
 136     int x = get_global_id(0);
 137     int y = get_global_id(1);
 138
 139     if(x < src1_col && y < src1_row)
 140     {
 141         dst[y * dst_step + x] = diffSign(src1[y * src1_step + x], src2[y * src2_step + x]);
 142     }
 143 }
 144
 145 __kernel void calcBtvRegularizationKernel(__global float* src,
 146         __global float* dst,
 147         int src_step,
 148         int dst_step,
 149         int src_row,
 150         int src_col,
 151         int ksize,
 152         int channels,
 153         __constant float* c_btvRegWeights
 154                                          )
 155 {
 156     int x = get_global_id(0) + ksize;
 157     int y = get_global_id(1) + ksize;
 158
 159     if ((y < src_row - ksize) && (x < src_col - ksize))
 160     {
 161         if(channels == 1)
 162         {
 163             const float srcVal = src[y * src_step + x];
 164             float dstVal = 0.0f;
 165
 166             for (int m = 0, count = 0; m <= ksize; ++m)
 167             {
 168                 for (int l = ksize; l + m >= 0; --l, ++count)
 169                 {
 170                     dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src[(y + m) * src_step + (x + l)]) - diffSign(src[(y - m) * src_step + (x - l)], srcVal));
 171                 }
 172             }
 173             dst[y * dst_step + x] = dstVal;
 174         }
 175         else
 176         {
 177             float4 srcVal = vload4(0, src + y * src_step + 4 * x);
 178             float4 dstVal = 0.f;
 179
 180             for (int m = 0, count = 0; m <= ksize; ++m)
 181             {
 182                 for (int l = ksize; l + m >= 0; --l, ++count)
 183                 {
 184                     float4 src1;
 185                     src1.x = src[(y + m) * src_step + 4 * (x + l) + 0];
 186                     src1.y = src[(y + m) * src_step + 4 * (x + l) + 1];
 187                     src1.z = src[(y + m) * src_step + 4 * (x + l) + 2];
 188                     src1.w = src[(y + m) * src_step + 4 * (x + l) + 3];
 189
 190                     float4 src2;
 191                     src2.x = src[(y - m) * src_step + 4 * (x - l) + 0];
 192                     src2.y = src[(y - m) * src_step + 4 * (x - l) + 1];
 193                     src2.z = src[(y - m) * src_step + 4 * (x - l) + 2];
 194                     src2.w = src[(y - m) * src_step + 4 * (x - l) + 3];
 195
 196                     dstVal = dstVal + c_btvRegWeights[count] * (diffSign4(srcVal, src1) - diffSign4(src2, srcVal));
 197                 }
 198             }
 199             vstore4(dstVal, 0, dst + y * dst_step + 4 * x);
 200         }
 201     }
 202 }