modules/gpu/src/imgproc.cpp

   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 "precomp.hpp"\r
  44 \r
  45 using namespace cv;\r
  46 using namespace cv::gpu;\r
  47 \r
  48 #if !defined (HAVE_CUDA)\r
  49 \r
  50 void cv::gpu::remap(const GpuMat&, GpuMat&, const GpuMat&, const GpuMat&, int, int, const Scalar&, Stream&){ throw_nogpu(); }\r
  51 void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria) { throw_nogpu(); }\r
  52 void cv::gpu::meanShiftProc(const GpuMat&, GpuMat&, GpuMat&, int, int, TermCriteria) { throw_nogpu(); }\r
  53 void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, Stream&) { throw_nogpu(); }\r
  54 void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, Stream&) { throw_nogpu(); }\r
  55 void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int, Stream&) { throw_nogpu(); }\r
  56 void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const Scalar&, Stream&) { throw_nogpu(); }\r
  57 void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int, Stream&) { throw_nogpu(); }\r
  58 void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int, Stream&) { throw_nogpu(); }\r
  59 void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  60 void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  61 void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  62 void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int, Stream&) { throw_nogpu(); }\r
  63 void cv::gpu::integral(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  64 void cv::gpu::integralBuffered(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  65 void cv::gpu::integral(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  66 void cv::gpu::sqrIntegral(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  67 void cv::gpu::columnSum(const GpuMat&, GpuMat&) { throw_nogpu(); }\r
  68 void cv::gpu::rectStdDev(const GpuMat&, const GpuMat&, GpuMat&, const Rect&, Stream&) { throw_nogpu(); }\r
  69 void cv::gpu::evenLevels(GpuMat&, int, int, int) { throw_nogpu(); }\r
  70 void cv::gpu::histEven(const GpuMat&, GpuMat&, int, int, int, Stream&) { throw_nogpu(); }\r
  71 void cv::gpu::histEven(const GpuMat&, GpuMat&, GpuMat&, int, int, int, Stream&) { throw_nogpu(); }\r
  72 void cv::gpu::histEven(const GpuMat&, GpuMat*, int*, int*, int*, Stream&) { throw_nogpu(); }\r
  73 void cv::gpu::histEven(const GpuMat&, GpuMat*, GpuMat&, int*, int*, int*, Stream&) { throw_nogpu(); }\r
  74 void cv::gpu::histRange(const GpuMat&, GpuMat&, const GpuMat&, Stream&) { throw_nogpu(); }\r
  75 void cv::gpu::histRange(const GpuMat&, GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  76 void cv::gpu::histRange(const GpuMat&, GpuMat*, const GpuMat*, Stream&) { throw_nogpu(); }\r
  77 void cv::gpu::histRange(const GpuMat&, GpuMat*, const GpuMat*, GpuMat&, Stream&) { throw_nogpu(); }\r
  78 void cv::gpu::calcHist(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  79 void cv::gpu::calcHist(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  80 void cv::gpu::equalizeHist(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  81 void cv::gpu::equalizeHist(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  82 void cv::gpu::equalizeHist(const GpuMat&, GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  83 void cv::gpu::cornerHarris(const GpuMat&, GpuMat&, int, int, double, int) { throw_nogpu(); }\r
  84 void cv::gpu::cornerHarris(const GpuMat&, GpuMat&, GpuMat&, GpuMat&, int, int, double, int) { throw_nogpu(); }\r
  85 void cv::gpu::cornerMinEigenVal(const GpuMat&, GpuMat&, int, int, int) { throw_nogpu(); }\r
  86 void cv::gpu::cornerMinEigenVal(const GpuMat&, GpuMat&, GpuMat&, GpuMat&, int, int, int) { throw_nogpu(); }\r
  87 void cv::gpu::mulSpectrums(const GpuMat&, const GpuMat&, GpuMat&, int, bool) { throw_nogpu(); }\r
  88 void cv::gpu::mulAndScaleSpectrums(const GpuMat&, const GpuMat&, GpuMat&, int, float, bool) { throw_nogpu(); }\r
  89 void cv::gpu::dft(const GpuMat&, GpuMat&, Size, int) { throw_nogpu(); }\r
  90 void cv::gpu::ConvolveBuf::create(Size, Size) { throw_nogpu(); }\r
  91 void cv::gpu::convolve(const GpuMat&, const GpuMat&, GpuMat&, bool) { throw_nogpu(); }\r
  92 void cv::gpu::convolve(const GpuMat&, const GpuMat&, GpuMat&, bool, ConvolveBuf&) { throw_nogpu(); }\r
  93 void cv::gpu::pyrDown(const GpuMat&, GpuMat&, int, Stream&) { throw_nogpu(); }\r
  94 void cv::gpu::pyrUp(const GpuMat&, GpuMat&, int, Stream&) { throw_nogpu(); }\r
  95 void cv::gpu::Canny(const GpuMat&, GpuMat&, double, double, int, bool) { throw_nogpu(); }\r
  96 void cv::gpu::Canny(const GpuMat&, CannyBuf&, GpuMat&, double, double, int, bool) { throw_nogpu(); }\r
  97 void cv::gpu::Canny(const GpuMat&, const GpuMat&, GpuMat&, double, double, bool) { throw_nogpu(); }\r
  98 void cv::gpu::Canny(const GpuMat&, const GpuMat&, CannyBuf&, GpuMat&, double, double, bool) { throw_nogpu(); }\r
  99 cv::gpu::CannyBuf::CannyBuf(const GpuMat&, const GpuMat&) { throw_nogpu(); }\r
 100 void cv::gpu::CannyBuf::create(const Size&, int) { throw_nogpu(); }\r
 101 void cv::gpu::CannyBuf::release() { throw_nogpu(); }\r
 102 \r
 103 #else /* !defined (HAVE_CUDA) */\r
 104 \r
 105 ////////////////////////////////////////////////////////////////////////\r
 106 // remap\r
 107 \r
 108 namespace cv { namespace gpu {  namespace imgproc\r
 109 {\r
 110     template <typename T> void remap_gpu(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2D& dst, \r
 111                                          int interpolation, int borderMode, const float* borderValue, cudaStream_t stream);\r
 112 }}}\r
 113 \r
 114 void cv::gpu::remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap, int interpolation, int borderMode, const Scalar& borderValue, Stream& stream)\r
 115 {\r
 116     using namespace cv::gpu::imgproc;\r
 117 \r
 118     typedef void (*caller_t)(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2D& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream);\r
 119     static const caller_t callers[6][4] = \r
 120     {\r
 121         {remap_gpu<uchar>, 0/*remap_gpu<uchar2>*/, remap_gpu<uchar3>, remap_gpu<uchar4>},\r
 122         {0/*remap_gpu<schar>*/, 0/*remap_gpu<char2>*/, 0/*remap_gpu<char3>*/, 0/*remap_gpu<char4>*/},\r
 123         {remap_gpu<ushort>, 0/*remap_gpu<ushort2>*/, remap_gpu<ushort3>, remap_gpu<ushort4>},\r
 124         {remap_gpu<short>, 0/*remap_gpu<short2>*/, remap_gpu<short3>, remap_gpu<short4>},\r
 125         {0/*remap_gpu<int>*/, 0/*remap_gpu<int2>*/, 0/*remap_gpu<int3>*/, 0/*remap_gpu<int4>*/},\r
 126         {remap_gpu<float>, 0/*remap_gpu<float2>*/, remap_gpu<float3>, remap_gpu<float4>}\r
 127     };\r
 128 \r
 129     CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);\r
 130     CV_Assert(xmap.type() == CV_32F && ymap.type() == CV_32F && xmap.size() == ymap.size());\r
 131 \r
 132     caller_t func = callers[src.depth()][src.channels() - 1];\r
 133     CV_Assert(func != 0);\r
 134 \r
 135     CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
 136 \r
 137     CV_Assert(borderMode == BORDER_REFLECT101 || borderMode == BORDER_REPLICATE || borderMode == BORDER_CONSTANT || borderMode == BORDER_REFLECT || borderMode == BORDER_WRAP);\r
 138     int gpuBorderType;\r
 139     CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType));\r
 140 \r
 141     dst.create(xmap.size(), src.type());\r
 142     \r
 143     Scalar_<float> borderValueFloat;\r
 144     borderValueFloat = borderValue;\r
 145 \r
 146     func(src, xmap, ymap, dst, interpolation, gpuBorderType, borderValueFloat.val, StreamAccessor::getStream(stream));\r
 147 }\r
 148 \r
 149 ////////////////////////////////////////////////////////////////////////\r
 150 // meanShiftFiltering_GPU\r
 151 \r
 152 namespace cv { namespace gpu {  namespace imgproc\r
 153 {\r
 154     extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps);\r
 155 }}}\r
 156 \r
 157 void cv::gpu::meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria)\r
 158 {\r
 159     if( src.empty() )\r
 160         CV_Error( CV_StsBadArg, "The input image is empty" );\r
 161 \r
 162     if( src.depth() != CV_8U || src.channels() != 4 )\r
 163         CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" );\r
 164 \r
 165     dst.create( src.size(), CV_8UC4 );\r
 166 \r
 167     if( !(criteria.type & TermCriteria::MAX_ITER) )\r
 168         criteria.maxCount = 5;\r
 169 \r
 170     int maxIter = std::min(std::max(criteria.maxCount, 1), 100);\r
 171 \r
 172     float eps;\r
 173     if( !(criteria.type & TermCriteria::EPS) )\r
 174         eps = 1.f;\r
 175     eps = (float)std::max(criteria.epsilon, 0.0);\r
 176 \r
 177     imgproc::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps);\r
 178 }\r
 179 \r
 180 ////////////////////////////////////////////////////////////////////////\r
 181 // meanShiftProc_GPU\r
 182 \r
 183 namespace cv { namespace gpu {  namespace imgproc\r
 184 {\r
 185     extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps);\r
 186 }}}\r
 187 \r
 188 void cv::gpu::meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int sp, int sr, TermCriteria criteria)\r
 189 {\r
 190     if( src.empty() )\r
 191         CV_Error( CV_StsBadArg, "The input image is empty" );\r
 192 \r
 193     if( src.depth() != CV_8U || src.channels() != 4 )\r
 194         CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" );\r
 195 \r
 196     dstr.create( src.size(), CV_8UC4 );\r
 197     dstsp.create( src.size(), CV_16SC2 );\r
 198 \r
 199     if( !(criteria.type & TermCriteria::MAX_ITER) )\r
 200         criteria.maxCount = 5;\r
 201 \r
 202     int maxIter = std::min(std::max(criteria.maxCount, 1), 100);\r
 203 \r
 204     float eps;\r
 205     if( !(criteria.type & TermCriteria::EPS) )\r
 206         eps = 1.f;\r
 207     eps = (float)std::max(criteria.epsilon, 0.0);\r
 208 \r
 209     imgproc::meanShiftProc_gpu(src, dstr, dstsp, sp, sr, maxIter, eps);\r
 210 }\r
 211 \r
 212 ////////////////////////////////////////////////////////////////////////\r
 213 // drawColorDisp\r
 214 \r
 215 namespace cv { namespace gpu {  namespace imgproc\r
 216 {\r
 217     void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);\r
 218     void drawColorDisp_gpu(const DevMem2D_<short>& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);\r
 219 }}}\r
 220 \r
 221 namespace\r
 222 {\r
 223     template <typename T>\r
 224     void drawColorDisp_caller(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream)\r
 225     {\r
 226         dst.create(src.size(), CV_8UC4);\r
 227 \r
 228         imgproc::drawColorDisp_gpu((DevMem2D_<T>)src, dst, ndisp, stream);\r
 229     }\r
 230 \r
 231     typedef void (*drawColorDisp_caller_t)(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream);\r
 232 \r
 233     const drawColorDisp_caller_t drawColorDisp_callers[] = {drawColorDisp_caller<unsigned char>, 0, 0, drawColorDisp_caller<short>, 0, 0, 0, 0};\r
 234 }\r
 235 \r
 236 void cv::gpu::drawColorDisp(const GpuMat& src, GpuMat& dst, int ndisp, Stream& stream)\r
 237 {\r
 238     CV_Assert(src.type() == CV_8U || src.type() == CV_16S);\r
 239 \r
 240     drawColorDisp_callers[src.type()](src, dst, ndisp, StreamAccessor::getStream(stream));\r
 241 }\r
 242 \r
 243 ////////////////////////////////////////////////////////////////////////\r
 244 // reprojectImageTo3D\r
 245 \r
 246 namespace cv { namespace gpu {  namespace imgproc\r
 247 {\r
 248     void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
 249     void reprojectImageTo3D_gpu(const DevMem2D_<short>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
 250 }}}\r
 251 \r
 252 namespace\r
 253 {\r
 254     template <typename T>\r
 255     void reprojectImageTo3D_caller(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream)\r
 256     {\r
 257         xyzw.create(disp.rows, disp.cols, CV_32FC4);\r
 258         imgproc::reprojectImageTo3D_gpu((DevMem2D_<T>)disp, xyzw, Q.ptr<float>(), stream);\r
 259     }\r
 260 \r
 261     typedef void (*reprojectImageTo3D_caller_t)(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream);\r
 262 \r
 263     const reprojectImageTo3D_caller_t reprojectImageTo3D_callers[] = {reprojectImageTo3D_caller<unsigned char>, 0, 0, reprojectImageTo3D_caller<short>, 0, 0, 0, 0};\r
 264 }\r
 265 \r
 266 void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, Stream& stream)\r
 267 {\r
 268     CV_Assert((disp.type() == CV_8U || disp.type() == CV_16S) && Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4);\r
 269 \r
 270     reprojectImageTo3D_callers[disp.type()](disp, xyzw, Q, StreamAccessor::getStream(stream));\r
 271 }\r
 272 \r
 273 ////////////////////////////////////////////////////////////////////////\r
 274 // resize\r
 275 \r
 276 namespace cv { namespace gpu {  namespace imgproc\r
 277 {\r
 278     template <typename T> void resize_gpu(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);\r
 279 }}}\r
 280 \r
 281 void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, double fy, int interpolation, Stream& s)\r
 282 {\r
 283     CV_Assert( src.depth() <= CV_32F && src.channels() <= 4 );\r
 284     CV_Assert( interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC );\r
 285     CV_Assert( !(dsize == Size()) || (fx > 0 && fy > 0) );\r
 286 \r
 287     if( dsize == Size() )\r
 288     {\r
 289         dsize = Size(saturate_cast<int>(src.cols * fx), saturate_cast<int>(src.rows * fy));\r
 290     }\r
 291     else\r
 292     {\r
 293         fx = (double)dsize.width / src.cols;\r
 294         fy = (double)dsize.height / src.rows;\r
 295     }\r
 296 \r
 297     dst.create(dsize, src.type());\r
 298 \r
 299     if (dsize == src.size())\r
 300     {\r
 301         if (s)\r
 302             s.enqueueCopy(src, dst);\r
 303         else\r
 304             src.copyTo(dst);\r
 305         return;\r
 306     }\r
 307 \r
 308     cudaStream_t stream = StreamAccessor::getStream(s);\r
 309 \r
 310     if ((src.type() == CV_8UC1 || src.type() == CV_8UC4) && (interpolation == INTER_NEAREST || interpolation == INTER_LINEAR))\r
 311     {\r
 312         static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC, 0, NPPI_INTER_LANCZOS};\r
 313 \r
 314         NppiSize srcsz;\r
 315         srcsz.width  = src.cols;\r
 316         srcsz.height = src.rows;\r
 317         NppiRect srcrect;\r
 318         srcrect.x = srcrect.y = 0;\r
 319         srcrect.width  = src.cols;\r
 320         srcrect.height = src.rows;\r
 321         NppiSize dstsz;\r
 322         dstsz.width  = dst.cols;\r
 323         dstsz.height = dst.rows;\r
 324 \r
 325         NppStreamHandler h(stream);\r
 326 \r
 327         if (src.type() == CV_8UC1)\r
 328         {\r
 329             nppSafeCall( nppiResize_8u_C1R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcrect,\r
 330                 dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, fx, fy, npp_inter[interpolation]) );\r
 331         }\r
 332         else\r
 333         {\r
 334             nppSafeCall( nppiResize_8u_C4R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcrect,\r
 335                 dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, fx, fy, npp_inter[interpolation]) );\r
 336         }\r
 337 \r
 338         if (stream == 0)\r
 339             cudaSafeCall( cudaDeviceSynchronize() );\r
 340     }\r
 341     else\r
 342     {\r
 343         using namespace cv::gpu::imgproc;\r
 344 \r
 345         typedef void (*caller_t)(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);\r
 346         static const caller_t callers[6][4] = \r
 347         {\r
 348             {resize_gpu<uchar>, 0/*resize_gpu<uchar2>*/, resize_gpu<uchar3>, resize_gpu<uchar4>},\r
 349             {0/*resize_gpu<schar>*/, 0/*resize_gpu<char2>*/, 0/*resize_gpu<char3>*/, 0/*resize_gpu<char4>*/},\r
 350             {resize_gpu<ushort>, 0/*resize_gpu<ushort2>*/, resize_gpu<ushort3>, resize_gpu<ushort4>},\r
 351             {resize_gpu<short>, 0/*resize_gpu<short2>*/, resize_gpu<short3>, resize_gpu<short4>},\r
 352             {0/*resize_gpu<int>*/, 0/*resize_gpu<int2>*/, 0/*resize_gpu<int3>*/, 0/*resize_gpu<int4>*/},\r
 353             {resize_gpu<float>, 0/*resize_gpu<float2>*/, resize_gpu<float3>, resize_gpu<float4>}\r
 354         };\r
 355 \r
 356         callers[src.depth()][src.channels() - 1](src, static_cast<float>(fx), static_cast<float>(fy), dst, interpolation, stream);\r
 357     }\r
 358 }\r
 359 \r
 360 ////////////////////////////////////////////////////////////////////////\r
 361 // copyMakeBorder\r
 362 \r
 363 void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value, Stream& s)\r
 364 {\r
 365     CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1);\r
 366 \r
 367     dst.create(src.rows + top + bottom, src.cols + left + right, src.type());\r
 368 \r
 369     NppiSize srcsz;\r
 370     srcsz.width  = src.cols;\r
 371     srcsz.height = src.rows;\r
 372     NppiSize dstsz;\r
 373     dstsz.width  = dst.cols;\r
 374     dstsz.height = dst.rows;\r
 375 \r
 376     cudaStream_t stream = StreamAccessor::getStream(s);\r
 377 \r
 378     NppStreamHandler h(stream);\r
 379 \r
 380     switch (src.type())\r
 381     {\r
 382     case CV_8UC1:\r
 383         {\r
 384             Npp8u nVal = static_cast<Npp8u>(value[0]);\r
 385             nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step), srcsz,\r
 386                 dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
 387             break;\r
 388         }\r
 389     case CV_8UC4:\r
 390         {\r
 391             Npp8u nVal[] = {static_cast<Npp8u>(value[0]), static_cast<Npp8u>(value[1]), static_cast<Npp8u>(value[2]), static_cast<Npp8u>(value[3])};\r
 392             nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr<Npp8u>(), static_cast<int>(src.step), srcsz,\r
 393                 dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
 394             break;\r
 395         }\r
 396     case CV_32SC1:\r
 397         {\r
 398             Npp32s nVal = static_cast<Npp32s>(value[0]);\r
 399             nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), srcsz,\r
 400                 dst.ptr<Npp32s>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
 401             break;\r
 402         }\r
 403     case CV_32FC1:\r
 404         {\r
 405             Npp32f val = static_cast<Npp32f>(value[0]);\r
 406             Npp32s nVal = *(reinterpret_cast<Npp32s*>(&val));\r
 407             nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), srcsz,\r
 408                 dst.ptr<Npp32s>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
 409             break;\r
 410         }\r
 411     default:\r
 412         CV_Assert(!"Unsupported source type");\r
 413     }\r
 414 \r
 415     if (stream == 0)\r
 416         cudaSafeCall( cudaDeviceSynchronize() );\r
 417 }\r
 418 \r
 419 ////////////////////////////////////////////////////////////////////////\r
 420 // warp\r
 421 \r
 422 namespace\r
 423 {\r
 424     typedef NppStatus (*npp_warp_8u_t)(const Npp8u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp8u* pDst,\r
 425                                        int dstStep, NppiRect dstRoi, const double coeffs[][3],\r
 426                                        int interpolation);\r
 427     typedef NppStatus (*npp_warp_16u_t)(const Npp16u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp16u* pDst,\r
 428                                        int dstStep, NppiRect dstRoi, const double coeffs[][3],\r
 429                                        int interpolation);\r
 430     typedef NppStatus (*npp_warp_32s_t)(const Npp32s* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32s* pDst,\r
 431                                        int dstStep, NppiRect dstRoi, const double coeffs[][3],\r
 432                                        int interpolation);\r
 433     typedef NppStatus (*npp_warp_32f_t)(const Npp32f* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32f* pDst,\r
 434                                        int dstStep, NppiRect dstRoi, const double coeffs[][3],\r
 435                                        int interpolation);\r
 436 \r
 437     void nppWarpCaller(const GpuMat& src, GpuMat& dst, double coeffs[][3], const Size& dsize, int flags,\r
 438                        npp_warp_8u_t npp_warp_8u[][2], npp_warp_16u_t npp_warp_16u[][2],\r
 439                        npp_warp_32s_t npp_warp_32s[][2], npp_warp_32f_t npp_warp_32f[][2], cudaStream_t stream)\r
 440     {\r
 441         static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};\r
 442 \r
 443         int interpolation = flags & INTER_MAX;\r
 444 \r
 445         CV_Assert((src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S || src.depth() == CV_32F) && src.channels() != 2);\r
 446         CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
 447 \r
 448         dst.create(dsize, src.type());\r
 449 \r
 450         NppiSize srcsz;\r
 451         srcsz.height = src.rows;\r
 452         srcsz.width = src.cols;\r
 453         NppiRect srcroi;\r
 454         srcroi.x = srcroi.y = 0;\r
 455         srcroi.height = src.rows;\r
 456         srcroi.width = src.cols;\r
 457         NppiRect dstroi;\r
 458         dstroi.x = dstroi.y = 0;\r
 459         dstroi.height = dst.rows;\r
 460         dstroi.width = dst.cols;\r
 461 \r
 462         int warpInd = (flags & WARP_INVERSE_MAP) >> 4;\r
 463 \r
 464         NppStreamHandler h(stream);\r
 465 \r
 466         switch (src.depth())\r
 467         {\r
 468         case CV_8U:\r
 469             nppSafeCall( npp_warp_8u[src.channels()][warpInd](src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcroi,\r
 470                 dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstroi, coeffs, npp_inter[interpolation]) );\r
 471             break;\r
 472         case CV_16U:\r
 473             nppSafeCall( npp_warp_16u[src.channels()][warpInd](src.ptr<Npp16u>(), srcsz, static_cast<int>(src.step), srcroi,\r
 474                 dst.ptr<Npp16u>(), static_cast<int>(dst.step), dstroi, coeffs, npp_inter[interpolation]) );\r
 475             break;\r
 476         case CV_32S:\r
 477             nppSafeCall( npp_warp_32s[src.channels()][warpInd](src.ptr<Npp32s>(), srcsz, static_cast<int>(src.step), srcroi,\r
 478                 dst.ptr<Npp32s>(), static_cast<int>(dst.step), dstroi, coeffs, npp_inter[interpolation]) );\r
 479             break;\r
 480         case CV_32F:\r
 481             nppSafeCall( npp_warp_32f[src.channels()][warpInd](src.ptr<Npp32f>(), srcsz, static_cast<int>(src.step), srcroi,\r
 482                 dst.ptr<Npp32f>(), static_cast<int>(dst.step), dstroi, coeffs, npp_inter[interpolation]) );\r
 483             break;\r
 484         default:\r
 485             CV_Assert(!"Unsupported source type");\r
 486         }\r
 487 \r
 488         if (stream == 0)\r
 489             cudaSafeCall( cudaDeviceSynchronize() );\r
 490     }\r
 491 }\r
 492 \r
 493 void cv::gpu::warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags, Stream& s)\r
 494 {\r
 495     static npp_warp_8u_t npp_warpAffine_8u[][2] =\r
 496         {\r
 497             {0, 0},\r
 498             {nppiWarpAffine_8u_C1R, nppiWarpAffineBack_8u_C1R},\r
 499             {0, 0},\r
 500             {nppiWarpAffine_8u_C3R, nppiWarpAffineBack_8u_C3R},\r
 501             {nppiWarpAffine_8u_C4R, nppiWarpAffineBack_8u_C4R}\r
 502         };\r
 503     static npp_warp_16u_t npp_warpAffine_16u[][2] =\r
 504         {\r
 505             {0, 0},\r
 506             {nppiWarpAffine_16u_C1R, nppiWarpAffineBack_16u_C1R},\r
 507             {0, 0},\r
 508             {nppiWarpAffine_16u_C3R, nppiWarpAffineBack_16u_C3R},\r
 509             {nppiWarpAffine_16u_C4R, nppiWarpAffineBack_16u_C4R}\r
 510         };\r
 511     static npp_warp_32s_t npp_warpAffine_32s[][2] =\r
 512         {\r
 513             {0, 0},\r
 514             {nppiWarpAffine_32s_C1R, nppiWarpAffineBack_32s_C1R},\r
 515             {0, 0},\r
 516             {nppiWarpAffine_32s_C3R, nppiWarpAffineBack_32s_C3R},\r
 517             {nppiWarpAffine_32s_C4R, nppiWarpAffineBack_32s_C4R}\r
 518         };\r
 519     static npp_warp_32f_t npp_warpAffine_32f[][2] =\r
 520         {\r
 521             {0, 0},\r
 522             {nppiWarpAffine_32f_C1R, nppiWarpAffineBack_32f_C1R},\r
 523             {0, 0},\r
 524             {nppiWarpAffine_32f_C3R, nppiWarpAffineBack_32f_C3R},\r
 525             {nppiWarpAffine_32f_C4R, nppiWarpAffineBack_32f_C4R}\r
 526         };\r
 527 \r
 528     CV_Assert(M.rows == 2 && M.cols == 3);\r
 529 \r
 530     double coeffs[2][3];\r
 531     Mat coeffsMat(2, 3, CV_64F, (void*)coeffs);\r
 532     M.convertTo(coeffsMat, coeffsMat.type());\r
 533 \r
 534     nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpAffine_8u, npp_warpAffine_16u, npp_warpAffine_32s, npp_warpAffine_32f, StreamAccessor::getStream(s));\r
 535 }\r
 536 \r
 537 void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags, Stream& s)\r
 538 {\r
 539     static npp_warp_8u_t npp_warpPerspective_8u[][2] =\r
 540         {\r
 541             {0, 0},\r
 542             {nppiWarpPerspective_8u_C1R, nppiWarpPerspectiveBack_8u_C1R},\r
 543             {0, 0},\r
 544             {nppiWarpPerspective_8u_C3R, nppiWarpPerspectiveBack_8u_C3R},\r
 545             {nppiWarpPerspective_8u_C4R, nppiWarpPerspectiveBack_8u_C4R}\r
 546         };\r
 547     static npp_warp_16u_t npp_warpPerspective_16u[][2] =\r
 548         {\r
 549             {0, 0},\r
 550             {nppiWarpPerspective_16u_C1R, nppiWarpPerspectiveBack_16u_C1R},\r
 551             {0, 0},\r
 552             {nppiWarpPerspective_16u_C3R, nppiWarpPerspectiveBack_16u_C3R},\r
 553             {nppiWarpPerspective_16u_C4R, nppiWarpPerspectiveBack_16u_C4R}\r
 554         };\r
 555     static npp_warp_32s_t npp_warpPerspective_32s[][2] =\r
 556         {\r
 557             {0, 0},\r
 558             {nppiWarpPerspective_32s_C1R, nppiWarpPerspectiveBack_32s_C1R},\r
 559             {0, 0},\r
 560             {nppiWarpPerspective_32s_C3R, nppiWarpPerspectiveBack_32s_C3R},\r
 561             {nppiWarpPerspective_32s_C4R, nppiWarpPerspectiveBack_32s_C4R}\r
 562         };\r
 563     static npp_warp_32f_t npp_warpPerspective_32f[][2] =\r
 564         {\r
 565             {0, 0},\r
 566             {nppiWarpPerspective_32f_C1R, nppiWarpPerspectiveBack_32f_C1R},\r
 567             {0, 0},\r
 568             {nppiWarpPerspective_32f_C3R, nppiWarpPerspectiveBack_32f_C3R},\r
 569             {nppiWarpPerspective_32f_C4R, nppiWarpPerspectiveBack_32f_C4R}\r
 570         };\r
 571 \r
 572     CV_Assert(M.rows == 3 && M.cols == 3);\r
 573 \r
 574     double coeffs[3][3];\r
 575     Mat coeffsMat(3, 3, CV_64F, (void*)coeffs);\r
 576     M.convertTo(coeffsMat, coeffsMat.type());\r
 577 \r
 578     nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpPerspective_8u, npp_warpPerspective_16u, npp_warpPerspective_32s, npp_warpPerspective_32f, StreamAccessor::getStream(s));\r
 579 }\r
 580 \r
 581 //////////////////////////////////////////////////////////////////////////////\r
 582 // buildWarpPlaneMaps\r
 583 \r
 584 namespace cv { namespace gpu { namespace imgproc\r
 585 {\r
 586     void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
 587                             const float k_rinv[9], const float r_kinv[9], float scale,\r
 588                             cudaStream_t stream);\r
 589 }}}\r
 590 \r
 591 void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale,\r
 592                                  GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
 593 {\r
 594     CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F);\r
 595     CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F);\r
 596 \r
 597     Mat K_Rinv = K * R.t();\r
 598     Mat R_Kinv = R * K.inv();\r
 599     CV_Assert(K_Rinv.isContinuous());\r
 600     CV_Assert(R_Kinv.isContinuous());\r
 601 \r
 602     map_x.create(dst_roi.size(), CV_32F);\r
 603     map_y.create(dst_roi.size(), CV_32F);\r
 604     imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(),\r
 605                                 scale, StreamAccessor::getStream(stream));\r
 606 }\r
 607 \r
 608 //////////////////////////////////////////////////////////////////////////////\r
 609 // buildWarpCylyndricalMaps\r
 610 \r
 611 namespace cv { namespace gpu { namespace imgproc\r
 612 {\r
 613     void buildWarpCylindricalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
 614                                   const float k_rinv[9], const float r_kinv[9], float scale,\r
 615                                   cudaStream_t stream);\r
 616 }}}\r
 617 \r
 618 void cv::gpu::buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale,\r
 619                                        GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
 620 {\r
 621     CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F);\r
 622     CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F);\r
 623 \r
 624     Mat K_Rinv = K * R.t();\r
 625     Mat R_Kinv = R * K.inv();\r
 626     CV_Assert(K_Rinv.isContinuous());\r
 627     CV_Assert(R_Kinv.isContinuous());\r
 628 \r
 629     map_x.create(dst_roi.size(), CV_32F);\r
 630     map_y.create(dst_roi.size(), CV_32F);\r
 631     imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(),\r
 632                                       scale, StreamAccessor::getStream(stream));\r
 633 }\r
 634 \r
 635 \r
 636 //////////////////////////////////////////////////////////////////////////////\r
 637 // buildWarpSphericalMaps\r
 638 \r
 639 namespace cv { namespace gpu { namespace imgproc\r
 640 {\r
 641     void buildWarpSphericalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
 642                                 const float k_rinv[9], const float r_kinv[9], float scale,\r
 643                                 cudaStream_t stream);\r
 644 }}}\r
 645 \r
 646 void cv::gpu::buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale,\r
 647                                      GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
 648 {\r
 649     CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F);\r
 650     CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F);\r
 651 \r
 652     Mat K_Rinv = K * R.t();\r
 653     Mat R_Kinv = R * K.inv();\r
 654     CV_Assert(K_Rinv.isContinuous());\r
 655     CV_Assert(R_Kinv.isContinuous());\r
 656 \r
 657     map_x.create(dst_roi.size(), CV_32F);\r
 658     map_y.create(dst_roi.size(), CV_32F);\r
 659     imgproc::buildWarpSphericalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(),\r
 660                                     scale, StreamAccessor::getStream(stream));\r
 661 }\r
 662 \r
 663 ////////////////////////////////////////////////////////////////////////\r
 664 // rotate\r
 665 \r
 666 void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, Stream& s)\r
 667 {\r
 668     static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};\r
 669 \r
 670     CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
 671     CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
 672 \r
 673     dst.create(dsize, src.type());\r
 674 \r
 675     NppiSize srcsz;\r
 676     srcsz.height = src.rows;\r
 677     srcsz.width = src.cols;\r
 678     NppiRect srcroi;\r
 679     srcroi.x = srcroi.y = 0;\r
 680     srcroi.height = src.rows;\r
 681     srcroi.width = src.cols;\r
 682     NppiRect dstroi;\r
 683     dstroi.x = dstroi.y = 0;\r
 684     dstroi.height = dst.rows;\r
 685     dstroi.width = dst.cols;\r
 686 \r
 687     cudaStream_t stream = StreamAccessor::getStream(s);\r
 688 \r
 689     NppStreamHandler h(stream);\r
 690 \r
 691     if (src.type() == CV_8UC1)\r
 692     {\r
 693         nppSafeCall( nppiRotate_8u_C1R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcroi,\r
 694             dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstroi, angle, xShift, yShift, npp_inter[interpolation]) );\r
 695     }\r
 696     else\r
 697     {\r
 698         nppSafeCall( nppiRotate_8u_C4R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcroi,\r
 699             dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstroi, angle, xShift, yShift, npp_inter[interpolation]) );\r
 700     }\r
 701 \r
 702     if (stream == 0)\r
 703         cudaSafeCall( cudaDeviceSynchronize() );\r
 704 }\r
 705 \r
 706 ////////////////////////////////////////////////////////////////////////\r
 707 // integral\r
 708 \r
 709 void cv::gpu::integral(const GpuMat& src, GpuMat& sum, Stream& s)\r
 710 {\r
 711     GpuMat buffer;\r
 712     integralBuffered(src, sum, buffer, s);\r
 713 }\r
 714 \r
 715 void cv::gpu::integralBuffered(const GpuMat& src, GpuMat& sum, GpuMat& buffer, Stream& s)\r
 716 {\r
 717     CV_Assert(src.type() == CV_8UC1);\r
 718 \r
 719     sum.create(src.rows + 1, src.cols + 1, CV_32S);\r
 720     \r
 721     NcvSize32u roiSize;\r
 722     roiSize.width = src.cols;\r
 723     roiSize.height = src.rows;\r
 724 \r
 725         cudaDeviceProp prop;\r
 726         cudaSafeCall( cudaGetDeviceProperties(&prop, cv::gpu::getDevice()) );\r
 727 \r
 728     Ncv32u bufSize;\r
 729     nppSafeCall( nppiStIntegralGetSize_8u32u(roiSize, &bufSize, prop) );\r
 730     ensureSizeIsEnough(1, bufSize, CV_8UC1, buffer);\r
 731 \r
 732     cudaStream_t stream = StreamAccessor::getStream(s);\r
 733 \r
 734     NppStStreamHandler h(stream);\r
 735 \r
 736     nppSafeCall( nppiStIntegral_8u32u_C1R(const_cast<Ncv8u*>(src.ptr<Ncv8u>()), static_cast<int>(src.step), \r
 737         sum.ptr<Ncv32u>(), static_cast<int>(sum.step), roiSize, buffer.ptr<Ncv8u>(), bufSize, prop) );\r
 738 \r
 739     if (stream == 0)\r
 740         cudaSafeCall( cudaDeviceSynchronize() );\r
 741 }\r
 742 \r
 743 void cv::gpu::integral(const GpuMat& src, GpuMat& sum, GpuMat& sqsum, Stream& s)\r
 744 {\r
 745     CV_Assert(src.type() == CV_8UC1);\r
 746 \r
 747     int width = src.cols + 1, height = src.rows + 1;\r
 748 \r
 749     sum.create(height, width, CV_32S);\r
 750     sqsum.create(height, width, CV_32F);\r
 751 \r
 752     NppiSize sz;\r
 753     sz.width = src.cols;\r
 754     sz.height = src.rows;\r
 755 \r
 756     cudaStream_t stream = StreamAccessor::getStream(s);\r
 757 \r
 758     NppStreamHandler h(stream);\r
 759 \r
 760     nppSafeCall( nppiSqrIntegral_8u32s32f_C1R(const_cast<Npp8u*>(src.ptr<Npp8u>()), static_cast<int>(src.step), \r
 761         sum.ptr<Npp32s>(), static_cast<int>(sum.step), sqsum.ptr<Npp32f>(), static_cast<int>(sqsum.step), sz, 0, 0.0f, height) );\r
 762 \r
 763     if (stream == 0)\r
 764         cudaSafeCall( cudaDeviceSynchronize() );\r
 765 }\r
 766 \r
 767 //////////////////////////////////////////////////////////////////////////////\r
 768 // sqrIntegral\r
 769 \r
 770 void cv::gpu::sqrIntegral(const GpuMat& src, GpuMat& sqsum, Stream& s)\r
 771 {\r
 772     CV_Assert(src.type() == CV_8U);\r
 773 \r
 774     NcvSize32u roiSize;\r
 775     roiSize.width = src.cols;\r
 776     roiSize.height = src.rows;\r
 777 \r
 778         cudaDeviceProp prop;\r
 779         cudaSafeCall( cudaGetDeviceProperties(&prop, cv::gpu::getDevice()) );\r
 780 \r
 781     Ncv32u bufSize;\r
 782     nppSafeCall(nppiStSqrIntegralGetSize_8u64u(roiSize, &bufSize, prop));       \r
 783     GpuMat buf(1, bufSize, CV_8U);\r
 784 \r
 785     cudaStream_t stream = StreamAccessor::getStream(s);\r
 786 \r
 787     NppStStreamHandler h(stream);\r
 788 \r
 789     sqsum.create(src.rows + 1, src.cols + 1, CV_64F);\r
 790     nppSafeCall(nppiStSqrIntegral_8u64u_C1R(const_cast<Ncv8u*>(src.ptr<Ncv8u>(0)), static_cast<int>(src.step), \r
 791             sqsum.ptr<Ncv64u>(0), static_cast<int>(sqsum.step), roiSize, buf.ptr<Ncv8u>(0), bufSize, prop));\r
 792 \r
 793     if (stream == 0)\r
 794         cudaSafeCall( cudaDeviceSynchronize() );\r
 795 }\r
 796 \r
 797 //////////////////////////////////////////////////////////////////////////////\r
 798 // columnSum\r
 799 \r
 800 namespace cv { namespace gpu { namespace imgproc\r
 801 {\r
 802     void columnSum_32F(const DevMem2D src, const DevMem2D dst);\r
 803 }}}\r
 804 \r
 805 void cv::gpu::columnSum(const GpuMat& src, GpuMat& dst)\r
 806 {\r
 807     CV_Assert(src.type() == CV_32F);\r
 808 \r
 809     dst.create(src.size(), CV_32F);\r
 810     imgproc::columnSum_32F(src, dst);\r
 811 }\r
 812 \r
 813 void cv::gpu::rectStdDev(const GpuMat& src, const GpuMat& sqr, GpuMat& dst, const Rect& rect, Stream& s)\r
 814 {\r
 815     CV_Assert(src.type() == CV_32SC1 && sqr.type() == CV_32FC1);\r
 816 \r
 817     dst.create(src.size(), CV_32FC1);\r
 818 \r
 819     NppiSize sz;\r
 820     sz.width = src.cols;\r
 821     sz.height = src.rows;\r
 822 \r
 823     NppiRect nppRect;\r
 824     nppRect.height = rect.height;\r
 825     nppRect.width = rect.width;\r
 826     nppRect.x = rect.x;\r
 827     nppRect.y = rect.y;\r
 828 \r
 829     cudaStream_t stream = StreamAccessor::getStream(s);\r
 830 \r
 831     NppStreamHandler h(stream);\r
 832 \r
 833     nppSafeCall( nppiRectStdDev_32s32f_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), sqr.ptr<Npp32f>(), static_cast<int>(sqr.step),\r
 834                 dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, nppRect) );\r
 835 \r
 836     if (stream == 0)\r
 837         cudaSafeCall( cudaDeviceSynchronize() );\r
 838 }\r
 839 \r
 840 \r
 841 ////////////////////////////////////////////////////////////////////////\r
 842 // Histogram\r
 843 \r
 844 namespace\r
 845 {\r
 846     template<int n> struct NPPTypeTraits;\r
 847     template<> struct NPPTypeTraits<CV_8U>  { typedef Npp8u npp_type; };\r
 848     template<> struct NPPTypeTraits<CV_16U> { typedef Npp16u npp_type; };\r
 849     template<> struct NPPTypeTraits<CV_16S> { typedef Npp16s npp_type; };\r
 850     template<> struct NPPTypeTraits<CV_32F> { typedef Npp32f npp_type; };\r
 851 \r
 852     typedef NppStatus (*get_buf_size_c1_t)(NppiSize oSizeROI, int nLevels, int* hpBufferSize);\r
 853     typedef NppStatus (*get_buf_size_c4_t)(NppiSize oSizeROI, int nLevels[], int* hpBufferSize);\r
 854 \r
 855     template<int SDEPTH> struct NppHistogramEvenFuncC1\r
 856     {\r
 857         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;\r
 858 \r
 859         typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s * pHist,\r
 860                     int nLevels, Npp32s nLowerLevel, Npp32s nUpperLevel, Npp8u * pBuffer);\r
 861     };\r
 862     template<int SDEPTH> struct NppHistogramEvenFuncC4\r
 863     {\r
 864         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;\r
 865 \r
 866         typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI,\r
 867             Npp32s * pHist[4], int nLevels[4], Npp32s nLowerLevel[4], Npp32s nUpperLevel[4], Npp8u * pBuffer);\r
 868     };\r
 869 \r
 870     template<int SDEPTH, typename NppHistogramEvenFuncC1<SDEPTH>::func_ptr func, get_buf_size_c1_t get_buf_size>\r
 871     struct NppHistogramEvenC1\r
 872     {\r
 873         typedef typename NppHistogramEvenFuncC1<SDEPTH>::src_t src_t;\r
 874 \r
 875         static void hist(const GpuMat& src, GpuMat& hist, GpuMat& buffer, int histSize, int lowerLevel, int upperLevel, cudaStream_t stream)\r
 876         {\r
 877             int levels = histSize + 1;\r
 878             hist.create(1, histSize, CV_32S);\r
 879 \r
 880             NppiSize sz;\r
 881             sz.width = src.cols;\r
 882             sz.height = src.rows;\r
 883 \r
 884             int buf_size;\r
 885             get_buf_size(sz, levels, &buf_size);\r
 886 \r
 887             ensureSizeIsEnough(1, buf_size, CV_8U, buffer);\r
 888 \r
 889             NppStreamHandler h(stream);\r
 890 \r
 891             nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), sz, hist.ptr<Npp32s>(), levels,\r
 892                 lowerLevel, upperLevel, buffer.ptr<Npp8u>()) );\r
 893 \r
 894             if (stream == 0)\r
 895                 cudaSafeCall( cudaDeviceSynchronize() );\r
 896         }\r
 897     };\r
 898     template<int SDEPTH, typename NppHistogramEvenFuncC4<SDEPTH>::func_ptr func, get_buf_size_c4_t get_buf_size>\r
 899     struct NppHistogramEvenC4\r
 900     {\r
 901         typedef typename NppHistogramEvenFuncC4<SDEPTH>::src_t src_t;\r
 902 \r
 903         static void hist(const GpuMat& src, GpuMat hist[4], GpuMat& buffer, int histSize[4], int lowerLevel[4], int upperLevel[4], cudaStream_t stream)\r
 904         {\r
 905             int levels[] = {histSize[0] + 1, histSize[1] + 1, histSize[2] + 1, histSize[3] + 1};\r
 906             hist[0].create(1, histSize[0], CV_32S);\r
 907             hist[1].create(1, histSize[1], CV_32S);\r
 908             hist[2].create(1, histSize[2], CV_32S);\r
 909             hist[3].create(1, histSize[3], CV_32S);\r
 910 \r
 911             NppiSize sz;\r
 912             sz.width = src.cols;\r
 913             sz.height = src.rows;\r
 914 \r
 915             Npp32s* pHist[] = {hist[0].ptr<Npp32s>(), hist[1].ptr<Npp32s>(), hist[2].ptr<Npp32s>(), hist[3].ptr<Npp32s>()};\r
 916 \r
 917             int buf_size;\r
 918             get_buf_size(sz, levels, &buf_size);\r
 919 \r
 920             ensureSizeIsEnough(1, buf_size, CV_8U, buffer);\r
 921 \r
 922             NppStreamHandler h(stream);\r
 923 \r
 924             nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), sz, pHist, levels, lowerLevel, upperLevel, buffer.ptr<Npp8u>()) );\r
 925 \r
 926             if (stream == 0)\r
 927                 cudaSafeCall( cudaDeviceSynchronize() );\r
 928         }\r
 929     };\r
 930 \r
 931     template<int SDEPTH> struct NppHistogramRangeFuncC1\r
 932     {\r
 933         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;\r
 934         typedef Npp32s level_t;\r
 935         enum {LEVEL_TYPE_CODE=CV_32SC1};\r
 936 \r
 937         typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s* pHist,\r
 938             const Npp32s* pLevels, int nLevels, Npp8u* pBuffer);\r
 939     };\r
 940     template<> struct NppHistogramRangeFuncC1<CV_32F>\r
 941     {\r
 942         typedef Npp32f src_t;\r
 943         typedef Npp32f level_t;\r
 944         enum {LEVEL_TYPE_CODE=CV_32FC1};\r
 945 \r
 946         typedef NppStatus (*func_ptr)(const Npp32f* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s* pHist,\r
 947             const Npp32f* pLevels, int nLevels, Npp8u* pBuffer);\r
 948     };\r
 949     template<int SDEPTH> struct NppHistogramRangeFuncC4\r
 950     {\r
 951         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;\r
 952         typedef Npp32s level_t;\r
 953         enum {LEVEL_TYPE_CODE=CV_32SC1};\r
 954 \r
 955         typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s* pHist[4],\r
 956             const Npp32s* pLevels[4], int nLevels[4], Npp8u* pBuffer);\r
 957     };\r
 958     template<> struct NppHistogramRangeFuncC4<CV_32F>\r
 959     {\r
 960         typedef Npp32f src_t;\r
 961         typedef Npp32f level_t;\r
 962         enum {LEVEL_TYPE_CODE=CV_32FC1};\r
 963 \r
 964         typedef NppStatus (*func_ptr)(const Npp32f* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s* pHist[4],\r
 965             const Npp32f* pLevels[4], int nLevels[4], Npp8u* pBuffer);\r
 966     };\r
 967 \r
 968     template<int SDEPTH, typename NppHistogramRangeFuncC1<SDEPTH>::func_ptr func, get_buf_size_c1_t get_buf_size>\r
 969     struct NppHistogramRangeC1\r
 970     {\r
 971         typedef typename NppHistogramRangeFuncC1<SDEPTH>::src_t src_t;\r
 972         typedef typename NppHistogramRangeFuncC1<SDEPTH>::level_t level_t;\r
 973         enum {LEVEL_TYPE_CODE=NppHistogramRangeFuncC1<SDEPTH>::LEVEL_TYPE_CODE};\r
 974 \r
 975         static void hist(const GpuMat& src, GpuMat& hist, const GpuMat& levels, GpuMat& buffer, cudaStream_t stream)\r
 976         {\r
 977             CV_Assert(levels.type() == LEVEL_TYPE_CODE && levels.rows == 1);\r
 978 \r
 979             hist.create(1, levels.cols - 1, CV_32S);\r
 980 \r
 981             NppiSize sz;\r
 982             sz.width = src.cols;\r
 983             sz.height = src.rows;\r
 984 \r
 985             int buf_size;\r
 986             get_buf_size(sz, levels.cols, &buf_size);\r
 987             \r
 988             ensureSizeIsEnough(1, buf_size, CV_8U, buffer);\r
 989 \r
 990             NppStreamHandler h(stream);\r
 991 \r
 992             nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), sz, hist.ptr<Npp32s>(), levels.ptr<level_t>(), levels.cols, buffer.ptr<Npp8u>()) );\r
 993 \r
 994             if (stream == 0)\r
 995                 cudaSafeCall( cudaDeviceSynchronize() );\r
 996         }\r
 997     };\r
 998     template<int SDEPTH, typename NppHistogramRangeFuncC4<SDEPTH>::func_ptr func, get_buf_size_c4_t get_buf_size>\r
 999     struct NppHistogramRangeC4\r
