modules/gpu/src/cuda/remap.cu

   1 /*M///////////////////////////////////////////////////////////////////////////////////////\r
   2 //\r
   3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
   4 //\r
   5 //  By downloading, copying, installing or using the software you agree to this license.\r
   6 //  If you do not agree to this license, do not download, install,\r
   7 //  copy or use the software.\r
   8 //\r
   9 //\r
  10 //                           License Agreement\r
  11 //                For Open Source Computer Vision Library\r
  12 //\r
  13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
  14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
  15 // Third party copyrights are property of their respective owners.\r
  16 //\r
  17 // Redistribution and use in source and binary forms, with or without modification,\r
  18 // are permitted provided that the following conditions are met:\r
  19 //\r
  20 //   * Redistribution's of source code must retain the above copyright notice,\r
  21 //     this list of conditions and the following disclaimer.\r
  22 //\r
  23 //   * Redistribution's in binary form must reproduce the above copyright notice,\r
  24 //     this list of conditions and the following disclaimer in the documentation\r
  25 //     and/or other materials provided with the distribution.\r
  26 //\r
  27 //   * The name of the copyright holders may not be used to endorse or promote products\r
  28 //     derived from this software without specific prior written permission.\r
  29 //\r
  30 // This software is provided by the copyright holders and contributors "as is" and\r
  31 // any express or implied warranties, including, but not limited to, the implied\r
  32 // warranties of merchantability and fitness for a particular purpose are disclaimed.\r
  33 // In no event shall the Intel Corporation or contributors be liable for any direct,\r
  34 // indirect, incidental, special, exemplary, or consequential damages\r
  35 // (including, but not limited to, procurement of substitute goods or services;\r
  36 // loss of use, data, or profits; or business interruption) however caused\r
  37 // and on any theory of liability, whether in contract, strict liability,\r
  38 // or tort (including negligence or otherwise) arising in any way out of\r
  39 // the use of this software, even if advised of the possibility of such damage.\r
  40 //\r
  41 //M*/\r
  42 \r
  43 #include "internal_shared.hpp"\r
  44 #include "opencv2/gpu/device/border_interpolate.hpp"\r
  45 #include "opencv2/gpu/device/vec_traits.hpp"\r
  46 #include "opencv2/gpu/device/vec_math.hpp"\r
  47 #include "opencv2/gpu/device/saturate_cast.hpp"\r
  48 #include "opencv2/gpu/device/filters.hpp"\r
  49 \r
  50 BEGIN_OPENCV_DEVICE_NAMESPACE\r
  51 \r
  52 namespace imgproc {\r
  53     \r
  54 template <typename Ptr2D, typename T> __global__ void remap(const Ptr2D src, const PtrStepf mapx, const PtrStepf mapy, DevMem2D_<T> dst)\r
  55 {\r
  56     const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
  57     const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
  58 \r
  59     if (x < dst.cols && y < dst.rows)\r
  60     {\r
  61         const float xcoo = mapx.ptr(y)[x];\r
  62         const float ycoo = mapy.ptr(y)[x];\r
  63 \r
  64         dst.ptr(y)[x] = saturate_cast<T>(src(ycoo, xcoo));\r
  65     }\r
  66 }\r
  67 \r
  68 template <template <typename> class Filter, template <typename> class B, typename T> struct RemapDispatcherStream\r
  69 {\r
  70     static void call(const DevMem2D_<T>& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_<T>& dst, \r
  71         const float* borderValue, cudaStream_t stream, int)\r
  72     {\r
  73         typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type; \r
  74         \r
  75         dim3 block(32, 8);\r
  76         dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));\r
  77 \r
  78         B<work_type> brd(src.rows, src.cols, VecTraits<work_type>::make(borderValue));\r
  79         BorderReader< PtrStep<T>, B<work_type> > brdSrc(src, brd);\r
  80         Filter< BorderReader< PtrStep<T>, B<work_type> > > filter_src(brdSrc);\r
  81 \r
  82         remap<<<grid, block, 0, stream>>>(filter_src, mapx, mapy, dst);\r
  83         cudaSafeCall( cudaGetLastError() );\r
  84     }\r
  85 };\r
  86 \r
  87 template <template <typename> class Filter, template <typename> class B, typename T> struct RemapDispatcherNonStream\r
  88 {\r
  89     static void call(const DevMem2D_<T>& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_<T>& dst, const float* borderValue, int)\r
  90     {\r
  91         typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type; \r
  92         \r
  93         dim3 block(32, 8);\r
  94         dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));\r
  95 \r
  96         B<work_type> brd(src.rows, src.cols, VecTraits<work_type>::make(borderValue));\r
  97         BorderReader< PtrStep<T>, B<work_type> > brdSrc(src, brd);\r
  98         Filter< BorderReader< PtrStep<T>, B<work_type> > > filter_src(brdSrc);\r
  99 \r
 100         remap<<<grid, block>>>(filter_src, mapx, mapy, dst);\r
 101         cudaSafeCall( cudaGetLastError() );\r
 102 \r
 103         cudaSafeCall( cudaDeviceSynchronize() );\r
 104     }\r
 105 };\r
 106 \r
 107 #define OPENCV_GPU_IMPLEMENT_REMAP_TEX(type) \\r
 108     texture< type , cudaTextureType2D> tex_remap_ ## type (0, cudaFilterModePoint, cudaAddressModeClamp); \\r
 109     struct tex_remap_ ## type ## _reader \\r
 110     { \\r
 111         typedef type elem_type; \\r
 112         typedef int index_type; \\r
 113         __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const \\r
 114         { \\r
 115             return tex2D(tex_remap_ ## type , x, y); \\r
 116         } \\r
 117     }; \\r
 118     template <template <typename> class Filter, template <typename> class B> struct RemapDispatcherNonStream<Filter, B, type> \\r
 119     { \\r
 120         static void call(const DevMem2D_< type >& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_< type >& dst, const float* borderValue, int cc) \\r
 121         { \\r
 122             typedef typename TypeVec<float, VecTraits< type >::cn>::vec_type work_type; \\r
 123             dim3 block(32, cc >= 20 ? 8 : 4); \\r
 124             dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \\r
 125             bindTexture(&tex_remap_ ## type , src); \\r
 126             tex_remap_ ## type ##_reader texSrc; \\r
 127             B<work_type> brd(src.rows, src.cols, VecTraits<work_type>::make(borderValue)); \\r
 128             BorderReader< tex_remap_ ## type ##_reader, B<work_type> > brdSrc(texSrc, brd); \\r
 129             Filter< BorderReader< tex_remap_ ## type ##_reader, B<work_type> > > filter_src(brdSrc); \\r
 130             remap<<<grid, block>>>(filter_src, mapx, mapy, dst); \\r
 131             cudaSafeCall( cudaGetLastError() ); \\r
 132             cudaSafeCall( cudaDeviceSynchronize() ); \\r
 133         } \\r
 134     }; \\r
 135     template <template <typename> class Filter> struct RemapDispatcherNonStream<Filter, BrdReplicate, type> \\r
 136     { \\r
 137         static void call(const DevMem2D_< type >& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_< type >& dst, const float*, int) \\r
 138         { \\r
 139             dim3 block(32, 8); \\r
 140             dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \\r
 141             bindTexture(&tex_remap_ ## type , src); \\r
 142             tex_remap_ ## type ##_reader texSrc; \\r
 143             Filter< tex_remap_ ## type ##_reader > filter_src(texSrc); \\r
 144             remap<<<grid, block>>>(filter_src, mapx, mapy, dst); \\r
 145             cudaSafeCall( cudaGetLastError() ); \\r
 146             cudaSafeCall( cudaDeviceSynchronize() ); \\r
 147         } \\r
 148     };\r
 149     \r
 150 OPENCV_GPU_IMPLEMENT_REMAP_TEX(uchar)\r
 151 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(uchar2)\r
 152 OPENCV_GPU_IMPLEMENT_REMAP_TEX(uchar4)\r
 153 \r
 154 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(schar)\r
 155 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(char2)\r
 156 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(char4)\r
 157 \r
 158 OPENCV_GPU_IMPLEMENT_REMAP_TEX(ushort)\r
 159 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(ushort2)\r
 160 OPENCV_GPU_IMPLEMENT_REMAP_TEX(ushort4)\r
 161 \r
 162 OPENCV_GPU_IMPLEMENT_REMAP_TEX(short)\r
 163 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(short2)\r
 164 OPENCV_GPU_IMPLEMENT_REMAP_TEX(short4)\r
 165 \r
 166 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(int)\r
 167 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(int2)\r
 168 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(int4)\r
 169 \r
 170 OPENCV_GPU_IMPLEMENT_REMAP_TEX(float)\r
 171 //OPENCV_GPU_IMPLEMENT_REMAP_TEX(float2)\r
 172 OPENCV_GPU_IMPLEMENT_REMAP_TEX(float4)\r
 173 \r
 174 #undef OPENCV_GPU_IMPLEMENT_REMAP_TEX\r
 175 \r
 176 template <template <typename> class Filter, template <typename> class B, typename T> struct RemapDispatcher\r
 177 { \r
 178     static void call(const DevMem2D_<T>& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_<T>& dst, \r
 179         const float* borderValue, cudaStream_t stream, int cc)\r
 180     {\r
 181         if (stream == 0)\r
 182             RemapDispatcherNonStream<Filter, B, T>::call(src, mapx, mapy, dst, borderValue, cc);\r
 183         else\r
 184             RemapDispatcherStream<Filter, B, T>::call(src, mapx, mapy, dst, borderValue, stream, cc);\r
 185     }\r
 186 };\r
 187 \r
 188 template <typename T> void remap_gpu(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, \r
 189     int borderMode, const float* borderValue, cudaStream_t stream, int cc)\r
 190 {\r
 191     typedef void (*caller_t)(const DevMem2D_<T>& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2D_<T>& dst, \r
 192         const float* borderValue, cudaStream_t stream, int cc);\r
 193 \r
 194     static const caller_t callers[3][5] = \r
 195     {\r
 196         { \r
 197             RemapDispatcher<PointFilter, BrdReflect101, T>::call, \r
 198             RemapDispatcher<PointFilter, BrdReplicate, T>::call, \r
 199             RemapDispatcher<PointFilter, BrdConstant, T>::call, \r
 200             RemapDispatcher<PointFilter, BrdReflect, T>::call, \r
 201             RemapDispatcher<PointFilter, BrdWrap, T>::call \r
 202         },\r
 203         { \r
 204             RemapDispatcher<LinearFilter, BrdReflect101, T>::call, \r
 205             RemapDispatcher<LinearFilter, BrdReplicate, T>::call, \r
 206             RemapDispatcher<LinearFilter, BrdConstant, T>::call, \r
 207             RemapDispatcher<LinearFilter, BrdReflect, T>::call, \r
 208             RemapDispatcher<LinearFilter, BrdWrap, T>::call \r
 209         },\r
 210         { \r
 211             RemapDispatcher<CubicFilter, BrdReflect101, T>::call, \r
 212             RemapDispatcher<CubicFilter, BrdReplicate, T>::call, \r
 213             RemapDispatcher<CubicFilter, BrdConstant, T>::call, \r
 214             RemapDispatcher<CubicFilter, BrdReflect, T>::call, \r
 215             RemapDispatcher<CubicFilter, BrdWrap, T>::call \r
 216         }\r
 217     };\r
 218 \r
 219     callers[interpolation][borderMode](static_cast< DevMem2D_<T> >(src), xmap, ymap, static_cast< DevMem2D_<T> >(dst), borderValue, stream, cc);\r
 220 }\r
 221 \r
 222 template void remap_gpu<uchar >(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 223 //template void remap_gpu<uchar2>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 224 template void remap_gpu<uchar3>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 225 template void remap_gpu<uchar4>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 226 \r
 227 //template void remap_gpu<schar>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 228 //template void remap_gpu<char2>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 229 //template void remap_gpu<char3>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 230 //template void remap_gpu<char4>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 231 \r
 232 template void remap_gpu<ushort >(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 233 //template void remap_gpu<ushort2>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 234 template void remap_gpu<ushort3>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 235 template void remap_gpu<ushort4>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 236 \r
 237 template void remap_gpu<short >(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 238 //template void remap_gpu<short2>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 239 template void remap_gpu<short3>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 240 template void remap_gpu<short4>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 241 \r
 242 //template void remap_gpu<int >(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 243 //template void remap_gpu<int2>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 244 //template void remap_gpu<int3>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 245 //template void remap_gpu<int4>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 246 \r
 247 template void remap_gpu<float >(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 248 //template void remap_gpu<float2>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 249 template void remap_gpu<float3>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 250 template void remap_gpu<float4>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);\r
 251 \r
 252 } // namespace imgproc\r
 253 \r
 254 END_OPENCV_DEVICE_NAMESPACE\r