1000     {\r
1001         typedef typename NppHistogramRangeFuncC4<SDEPTH>::src_t src_t;\r
1002         typedef typename NppHistogramRangeFuncC1<SDEPTH>::level_t level_t;\r
1003         enum {LEVEL_TYPE_CODE=NppHistogramRangeFuncC1<SDEPTH>::LEVEL_TYPE_CODE};\r
1004 \r
1005         static void hist(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4], GpuMat& buffer, cudaStream_t stream)\r
1006         {\r
1007             CV_Assert(levels[0].type() == LEVEL_TYPE_CODE && levels[0].rows == 1);\r
1008             CV_Assert(levels[1].type() == LEVEL_TYPE_CODE && levels[1].rows == 1);\r
1009             CV_Assert(levels[2].type() == LEVEL_TYPE_CODE && levels[2].rows == 1);\r
1010             CV_Assert(levels[3].type() == LEVEL_TYPE_CODE && levels[3].rows == 1);\r
1011 \r
1012             hist[0].create(1, levels[0].cols - 1, CV_32S);\r
1013             hist[1].create(1, levels[1].cols - 1, CV_32S);\r
1014             hist[2].create(1, levels[2].cols - 1, CV_32S);\r
1015             hist[3].create(1, levels[3].cols - 1, CV_32S);\r
1016 \r
1017             Npp32s* pHist[] = {hist[0].ptr<Npp32s>(), hist[1].ptr<Npp32s>(), hist[2].ptr<Npp32s>(), hist[3].ptr<Npp32s>()};\r
1018             int nLevels[] = {levels[0].cols, levels[1].cols, levels[2].cols, levels[3].cols};\r
1019             const level_t* pLevels[] = {levels[0].ptr<level_t>(), levels[1].ptr<level_t>(), levels[2].ptr<level_t>(), levels[3].ptr<level_t>()};\r
1020 \r
1021             NppiSize sz;\r
1022             sz.width = src.cols;\r
1023             sz.height = src.rows;\r
1024 \r
1025             int buf_size;\r
1026             get_buf_size(sz, nLevels, &buf_size);\r
1027 \r
1028             ensureSizeIsEnough(1, buf_size, CV_8U, buffer);\r
1029 \r
1030             NppStreamHandler h(stream);\r
1031 \r
1032             nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), sz, pHist, pLevels, nLevels, buffer.ptr<Npp8u>()) );\r
1033 \r
1034             if (stream == 0)\r
1035                 cudaSafeCall( cudaDeviceSynchronize() );\r
1036         }\r
1037     };\r
1038 }\r
1039 \r
1040 void cv::gpu::evenLevels(GpuMat& levels, int nLevels, int lowerLevel, int upperLevel)\r
1041 {\r
1042     Mat host_levels(1, nLevels, CV_32SC1);\r
1043     nppSafeCall( nppiEvenLevelsHost_32s(host_levels.ptr<Npp32s>(), nLevels, lowerLevel, upperLevel) );\r
1044     levels.upload(host_levels);\r
1045 }\r
1046 \r
1047 void cv::gpu::histEven(const GpuMat& src, GpuMat& hist, int histSize, int lowerLevel, int upperLevel, Stream& stream)\r
1048 {\r
1049     GpuMat buf;\r
1050     histEven(src, hist, buf, histSize, lowerLevel, upperLevel, stream);\r
1051 }\r
1052 \r
1053 void cv::gpu::histEven(const GpuMat& src, GpuMat& hist, GpuMat& buf, int histSize, int lowerLevel, int upperLevel, Stream& stream)\r
1054 {\r
1055     CV_Assert(src.type() == CV_8UC1 || src.type() == CV_16UC1 || src.type() == CV_16SC1 );\r
1056 \r
1057     typedef void (*hist_t)(const GpuMat& src, GpuMat& hist, GpuMat& buf, int levels, int lowerLevel, int upperLevel, cudaStream_t stream);\r
1058     static const hist_t hist_callers[] =\r
1059     {\r
1060         NppHistogramEvenC1<CV_8U , nppiHistogramEven_8u_C1R , nppiHistogramEvenGetBufferSize_8u_C1R >::hist,\r
1061         0,\r
1062         NppHistogramEvenC1<CV_16U, nppiHistogramEven_16u_C1R, nppiHistogramEvenGetBufferSize_16u_C1R>::hist,\r
1063         NppHistogramEvenC1<CV_16S, nppiHistogramEven_16s_C1R, nppiHistogramEvenGetBufferSize_16s_C1R>::hist\r
1064     };\r
1065 \r
1066     hist_callers[src.depth()](src, hist, buf, histSize, lowerLevel, upperLevel, StreamAccessor::getStream(stream));\r
1067 }\r
1068 \r
1069 void cv::gpu::histEven(const GpuMat& src, GpuMat hist[4], int histSize[4], int lowerLevel[4], int upperLevel[4], Stream& stream)\r
1070 {\r
1071     GpuMat buf;\r
1072     histEven(src, hist, buf, histSize, lowerLevel, upperLevel, stream);\r
1073 }\r
1074 \r
1075 void cv::gpu::histEven(const GpuMat& src, GpuMat hist[4], GpuMat& buf, int histSize[4], int lowerLevel[4], int upperLevel[4], Stream& stream)\r
1076 {\r
1077     CV_Assert(src.type() == CV_8UC4 || src.type() == CV_16UC4 || src.type() == CV_16SC4 );\r
1078 \r
1079     typedef void (*hist_t)(const GpuMat& src, GpuMat hist[4], GpuMat& buf, int levels[4], int lowerLevel[4], int upperLevel[4], cudaStream_t stream);\r
1080     static const hist_t hist_callers[] =\r
1081     {\r
1082         NppHistogramEvenC4<CV_8U , nppiHistogramEven_8u_C4R , nppiHistogramEvenGetBufferSize_8u_C4R >::hist,\r
1083         0,\r
1084         NppHistogramEvenC4<CV_16U, nppiHistogramEven_16u_C4R, nppiHistogramEvenGetBufferSize_16u_C4R>::hist,\r
1085         NppHistogramEvenC4<CV_16S, nppiHistogramEven_16s_C4R, nppiHistogramEvenGetBufferSize_16s_C4R>::hist\r
1086     };\r
1087 \r
1088     hist_callers[src.depth()](src, hist, buf, histSize, lowerLevel, upperLevel, StreamAccessor::getStream(stream));\r
1089 }\r
1090 \r
1091 void cv::gpu::histRange(const GpuMat& src, GpuMat& hist, const GpuMat& levels, Stream& stream)\r
1092 {\r
1093     GpuMat buf;\r
1094     histRange(src, hist, levels, buf, stream);\r
1095 }\r
1096 \r
1097 \r
1098 void cv::gpu::histRange(const GpuMat& src, GpuMat& hist, const GpuMat& levels, GpuMat& buf, Stream& stream)\r
1099 {\r
1100     CV_Assert(src.type() == CV_8UC1 || src.type() == CV_16UC1 || src.type() == CV_16SC1 || src.type() == CV_32FC1);\r
1101 \r
1102     typedef void (*hist_t)(const GpuMat& src, GpuMat& hist, const GpuMat& levels, GpuMat& buf, cudaStream_t stream);\r
1103     static const hist_t hist_callers[] =\r
1104     {\r
1105         NppHistogramRangeC1<CV_8U , nppiHistogramRange_8u_C1R , nppiHistogramRangeGetBufferSize_8u_C1R >::hist,\r
1106         0,\r
1107         NppHistogramRangeC1<CV_16U, nppiHistogramRange_16u_C1R, nppiHistogramRangeGetBufferSize_16u_C1R>::hist,\r
1108         NppHistogramRangeC1<CV_16S, nppiHistogramRange_16s_C1R, nppiHistogramRangeGetBufferSize_16s_C1R>::hist,\r
1109         0,\r
1110         NppHistogramRangeC1<CV_32F, nppiHistogramRange_32f_C1R, nppiHistogramRangeGetBufferSize_32f_C1R>::hist\r
1111     };\r
1112 \r
1113     hist_callers[src.depth()](src, hist, levels, buf, StreamAccessor::getStream(stream));\r
1114 }\r
1115 \r
1116 void cv::gpu::histRange(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4], Stream& stream)\r
1117 {\r
1118     GpuMat buf;\r
1119     histRange(src, hist, levels, buf, stream);\r
1120 }\r
1121 \r
1122 void cv::gpu::histRange(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4], GpuMat& buf, Stream& stream)\r
1123 {\r
1124     CV_Assert(src.type() == CV_8UC4 || src.type() == CV_16UC4 || src.type() == CV_16SC4 || src.type() == CV_32FC4);\r
1125 \r
1126     typedef void (*hist_t)(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4], GpuMat& buf, cudaStream_t stream);\r
1127     static const hist_t hist_callers[] =\r
1128     {\r
1129         NppHistogramRangeC4<CV_8U , nppiHistogramRange_8u_C4R , nppiHistogramRangeGetBufferSize_8u_C4R >::hist,\r
1130         0,\r
1131         NppHistogramRangeC4<CV_16U, nppiHistogramRange_16u_C4R, nppiHistogramRangeGetBufferSize_16u_C4R>::hist,\r
1132         NppHistogramRangeC4<CV_16S, nppiHistogramRange_16s_C4R, nppiHistogramRangeGetBufferSize_16s_C4R>::hist,\r
1133         0,\r
1134         NppHistogramRangeC4<CV_32F, nppiHistogramRange_32f_C4R, nppiHistogramRangeGetBufferSize_32f_C4R>::hist\r
1135     };\r
1136 \r
1137     hist_callers[src.depth()](src, hist, levels, buf, StreamAccessor::getStream(stream));\r
1138 }\r
1139 \r
1140 namespace cv { namespace gpu { namespace histograms\r
1141 {\r
1142     void histogram256_gpu(DevMem2D src, int* hist, unsigned int* buf, cudaStream_t stream);\r
1143 \r
1144     const int PARTIAL_HISTOGRAM256_COUNT = 240;\r
1145     const int HISTOGRAM256_BIN_COUNT     = 256;\r
1146 }}}\r
1147 \r
1148 void cv::gpu::calcHist(const GpuMat& src, GpuMat& hist, Stream& stream)\r
1149 {\r
1150     GpuMat buf;\r
1151     calcHist(src, hist, buf, stream);\r
1152 }\r
1153 \r
1154 void cv::gpu::calcHist(const GpuMat& src, GpuMat& hist, GpuMat& buf, Stream& stream)\r
1155 {\r
1156     using namespace cv::gpu::histograms;\r
1157 \r
1158     CV_Assert(src.type() == CV_8UC1);\r
1159 \r
1160     hist.create(1, 256, CV_32SC1);\r
1161 \r
1162     ensureSizeIsEnough(1, PARTIAL_HISTOGRAM256_COUNT * HISTOGRAM256_BIN_COUNT, CV_32SC1, buf);\r
1163 \r
1164     histogram256_gpu(src, hist.ptr<int>(), buf.ptr<unsigned int>(), StreamAccessor::getStream(stream));\r
1165 }\r
1166 \r
1167 void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, Stream& stream)\r
1168 {\r
1169     GpuMat hist;\r
1170     GpuMat buf;\r
1171     equalizeHist(src, dst, hist, buf, stream);\r
1172 }\r
1173 \r
1174 void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, Stream& stream)\r
1175 {\r
1176     GpuMat buf;\r
1177     equalizeHist(src, dst, hist, buf, stream);\r
1178 }\r
1179 \r
1180 namespace cv { namespace gpu { namespace histograms\r
1181 {\r
1182     void equalizeHist_gpu(DevMem2D src, DevMem2D dst, const int* lut, cudaStream_t stream);\r
1183 }}}\r
1184 \r
1185 void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, GpuMat& buf, Stream& s)\r
1186 {\r
1187     using namespace cv::gpu::histograms;\r
1188 \r
1189     CV_Assert(src.type() == CV_8UC1);\r
1190 \r
1191     dst.create(src.size(), src.type());\r
1192 \r
1193     int intBufSize;\r
1194     nppSafeCall( nppsIntegralGetBufferSize_32s(256, &intBufSize) );\r
1195 \r
1196     int bufSize = static_cast<int>(std::max(256 * 240 * sizeof(int), intBufSize + 256 * sizeof(int)));\r
1197 \r
1198     ensureSizeIsEnough(1, bufSize, CV_8UC1, buf);\r
1199 \r
1200     GpuMat histBuf(1, 256 * 240, CV_32SC1, buf.ptr());\r
1201     GpuMat intBuf(1, intBufSize, CV_8UC1, buf.ptr());\r
1202     GpuMat lut(1, 256, CV_32S, buf.ptr() + intBufSize);\r
1203 \r
1204     calcHist(src, hist, histBuf, s);\r
1205 \r
1206     cudaStream_t stream = StreamAccessor::getStream(s);\r
1207 \r
1208     NppStreamHandler h(stream);\r
1209 \r
1210     nppSafeCall( nppsIntegral_32s(hist.ptr<Npp32s>(), lut.ptr<Npp32s>(), 256, intBuf.ptr<Npp8u>()) );\r
1211     \r
1212     if (stream == 0)\r
1213         cudaSafeCall( cudaDeviceSynchronize() );\r
1214 \r
1215     equalizeHist_gpu(src, dst, lut.ptr<int>(), stream);\r
1216 }\r
1217 \r
1218 ////////////////////////////////////////////////////////////////////////\r
1219 // cornerHarris & minEgenVal\r
1220 \r
1221 namespace cv { namespace gpu { namespace imgproc {\r
1222 \r
1223     void extractCovData_caller(const DevMem2Df Dx, const DevMem2Df Dy, PtrStepf dst);\r
1224     void cornerHarris_caller(const int block_size, const float k, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst, int border_type);\r
1225     void cornerMinEigenVal_caller(const int block_size, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst, int border_type);\r
1226 \r
1227 }}}\r
1228 \r
1229 namespace \r
1230 {\r
1231     template <typename T>\r
1232     void extractCovData(const GpuMat& src, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, int borderType)\r
1233     {   \r
1234         double scale = (double)(1 << ((ksize > 0 ? ksize : 3) - 1)) * blockSize;\r
1235         if (ksize < 0) \r
1236             scale *= 2.;\r
1237         if (src.depth() == CV_8U)\r
1238             scale *= 255.;\r
1239         scale = 1./scale;\r
1240 \r
1241         Dx.create(src.size(), CV_32F);\r
1242         Dy.create(src.size(), CV_32F);\r
1243 \r
1244         if (ksize > 0)\r
1245         {\r
1246             Sobel(src, Dx, CV_32F, 1, 0, ksize, scale, borderType);\r
1247             Sobel(src, Dy, CV_32F, 0, 1, ksize, scale, borderType);\r
1248         }\r
1249         else\r
1250         {\r
1251             Scharr(src, Dx, CV_32F, 1, 0, scale, borderType);\r
1252             Scharr(src, Dy, CV_32F, 0, 1, scale, borderType);\r
1253         }\r
1254     }\r
1255 \r
1256     void extractCovData(const GpuMat& src, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, int borderType)\r
1257     {\r
1258         switch (src.type())\r
1259         {\r
1260         case CV_8U:\r
1261             extractCovData<unsigned char>(src, Dx, Dy, blockSize, ksize, borderType);\r
1262             break;\r
1263         case CV_32F:\r
1264             extractCovData<float>(src, Dx, Dy, blockSize, ksize, borderType);\r
1265             break;\r
1266         default:\r
1267             CV_Error(CV_StsBadArg, "extractCovData: unsupported type of the source matrix");\r
1268         }\r
1269     }\r
1270 \r
1271 } // Anonymous namespace\r
1272 \r
1273 \r
1274 bool cv::gpu::tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType)\r
1275 {\r
1276     switch (cpuBorderType)\r
1277     {\r
1278     case cv::BORDER_REFLECT101:\r
1279         gpuBorderType = cv::gpu::BORDER_REFLECT101_GPU;\r
1280         return true;\r
1281     case cv::BORDER_REPLICATE:\r
1282         gpuBorderType = cv::gpu::BORDER_REPLICATE_GPU;\r
1283         return true;\r
1284     case cv::BORDER_CONSTANT:\r
1285         gpuBorderType = cv::gpu::BORDER_CONSTANT_GPU;\r
1286         return true;\r
1287     case cv::BORDER_REFLECT:\r
1288         gpuBorderType = cv::gpu::BORDER_REFLECT_GPU;\r
1289         return true;\r
1290     case cv::BORDER_WRAP:\r
1291         gpuBorderType = cv::gpu::BORDER_WRAP_GPU;\r
1292         return true;\r
1293     default:\r
1294         return false;\r
1295     };\r
1296     return false;\r
1297 }\r
1298 \r
1299 void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, double k, int borderType)\r
1300 {\r
1301     GpuMat Dx, Dy;\r
1302     cornerHarris(src, dst, Dx, Dy, blockSize, ksize, k, borderType);\r
1303 }\r
1304 \r
1305 void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, double k, int borderType)\r
1306 {\r
1307     CV_Assert(borderType == cv::BORDER_REFLECT101 ||\r
1308               borderType == cv::BORDER_REPLICATE);\r
1309 \r
1310     int gpuBorderType;\r
1311     CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
1312 \r
1313     extractCovData(src, Dx, Dy, blockSize, ksize, borderType);\r
1314     dst.create(src.size(), CV_32F);\r
1315     imgproc::cornerHarris_caller(blockSize, (float)k, Dx, Dy, dst, gpuBorderType);\r
1316 }\r
1317 \r
1318 void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, int borderType)\r
1319 {  \r
1320     GpuMat Dx, Dy;\r
1321     cornerMinEigenVal(src, dst, Dx, Dy, blockSize, ksize, borderType);\r
1322 }\r
1323 \r
1324 void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, int borderType)\r
1325 {  \r
1326     CV_Assert(borderType == cv::BORDER_REFLECT101 ||\r
1327               borderType == cv::BORDER_REPLICATE);\r
1328 \r
1329     int gpuBorderType;\r
1330     CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
1331 \r
1332     extractCovData(src, Dx, Dy, blockSize, ksize, borderType);    \r
1333     dst.create(src.size(), CV_32F);\r
1334     imgproc::cornerMinEigenVal_caller(blockSize, Dx, Dy, dst, gpuBorderType);\r
1335 }\r
1336 \r
1337 //////////////////////////////////////////////////////////////////////////////\r
1338 // mulSpectrums\r
1339 \r
1340 namespace cv { namespace gpu { namespace imgproc \r
1341 {\r
1342     void mulSpectrums(const PtrStep_<cufftComplex> a, const PtrStep_<cufftComplex> b, \r
1343                       DevMem2D_<cufftComplex> c);\r
1344 \r
1345     void mulSpectrums_CONJ(const PtrStep_<cufftComplex> a, const PtrStep_<cufftComplex> b, \r
1346                            DevMem2D_<cufftComplex> c);\r
1347 }}}\r
1348 \r
1349 \r
1350 void cv::gpu::mulSpectrums(const GpuMat& a, const GpuMat& b, GpuMat& c, \r
1351                            int flags, bool conjB) \r
1352 {\r
1353     typedef void (*Caller)(const PtrStep_<cufftComplex>, const PtrStep_<cufftComplex>, \r
1354                            DevMem2D_<cufftComplex>);\r
1355     static Caller callers[] = { imgproc::mulSpectrums, \r
1356                                 imgproc::mulSpectrums_CONJ };\r
1357 \r
1358     CV_Assert(a.type() == b.type() && a.type() == CV_32FC2);\r
1359     CV_Assert(a.size() == b.size());\r
1360 \r
1361     c.create(a.size(), CV_32FC2);\r
1362 \r
1363     Caller caller = callers[(int)conjB];\r
1364     caller(a, b, c);\r
1365 }\r
1366 \r
1367 //////////////////////////////////////////////////////////////////////////////\r
1368 // mulAndScaleSpectrums\r
1369 \r
1370 namespace cv { namespace gpu { namespace imgproc \r
1371 {\r
1372     void mulAndScaleSpectrums(const PtrStep_<cufftComplex> a, const PtrStep_<cufftComplex> b,\r
1373                              float scale, DevMem2D_<cufftComplex> c);\r
1374 \r
1375     void mulAndScaleSpectrums_CONJ(const PtrStep_<cufftComplex> a, const PtrStep_<cufftComplex> b,\r
1376                                   float scale, DevMem2D_<cufftComplex> c);\r
1377 }}}\r
1378 \r
1379 \r
1380 void cv::gpu::mulAndScaleSpectrums(const GpuMat& a, const GpuMat& b, GpuMat& c,\r
1381                                   int flags, float scale, bool conjB) \r
1382 {\r
1383     typedef void (*Caller)(const PtrStep_<cufftComplex>, const PtrStep_<cufftComplex>,\r
1384                            float scale, DevMem2D_<cufftComplex>);\r
1385     static Caller callers[] = { imgproc::mulAndScaleSpectrums, \r
1386                                 imgproc::mulAndScaleSpectrums_CONJ };\r
1387 \r
1388     CV_Assert(a.type() == b.type() && a.type() == CV_32FC2);\r
1389     CV_Assert(a.size() == b.size());\r
1390 \r
1391     c.create(a.size(), CV_32FC2);\r
1392 \r
1393     Caller caller = callers[(int)conjB];\r
1394     caller(a, b, scale, c);\r
1395 }\r
1396 \r
1397 //////////////////////////////////////////////////////////////////////////////\r
1398 // dft\r
1399 \r
1400 void cv::gpu::dft(const GpuMat& src, GpuMat& dst, Size dft_size, int flags)\r
1401 {\r
1402     CV_Assert(src.type() == CV_32F || src.type() == CV_32FC2);\r
1403 \r
1404     // We don't support unpacked output (in the case of real input)\r
1405     CV_Assert(!(flags & DFT_COMPLEX_OUTPUT));\r
1406 \r
1407     bool is_1d_input = (dft_size.height == 1) || (dft_size.width == 1);\r
1408     int is_row_dft = flags & DFT_ROWS;\r
1409     int is_scaled_dft = flags & DFT_SCALE;\r
1410     int is_inverse = flags & DFT_INVERSE;\r
1411     bool is_complex_input = src.channels() == 2;\r
1412     bool is_complex_output = !(flags & DFT_REAL_OUTPUT);\r
1413 \r
1414     // We don't support real-to-real transform\r
1415     CV_Assert(is_complex_input || is_complex_output);\r
1416 \r
1417     GpuMat src_data;\r
1418 \r
1419     // Make sure here we work with the continuous input, \r
1420     // as CUFFT can't handle gaps\r
1421     src_data = src;\r
1422     createContinuous(src.rows, src.cols, src.type(), src_data);\r
1423     if (src_data.data != src.data)\r
1424         src.copyTo(src_data);\r
1425 \r
1426     Size dft_size_opt = dft_size;\r
1427     if (is_1d_input && !is_row_dft)\r
1428     {\r
1429         // If the source matrix is single column handle it as single row\r
1430         dft_size_opt.width = std::max(dft_size.width, dft_size.height);\r
1431         dft_size_opt.height = std::min(dft_size.width, dft_size.height);\r
1432     }\r
1433 \r
1434     cufftType dft_type = CUFFT_R2C;\r
1435     if (is_complex_input) \r
1436         dft_type = is_complex_output ? CUFFT_C2C : CUFFT_C2R;\r
1437 \r
1438     CV_Assert(dft_size_opt.width > 1);\r
1439 \r
1440     cufftHandle plan;\r
1441     if (is_1d_input || is_row_dft)\r
1442         cufftPlan1d(&plan, dft_size_opt.width, dft_type, dft_size_opt.height);\r
1443     else\r
1444         cufftPlan2d(&plan, dft_size_opt.height, dft_size_opt.width, dft_type);\r
1445 \r
1446     if (is_complex_input)\r
1447     {\r
1448         if (is_complex_output)\r
1449         {\r
1450             createContinuous(dft_size, CV_32FC2, dst);\r
1451             cufftSafeCall(cufftExecC2C(\r
1452                     plan, src_data.ptr<cufftComplex>(), dst.ptr<cufftComplex>(),\r
1453                     is_inverse ? CUFFT_INVERSE : CUFFT_FORWARD));\r
1454         }\r
1455         else\r
1456         {\r
1457             createContinuous(dft_size, CV_32F, dst);\r
1458             cufftSafeCall(cufftExecC2R(\r
1459                     plan, src_data.ptr<cufftComplex>(), dst.ptr<cufftReal>()));\r
1460         }\r
1461     }\r
1462     else\r
1463     {\r
1464         // We could swap dft_size for efficiency. Here we must reflect it\r
1465         if (dft_size == dft_size_opt)\r
1466             createContinuous(Size(dft_size.width / 2 + 1, dft_size.height), CV_32FC2, dst);\r
1467         else\r
1468             createContinuous(Size(dft_size.width, dft_size.height / 2 + 1), CV_32FC2, dst);\r
1469 \r
1470         cufftSafeCall(cufftExecR2C(\r
1471                 plan, src_data.ptr<cufftReal>(), dst.ptr<cufftComplex>()));\r
1472     }\r
1473 \r
1474     cufftSafeCall(cufftDestroy(plan));\r
1475 \r
1476     if (is_scaled_dft)\r
1477         multiply(dst, Scalar::all(1. / dft_size.area()), dst);\r
1478 }\r
1479 \r
1480 //////////////////////////////////////////////////////////////////////////////\r
1481 // convolve\r
1482 \r
1483 \r
1484 void cv::gpu::ConvolveBuf::create(Size image_size, Size templ_size)\r
1485 {\r
1486     result_size = Size(image_size.width - templ_size.width + 1,\r
1487                        image_size.height - templ_size.height + 1);\r
1488     block_size = estimateBlockSize(result_size, templ_size);\r
1489 \r
1490     dft_size.width = getOptimalDFTSize(block_size.width + templ_size.width - 1);\r
1491     dft_size.height = getOptimalDFTSize(block_size.width + templ_size.height - 1);\r
1492     createContinuous(dft_size, CV_32F, image_block);\r
1493     createContinuous(dft_size, CV_32F, templ_block);\r
1494     createContinuous(dft_size, CV_32F, result_data);\r
1495 \r
1496     spect_len = dft_size.height * (dft_size.width / 2 + 1);\r
1497     createContinuous(1, spect_len, CV_32FC2, image_spect);\r
1498     createContinuous(1, spect_len, CV_32FC2, templ_spect);\r
1499     createContinuous(1, spect_len, CV_32FC2, result_spect);\r
1500 \r
1501     block_size.width = std::min(dft_size.width - templ_size.width + 1, result_size.width);\r
1502     block_size.height = std::min(dft_size.height - templ_size.height + 1, result_size.height);\r
1503 }\r
1504 \r
1505 \r
1506 Size cv::gpu::ConvolveBuf::estimateBlockSize(Size result_size, Size templ_size)\r
1507 {\r
1508     int scale = 40;\r
1509     Size bsize_min(1024, 1024);\r
1510 \r
1511     // Check whether we use Fermi generation or newer GPU\r
1512     if (DeviceInfo().majorVersion() >= 2)\r
1513     {\r
1514         bsize_min.width = 2048;\r
1515         bsize_min.height = 2048;\r
1516     }\r
1517 \r
1518     Size bsize(std::max(templ_size.width * scale, bsize_min.width),\r
1519                std::max(templ_size.height * scale, bsize_min.height));\r
1520 \r
1521     bsize.width = std::min(bsize.width, result_size.width);\r
1522     bsize.height = std::min(bsize.height, result_size.height);\r
1523     return bsize;\r
1524 }\r
1525 \r
1526 \r
1527 void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, \r
1528                        bool ccorr)\r
1529 {\r
1530     ConvolveBuf buf;\r
1531     convolve(image, templ, result, ccorr, buf);\r
1532 }\r
1533 \r
1534 \r
1535 void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, \r
1536                        bool ccorr, ConvolveBuf& buf)\r
1537 {\r
1538     StaticAssert<sizeof(float) == sizeof(cufftReal)>::check();\r
1539     StaticAssert<sizeof(float) * 2 == sizeof(cufftComplex)>::check();\r
1540 \r
1541     CV_Assert(image.type() == CV_32F);\r
1542     CV_Assert(templ.type() == CV_32F);\r
1543 \r
1544     buf.create(image.size(), templ.size());\r
1545     result.create(buf.result_size, CV_32F);\r
1546 \r
1547     Size& block_size = buf.block_size;\r
1548     Size& dft_size = buf.dft_size;\r
1549 \r
1550     GpuMat& image_block = buf.image_block;\r
1551     GpuMat& templ_block = buf.templ_block;\r
1552     GpuMat& result_data = buf.result_data;\r
1553 \r
1554     GpuMat& image_spect = buf.image_spect;\r
1555     GpuMat& templ_spect = buf.templ_spect;\r
1556     GpuMat& result_spect = buf.result_spect;\r
1557 \r
1558     cufftHandle planR2C, planC2R;\r
1559     cufftSafeCall(cufftPlan2d(&planC2R, dft_size.height, dft_size.width, CUFFT_C2R));\r
1560     cufftSafeCall(cufftPlan2d(&planR2C, dft_size.height, dft_size.width, CUFFT_R2C));\r
1561 \r
1562     GpuMat templ_roi(templ.size(), CV_32F, templ.data, templ.step);\r
1563     copyMakeBorder(templ_roi, templ_block, 0, templ_block.rows - templ_roi.rows, 0, \r
1564                    templ_block.cols - templ_roi.cols, 0);\r
1565 \r
1566     cufftSafeCall(cufftExecR2C(planR2C, templ_block.ptr<cufftReal>(), \r
1567                                templ_spect.ptr<cufftComplex>()));\r
1568 \r
1569     // Process all blocks of the result matrix\r
1570     for (int y = 0; y < result.rows; y += block_size.height)\r
1571     {\r
1572         for (int x = 0; x < result.cols; x += block_size.width)\r
1573         {\r
1574             Size image_roi_size(std::min(x + dft_size.width, image.cols) - x,\r
1575                                 std::min(y + dft_size.height, image.rows) - y);\r
1576             GpuMat image_roi(image_roi_size, CV_32F, (void*)(image.ptr<float>(y) + x), \r
1577                              image.step);\r
1578             copyMakeBorder(image_roi, image_block, 0, image_block.rows - image_roi.rows,\r
1579                            0, image_block.cols - image_roi.cols, 0);\r
1580 \r
1581             cufftSafeCall(cufftExecR2C(planR2C, image_block.ptr<cufftReal>(), \r
1582                                        image_spect.ptr<cufftComplex>()));\r
1583             mulAndScaleSpectrums(image_spect, templ_spect, result_spect, 0,\r
1584                                  1.f / dft_size.area(), ccorr);\r
1585             cufftSafeCall(cufftExecC2R(planC2R, result_spect.ptr<cufftComplex>(), \r
1586                                        result_data.ptr<cufftReal>()));\r
1587 \r
1588             Size result_roi_size(std::min(x + block_size.width, result.cols) - x,\r
1589                                  std::min(y + block_size.height, result.rows) - y);\r
1590             GpuMat result_roi(result_roi_size, result.type(), \r
1591                               (void*)(result.ptr<float>(y) + x), result.step);\r
1592             GpuMat result_block(result_roi_size, result_data.type(), \r
1593                                 result_data.ptr(), result_data.step);\r
1594             result_block.copyTo(result_roi);\r
1595         }\r
1596     }\r
1597 \r
1598     cufftSafeCall(cufftDestroy(planR2C));\r
1599     cufftSafeCall(cufftDestroy(planC2R));\r
1600 }\r
1601 \r
1602 //////////////////////////////////////////////////////////////////////////////\r
1603 // pyrDown\r
1604 \r
1605 namespace cv { namespace gpu { namespace imgproc\r
1606 {\r
1607     template <typename T, int cn> void pyrDown_gpu(const DevMem2D& src, const DevMem2D& dst, int borderType, cudaStream_t stream);\r
1608 }}}\r
1609 \r
1610 void cv::gpu::pyrDown(const GpuMat& src, GpuMat& dst, int borderType, Stream& stream)\r
1611 {\r
1612     using namespace cv::gpu::imgproc;\r
1613 \r
1614     typedef void (*func_t)(const DevMem2D& src, const DevMem2D& dst, int borderType, cudaStream_t stream);\r
1615 \r
1616     static const func_t funcs[6][4] = \r
1617     {\r
1618         {pyrDown_gpu<uchar, 1>, pyrDown_gpu<uchar, 2>, pyrDown_gpu<uchar, 3>, pyrDown_gpu<uchar, 4>},\r
1619         {pyrDown_gpu<schar, 1>, pyrDown_gpu<schar, 2>, pyrDown_gpu<schar, 3>, pyrDown_gpu<schar, 4>},\r
1620         {pyrDown_gpu<ushort, 1>, pyrDown_gpu<ushort, 2>, pyrDown_gpu<ushort, 3>, pyrDown_gpu<ushort, 4>},\r
1621         {pyrDown_gpu<short, 1>, pyrDown_gpu<short, 2>, pyrDown_gpu<short, 3>, pyrDown_gpu<short, 4>},\r
1622         {pyrDown_gpu<int, 1>, pyrDown_gpu<int, 2>, pyrDown_gpu<int, 3>, pyrDown_gpu<int, 4>},\r
1623         {pyrDown_gpu<float, 1>, pyrDown_gpu<float, 2>, pyrDown_gpu<float, 3>, pyrDown_gpu<float, 4>},\r
1624     };\r
1625 \r
1626     CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);\r
1627 \r
1628     CV_Assert(borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP);\r
1629     int gpuBorderType;\r
1630     CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
1631 \r
1632     dst.create((src.rows + 1) / 2, (src.cols + 1) / 2, src.type());\r
1633 \r
1634     funcs[src.depth()][src.channels() - 1](src, dst, gpuBorderType, StreamAccessor::getStream(stream));\r
1635 }\r
1636 \r
1637 \r
1638 //////////////////////////////////////////////////////////////////////////////\r
1639 // pyrUp\r
1640 \r
1641 namespace cv { namespace gpu { namespace imgproc\r
1642 {\r
1643     template <typename T, int cn> void pyrUp_gpu(const DevMem2D& src, const DevMem2D& dst, int borderType, cudaStream_t stream);\r
1644 }}}\r
1645 \r
1646 void cv::gpu::pyrUp(const GpuMat& src, GpuMat& dst, int borderType, Stream& stream)\r
1647 {\r
1648     using namespace cv::gpu::imgproc;\r
1649 \r
1650     typedef void (*func_t)(const DevMem2D& src, const DevMem2D& dst, int borderType, cudaStream_t stream);\r
1651 \r
1652     static const func_t funcs[6][4] = \r
1653     {\r
1654         {pyrUp_gpu<uchar, 1>, pyrUp_gpu<uchar, 2>, pyrUp_gpu<uchar, 3>, pyrUp_gpu<uchar, 4>},\r
1655         {pyrUp_gpu<schar, 1>, pyrUp_gpu<schar, 2>, pyrUp_gpu<schar, 3>, pyrUp_gpu<schar, 4>},\r
1656         {pyrUp_gpu<ushort, 1>, pyrUp_gpu<ushort, 2>, pyrUp_gpu<ushort, 3>, pyrUp_gpu<ushort, 4>},\r
1657         {pyrUp_gpu<short, 1>, pyrUp_gpu<short, 2>, pyrUp_gpu<short, 3>, pyrUp_gpu<short, 4>},\r
1658         {pyrUp_gpu<int, 1>, pyrUp_gpu<int, 2>, pyrUp_gpu<int, 3>, pyrUp_gpu<int, 4>},\r
1659         {pyrUp_gpu<float, 1>, pyrUp_gpu<float, 2>, pyrUp_gpu<float, 3>, pyrUp_gpu<float, 4>},\r
1660     };\r
1661 \r
1662     CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);\r
1663 \r
1664     CV_Assert(borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP);\r
1665     int gpuBorderType;\r
1666     CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
1667 \r
1668     dst.create(src.rows*2, src.cols*2, src.type());\r
1669 \r
1670     funcs[src.depth()][src.channels() - 1](src, dst, gpuBorderType, StreamAccessor::getStream(stream));\r
1671 }\r
1672 \r
1673 \r
1674 //////////////////////////////////////////////////////////////////////////////\r
1675 // Canny\r
1676 \r
1677 cv::gpu::CannyBuf::CannyBuf(const GpuMat& dx_, const GpuMat& dy_) : dx(dx_), dy(dy_)\r
1678 {\r
1679     CV_Assert(dx_.type() == CV_32SC1 && dy_.type() == CV_32SC1 && dx_.size() == dy_.size());\r
1680 \r
1681     create(dx_.size(), -1);\r
1682 }\r
1683 \r
1684 void cv::gpu::CannyBuf::create(const Size& image_size, int apperture_size)\r
1685 {\r
1686     ensureSizeIsEnough(image_size, CV_32SC1, dx);\r
1687     ensureSizeIsEnough(image_size, CV_32SC1, dy);\r
1688 \r
1689     if (apperture_size == 3)\r
1690     {\r
1691         ensureSizeIsEnough(image_size, CV_32SC1, dx_buf);\r
1692         ensureSizeIsEnough(image_size, CV_32SC1, dy_buf);\r
1693     }\r
1694     else if(apperture_size > 0)\r
1695     {\r
1696         if (!filterDX)\r
1697             filterDX = createDerivFilter_GPU(CV_8UC1, CV_32S, 1, 0, apperture_size, BORDER_REPLICATE);\r
1698         if (!filterDY)\r
1699             filterDY = createDerivFilter_GPU(CV_8UC1, CV_32S, 0, 1, apperture_size, BORDER_REPLICATE);\r
1700     }\r
1701 \r
1702     ensureSizeIsEnough(image_size.height + 2, image_size.width + 2, CV_32FC1, edgeBuf);\r
1703 \r
1704     ensureSizeIsEnough(1, image_size.width * image_size.height, CV_16UC2, trackBuf1);\r
1705     ensureSizeIsEnough(1, image_size.width * image_size.height, CV_16UC2, trackBuf2);\r
1706 }\r
1707 \r
1708 void cv::gpu::CannyBuf::release()\r
1709 {\r
1710     dx.release();\r
1711     dy.release();\r
1712     dx_buf.release();\r
1713     dy_buf.release();\r
1714     edgeBuf.release();\r
1715     trackBuf1.release();\r
1716     trackBuf2.release();\r
1717 }\r
1718 \r
1719 namespace cv { namespace gpu { namespace canny\r
1720 {    \r
1721     void calcSobelRowPass_gpu(PtrStep src, PtrStepi dx_buf, PtrStepi dy_buf, int rows, int cols);\r
1722 \r
1723     void calcMagnitude_gpu(PtrStepi dx_buf, PtrStepi dy_buf, PtrStepi dx, PtrStepi dy, PtrStepf mag, int rows, int cols, bool L2Grad);\r
1724     void calcMagnitude_gpu(PtrStepi dx, PtrStepi dy, PtrStepf mag, int rows, int cols, bool L2Grad);\r
1725 \r
1726     void calcMap_gpu(PtrStepi dx, PtrStepi dy, PtrStepf mag, PtrStepi map, int rows, int cols, float low_thresh, float high_thresh);\r
1727     \r
1728     void edgesHysteresisLocal_gpu(PtrStepi map, ushort2* st1, int rows, int cols);\r
1729 \r
1730     void edgesHysteresisGlobal_gpu(PtrStepi map, ushort2* st1, ushort2* st2, int rows, int cols);\r
1731 \r
1732     void getEdges_gpu(PtrStepi map, PtrStep dst, int rows, int cols);\r
1733 }}}\r
1734 \r
1735 namespace\r
1736 {\r
1737     void CannyCaller(CannyBuf& buf, GpuMat& dst, float low_thresh, float high_thresh)\r
1738     {\r
1739         using namespace cv::gpu::canny;\r
1740 \r
1741         calcMap_gpu(buf.dx, buf.dy, buf.edgeBuf, buf.edgeBuf, dst.rows, dst.cols, low_thresh, high_thresh);\r
1742         \r
1743         edgesHysteresisLocal_gpu(buf.edgeBuf, buf.trackBuf1.ptr<ushort2>(), dst.rows, dst.cols);\r
1744         \r
1745         edgesHysteresisGlobal_gpu(buf.edgeBuf, buf.trackBuf1.ptr<ushort2>(), buf.trackBuf2.ptr<ushort2>(), dst.rows, dst.cols);\r
1746         \r
1747         getEdges_gpu(buf.edgeBuf, dst, dst.rows, dst.cols);\r
1748     }\r
1749 }\r
1750 \r
1751 void cv::gpu::Canny(const GpuMat& src, GpuMat& dst, double low_thresh, double high_thresh, int apperture_size, bool L2gradient)\r
1752 {\r
1753     CannyBuf buf(src.size(), apperture_size);\r
1754     Canny(src, buf, dst, low_thresh, high_thresh, apperture_size, L2gradient);\r
1755 }\r
1756 \r
1757 void cv::gpu::Canny(const GpuMat& src, CannyBuf& buf, GpuMat& dst, double low_thresh, double high_thresh, int apperture_size, bool L2gradient)\r
1758 {\r
1759     using namespace cv::gpu::canny;\r
1760 \r
1761     CV_Assert(TargetArchs::builtWith(SHARED_ATOMICS) && DeviceInfo().supports(SHARED_ATOMICS));\r
1762     CV_Assert(src.type() == CV_8UC1);\r
1763 \r
1764     if( low_thresh > high_thresh )\r
1765         std::swap( low_thresh, high_thresh);\r
1766 \r
1767     dst.create(src.size(), CV_8U);\r
1768     dst.setTo(Scalar::all(0));\r
1769     \r
1770     buf.create(src.size(), apperture_size);\r
1771     buf.edgeBuf.setTo(Scalar::all(0));\r
1772 \r
1773     if (apperture_size == 3)\r
1774     {\r
1775         calcSobelRowPass_gpu(src, buf.dx_buf, buf.dy_buf, src.rows, src.cols);\r
1776 \r
1777         calcMagnitude_gpu(buf.dx_buf, buf.dy_buf, buf.dx, buf.dy, buf.edgeBuf, src.rows, src.cols, L2gradient);\r
1778     }\r
1779     else\r
1780     {\r
1781         buf.filterDX->apply(src, buf.dx, Rect(0, 0, src.cols, src.rows));\r
1782         buf.filterDY->apply(src, buf.dy, Rect(0, 0, src.cols, src.rows));\r
1783 \r
1784         calcMagnitude_gpu(buf.dx, buf.dy, buf.edgeBuf, src.rows, src.cols, L2gradient);\r
1785     }\r
1786 \r
1787     CannyCaller(buf, dst, static_cast<float>(low_thresh), static_cast<float>(high_thresh));\r
1788 }\r
1789 \r
1790 void cv::gpu::Canny(const GpuMat& dx, const GpuMat& dy, GpuMat& dst, double low_thresh, double high_thresh, bool L2gradient)\r
1791 {\r
1792     CannyBuf buf(dx, dy);\r
1793     Canny(dx, dy, buf, dst, low_thresh, high_thresh, L2gradient);\r
1794 }\r
1795 \r
1796 void cv::gpu::Canny(const GpuMat& dx, const GpuMat& dy, CannyBuf& buf, GpuMat& dst, double low_thresh, double high_thresh, bool L2gradient)\r
1797 {\r
1798     using namespace cv::gpu::canny;\r
1799 \r
1800     CV_Assert(TargetArchs::builtWith(SHARED_ATOMICS) && DeviceInfo().supports(SHARED_ATOMICS));\r
1801     CV_Assert(dx.type() == CV_32SC1 && dy.type() == CV_32SC1 && dx.size() == dy.size());\r
1802 \r
1803     if( low_thresh > high_thresh )\r
1804         std::swap( low_thresh, high_thresh);\r
1805 \r
1806     dst.create(dx.size(), CV_8U);\r
1807     dst.setTo(Scalar::all(0));\r
1808     \r
1809     buf.dx = dx; buf.dy = dy;\r
1810     buf.create(dx.size(), -1);\r
1811     buf.edgeBuf.setTo(Scalar::all(0));\r
1812 \r
1813     calcMagnitude_gpu(dx, dy, buf.edgeBuf, dx.rows, dx.cols, L2gradient);\r
1814 \r
1815     CannyCaller(buf, dst, static_cast<float>(low_thresh), static_cast<float>(high_thresh));\r
1816 }\r
1817 \r
1818 #endif /* !defined (HAVE_CUDA) */\r
1819 \r
1820 \r