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&, double, double, double, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  60 void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, double, double, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
  61 void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, double, double, 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 r[9], const float rinv[9], float f, float s, float dist,\r
 588                             float half_w, float half_h, cudaStream_t stream);\r
 589 }}}\r
 590 \r
 591 void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s,\r
 592                                  double dist, GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
 593 {\r
 594     CV_Assert(R.size() == Size(3,3) && R.isContinuous() && R.type() == CV_32F);\r
 595     Mat Rinv = R.inv();\r
 596     CV_Assert(Rinv.isContinuous());\r
 597 \r
 598     map_x.create(dst_roi.size(), CV_32F);\r
 599     map_y.create(dst_roi.size(), CV_32F);\r
 600     imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, R.ptr<float>(), Rinv.ptr<float>(),\r
 601                                 static_cast<float>(f), static_cast<float>(s), static_cast<float>(dist), \r
 602                                 0.5f*src_size.width, 0.5f*src_size.height, StreamAccessor::getStream(stream));\r
 603 }\r
 604 \r
 605 //////////////////////////////////////////////////////////////////////////////\r
 606 // buildWarpCylyndricalMaps\r
 607 \r
 608 namespace cv { namespace gpu { namespace imgproc\r
 609 {\r
 610     void buildWarpCylindricalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
 611                                   const float r[9], const float rinv[9], float f, float s,\r
 612                                   float half_w, float half_h, cudaStream_t stream);\r
 613 }}}\r
 614 \r
 615 void cv::gpu::buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s,\r
 616                                        GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
 617 {\r
 618     CV_Assert(R.size() == Size(3,3) && R.isContinuous() && R.type() == CV_32F);\r
 619     Mat Rinv = R.inv();\r
 620     CV_Assert(Rinv.isContinuous());\r
 621 \r
 622     map_x.create(dst_roi.size(), CV_32F);\r
 623     map_y.create(dst_roi.size(), CV_32F);\r
 624     imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, R.ptr<float>(), Rinv.ptr<float>(),\r
 625                                       static_cast<float>(f), static_cast<float>(s), 0.5f*src_size.width, 0.5f*src_size.height, \r
 626                                       StreamAccessor::getStream(stream));\r
 627 }\r
 628 \r
 629 \r
 630 //////////////////////////////////////////////////////////////////////////////\r
 631 // buildWarpSphericalMaps\r
 632 \r
 633 namespace cv { namespace gpu { namespace imgproc\r
 634 {\r
 635     void buildWarpSphericalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
 636                                 const float r[9], const float rinv[9], float f, float s,\r
 637                                 float half_w, float half_h, cudaStream_t stream);\r
 638 }}}\r
 639 \r
 640 void cv::gpu::buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s,\r
 641                                      GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
 642 {\r
 643     CV_Assert(R.size() == Size(3,3) && R.isContinuous() && R.type() == CV_32F);\r
 644     Mat Rinv = R.inv();\r
 645     CV_Assert(Rinv.isContinuous());\r
 646 \r
 647     map_x.create(dst_roi.size(), CV_32F);\r
 648     map_y.create(dst_roi.size(), CV_32F);\r
 649     imgproc::buildWarpSphericalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, R.ptr<float>(), Rinv.ptr<float>(),\r
 650                                     static_cast<float>(f), static_cast<float>(s), 0.5f*src_size.width, 0.5f*src_size.height, \r
 651                                     StreamAccessor::getStream(stream));\r
 652 }\r
 653 \r
 654 ////////////////////////////////////////////////////////////////////////\r
 655 // rotate\r
 656 \r
 657 void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, Stream& s)\r
 658 {\r
 659     static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};\r
 660 \r
 661     CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
 662     CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
 663 \r
 664     dst.create(dsize, src.type());\r
 665 \r
 666     NppiSize srcsz;\r
 667     srcsz.height = src.rows;\r
 668     srcsz.width = src.cols;\r
 669     NppiRect srcroi;\r
 670     srcroi.x = srcroi.y = 0;\r
 671     srcroi.height = src.rows;\r
 672     srcroi.width = src.cols;\r
 673     NppiRect dstroi;\r
 674     dstroi.x = dstroi.y = 0;\r
 675     dstroi.height = dst.rows;\r
 676     dstroi.width = dst.cols;\r
 677 \r
 678     cudaStream_t stream = StreamAccessor::getStream(s);\r
 679 \r
 680     NppStreamHandler h(stream);\r
 681 \r
 682     if (src.type() == CV_8UC1)\r
 683     {\r
 684         nppSafeCall( nppiRotate_8u_C1R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcroi,\r
 685             dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstroi, angle, xShift, yShift, npp_inter[interpolation]) );\r
 686     }\r
 687     else\r
 688     {\r
 689         nppSafeCall( nppiRotate_8u_C4R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcroi,\r
 690             dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstroi, angle, xShift, yShift, npp_inter[interpolation]) );\r
 691     }\r
 692 \r
 693     if (stream == 0)\r
 694         cudaSafeCall( cudaDeviceSynchronize() );\r
 695 }\r
 696 \r
 697 ////////////////////////////////////////////////////////////////////////\r
 698 // integral\r
 699 \r
 700 void cv::gpu::integral(const GpuMat& src, GpuMat& sum, Stream& s)\r
 701 {\r
 702     GpuMat buffer;\r
 703     integralBuffered(src, sum, buffer, s);\r
 704 }\r
 705 \r
 706 void cv::gpu::integralBuffered(const GpuMat& src, GpuMat& sum, GpuMat& buffer, Stream& s)\r
 707 {\r
 708     CV_Assert(src.type() == CV_8UC1);\r
 709 \r
 710     sum.create(src.rows + 1, src.cols + 1, CV_32S);\r
 711     \r
 712     NcvSize32u roiSize;\r
 713     roiSize.width = src.cols;\r
 714     roiSize.height = src.rows;\r
 715 \r
 716         cudaDeviceProp prop;\r
 717         cudaSafeCall( cudaGetDeviceProperties(&prop, cv::gpu::getDevice()) );\r
 718 \r
 719     Ncv32u bufSize;\r
 720     nppSafeCall( nppiStIntegralGetSize_8u32u(roiSize, &bufSize, prop) );\r
 721     ensureSizeIsEnough(1, bufSize, CV_8UC1, buffer);\r
 722 \r
 723     cudaStream_t stream = StreamAccessor::getStream(s);\r
 724 \r
 725     NppStStreamHandler h(stream);\r
 726 \r
 727     nppSafeCall( nppiStIntegral_8u32u_C1R(const_cast<Ncv8u*>(src.ptr<Ncv8u>()), static_cast<int>(src.step), \r
 728         sum.ptr<Ncv32u>(), static_cast<int>(sum.step), roiSize, buffer.ptr<Ncv8u>(), bufSize, prop) );\r
 729 \r
 730     if (stream == 0)\r
 731         cudaSafeCall( cudaDeviceSynchronize() );\r
 732 }\r
 733 \r
 734 void cv::gpu::integral(const GpuMat& src, GpuMat& sum, GpuMat& sqsum, Stream& s)\r
 735 {\r
 736     CV_Assert(src.type() == CV_8UC1);\r
 737 \r
 738     int width = src.cols + 1, height = src.rows + 1;\r
 739 \r
 740     sum.create(height, width, CV_32S);\r
 741     sqsum.create(height, width, CV_32F);\r
 742 \r
 743     NppiSize sz;\r
 744     sz.width = src.cols;\r
 745     sz.height = src.rows;\r
 746 \r
 747     cudaStream_t stream = StreamAccessor::getStream(s);\r
 748 \r
 749     NppStreamHandler h(stream);\r
 750 \r
 751     nppSafeCall( nppiSqrIntegral_8u32s32f_C1R(const_cast<Npp8u*>(src.ptr<Npp8u>()), static_cast<int>(src.step), \r
 752         sum.ptr<Npp32s>(), static_cast<int>(sum.step), sqsum.ptr<Npp32f>(), static_cast<int>(sqsum.step), sz, 0, 0.0f, height) );\r
 753 \r
 754     if (stream == 0)\r
 755         cudaSafeCall( cudaDeviceSynchronize() );\r
 756 }\r
 757 \r
 758 //////////////////////////////////////////////////////////////////////////////\r
 759 // sqrIntegral\r
 760 \r
 761 void cv::gpu::sqrIntegral(const GpuMat& src, GpuMat& sqsum, Stream& s)\r
 762 {\r
 763     CV_Assert(src.type() == CV_8U);\r
 764 \r
 765     NcvSize32u roiSize;\r
 766     roiSize.width = src.cols;\r
 767     roiSize.height = src.rows;\r
 768 \r
 769         cudaDeviceProp prop;\r
 770         cudaSafeCall( cudaGetDeviceProperties(&prop, cv::gpu::getDevice()) );\r
 771 \r
 772     Ncv32u bufSize;\r
 773     nppSafeCall(nppiStSqrIntegralGetSize_8u64u(roiSize, &bufSize, prop));       \r
 774     GpuMat buf(1, bufSize, CV_8U);\r
 775 \r
 776     cudaStream_t stream = StreamAccessor::getStream(s);\r
 777 \r
 778     NppStStreamHandler h(stream);\r
 779 \r
 780     sqsum.create(src.rows + 1, src.cols + 1, CV_64F);\r
 781     nppSafeCall(nppiStSqrIntegral_8u64u_C1R(const_cast<Ncv8u*>(src.ptr<Ncv8u>(0)), static_cast<int>(src.step), \r
 782             sqsum.ptr<Ncv64u>(0), static_cast<int>(sqsum.step), roiSize, buf.ptr<Ncv8u>(0), bufSize, prop));\r
 783 \r
 784     if (stream == 0)\r
 785         cudaSafeCall( cudaDeviceSynchronize() );\r
 786 }\r
 787 \r
 788 //////////////////////////////////////////////////////////////////////////////\r
 789 // columnSum\r
 790 \r
 791 namespace cv { namespace gpu { namespace imgproc\r
 792 {\r
 793     void columnSum_32F(const DevMem2D src, const DevMem2D dst);\r
 794 }}}\r
 795 \r
 796 void cv::gpu::columnSum(const GpuMat& src, GpuMat& dst)\r
 797 {\r
 798     CV_Assert(src.type() == CV_32F);\r
 799 \r
 800     dst.create(src.size(), CV_32F);\r
 801     imgproc::columnSum_32F(src, dst);\r
 802 }\r
 803 \r
 804 void cv::gpu::rectStdDev(const GpuMat& src, const GpuMat& sqr, GpuMat& dst, const Rect& rect, Stream& s)\r
 805 {\r
 806     CV_Assert(src.type() == CV_32SC1 && sqr.type() == CV_32FC1);\r
 807 \r
 808     dst.create(src.size(), CV_32FC1);\r
 809 \r
 810     NppiSize sz;\r
 811     sz.width = src.cols;\r
 812     sz.height = src.rows;\r
 813 \r
 814     NppiRect nppRect;\r
 815     nppRect.height = rect.height;\r
 816     nppRect.width = rect.width;\r
 817     nppRect.x = rect.x;\r
 818     nppRect.y = rect.y;\r
 819 \r
 820     cudaStream_t stream = StreamAccessor::getStream(s);\r
 821 \r
 822     NppStreamHandler h(stream);\r
 823 \r
 824     nppSafeCall( nppiRectStdDev_32s32f_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), sqr.ptr<Npp32f>(), static_cast<int>(sqr.step),\r
 825                 dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, nppRect) );\r
 826 \r
 827     if (stream == 0)\r
 828         cudaSafeCall( cudaDeviceSynchronize() );\r
 829 }\r
 830 \r
 831 \r
 832 ////////////////////////////////////////////////////////////////////////\r
 833 // Histogram\r
 834 \r
 835 namespace\r
 836 {\r
 837     template<int n> struct NPPTypeTraits;\r
 838     template<> struct NPPTypeTraits<CV_8U>  { typedef Npp8u npp_type; };\r
 839     template<> struct NPPTypeTraits<CV_16U> { typedef Npp16u npp_type; };\r
 840     template<> struct NPPTypeTraits<CV_16S> { typedef Npp16s npp_type; };\r
 841     template<> struct NPPTypeTraits<CV_32F> { typedef Npp32f npp_type; };\r
 842 \r
 843     typedef NppStatus (*get_buf_size_c1_t)(NppiSize oSizeROI, int nLevels, int* hpBufferSize);\r
 844     typedef NppStatus (*get_buf_size_c4_t)(NppiSize oSizeROI, int nLevels[], int* hpBufferSize);\r
 845 \r
 846     template<int SDEPTH> struct NppHistogramEvenFuncC1\r
 847     {\r
 848         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;\r
 849 \r
 850         typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s * pHist,\r
 851                     int nLevels, Npp32s nLowerLevel, Npp32s nUpperLevel, Npp8u * pBuffer);\r
 852     };\r
 853     template<int SDEPTH> struct NppHistogramEvenFuncC4\r
 854     {\r
 855         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;\r
 856 \r
 857         typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI,\r
 858             Npp32s * pHist[4], int nLevels[4], Npp32s nLowerLevel[4], Npp32s nUpperLevel[4], Npp8u * pBuffer);\r
 859     };\r
 860 \r
 861     template<int SDEPTH, typename NppHistogramEvenFuncC1<SDEPTH>::func_ptr func, get_buf_size_c1_t get_buf_size>\r
 862     struct NppHistogramEvenC1\r
 863     {\r
 864         typedef typename NppHistogramEvenFuncC1<SDEPTH>::src_t src_t;\r
 865 \r
 866         static void hist(const GpuMat& src, GpuMat& hist, GpuMat& buffer, int histSize, int lowerLevel, int upperLevel, cudaStream_t stream)\r
 867         {\r
 868             int levels = histSize + 1;\r
 869             hist.create(1, histSize, CV_32S);\r
 870 \r
 871             NppiSize sz;\r
 872             sz.width = src.cols;\r
 873             sz.height = src.rows;\r
 874 \r
 875             int buf_size;\r
 876             get_buf_size(sz, levels, &buf_size);\r
 877 \r
 878             ensureSizeIsEnough(1, buf_size, CV_8U, buffer);\r
 879 \r
 880             NppStreamHandler h(stream);\r
 881 \r
 882             nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), sz, hist.ptr<Npp32s>(), levels,\r
 883                 lowerLevel, upperLevel, buffer.ptr<Npp8u>()) );\r
 884 \r
 885             if (stream == 0)\r
 886                 cudaSafeCall( cudaDeviceSynchronize() );\r
 887         }\r
 888     };\r
 889     template<int SDEPTH, typename NppHistogramEvenFuncC4<SDEPTH>::func_ptr func, get_buf_size_c4_t get_buf_size>\r
 890     struct NppHistogramEvenC4\r
 891     {\r
 892         typedef typename NppHistogramEvenFuncC4<SDEPTH>::src_t src_t;\r
 893 \r
 894         static void hist(const GpuMat& src, GpuMat hist[4], GpuMat& buffer, int histSize[4], int lowerLevel[4], int upperLevel[4], cudaStream_t stream)\r
 895         {\r
 896             int levels[] = {histSize[0] + 1, histSize[1] + 1, histSize[2] + 1, histSize[3] + 1};\r
 897             hist[0].create(1, histSize[0], CV_32S);\r
 898             hist[1].create(1, histSize[1], CV_32S);\r
 899             hist[2].create(1, histSize[2], CV_32S);\r
 900             hist[3].create(1, histSize[3], CV_32S);\r
 901 \r
 902             NppiSize sz;\r
 903             sz.width = src.cols;\r
 904             sz.height = src.rows;\r
 905 \r
 906             Npp32s* pHist[] = {hist[0].ptr<Npp32s>(), hist[1].ptr<Npp32s>(), hist[2].ptr<Npp32s>(), hist[3].ptr<Npp32s>()};\r
 907 \r
 908             int buf_size;\r
 909             get_buf_size(sz, levels, &buf_size);\r
 910 \r
 911             ensureSizeIsEnough(1, buf_size, CV_8U, buffer);\r
 912 \r
 913             NppStreamHandler h(stream);\r
 914 \r
 915             nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), sz, pHist, levels, lowerLevel, upperLevel, buffer.ptr<Npp8u>()) );\r
 916 \r
 917             if (stream == 0)\r
 918                 cudaSafeCall( cudaDeviceSynchronize() );\r
 919         }\r
 920     };\r
 921 \r
 922     template<int SDEPTH> struct NppHistogramRangeFuncC1\r
 923     {\r
 924         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;\r
 925         typedef Npp32s level_t;\r
 926         enum {LEVEL_TYPE_CODE=CV_32SC1};\r
 927 \r
 928         typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s* pHist,\r
 929             const Npp32s* pLevels, int nLevels, Npp8u* pBuffer);\r
 930     };\r
 931     template<> struct NppHistogramRangeFuncC1<CV_32F>\r
 932     {\r
 933         typedef Npp32f src_t;\r
 934         typedef Npp32f level_t;\r
 935         enum {LEVEL_TYPE_CODE=CV_32FC1};\r
 936 \r
 937         typedef NppStatus (*func_ptr)(const Npp32f* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s* pHist,\r
 938             const Npp32f* pLevels, int nLevels, Npp8u* pBuffer);\r
 939     };\r
 940     template<int SDEPTH> struct NppHistogramRangeFuncC4\r
 941     {\r
 942         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;\r
 943         typedef Npp32s level_t;\r
 944         enum {LEVEL_TYPE_CODE=CV_32SC1};\r
 945 \r
 946         typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s* pHist[4],\r
 947             const Npp32s* pLevels[4], int nLevels[4], Npp8u* pBuffer);\r
 948     };\r
 949     template<> struct NppHistogramRangeFuncC4<CV_32F>\r
 950     {\r
 951         typedef Npp32f src_t;\r
 952         typedef Npp32f level_t;\r
 953         enum {LEVEL_TYPE_CODE=CV_32FC1};\r
 954 \r
 955         typedef NppStatus (*func_ptr)(const Npp32f* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s* pHist[4],\r
 956             const Npp32f* pLevels[4], int nLevels[4], Npp8u* pBuffer);\r
 957     };\r
 958 \r
 959     template<int SDEPTH, typename NppHistogramRangeFuncC1<SDEPTH>::func_ptr func, get_buf_size_c1_t get_buf_size>\r
 960     struct NppHistogramRangeC1\r
 961     {\r
 962         typedef typename NppHistogramRangeFuncC1<SDEPTH>::src_t src_t;\r
 963         typedef typename NppHistogramRangeFuncC1<SDEPTH>::level_t level_t;\r
 964         enum {LEVEL_TYPE_CODE=NppHistogramRangeFuncC1<SDEPTH>::LEVEL_TYPE_CODE};\r
 965 \r
 966         static void hist(const GpuMat& src, GpuMat& hist, const GpuMat& levels, GpuMat& buffer, cudaStream_t stream)\r
 967         {\r
 968             CV_Assert(levels.type() == LEVEL_TYPE_CODE && levels.rows == 1);\r
 969 \r
 970             hist.create(1, levels.cols - 1, CV_32S);\r
 971 \r
 972             NppiSize sz;\r
 973             sz.width = src.cols;\r
 974             sz.height = src.rows;\r
 975 \r
 976             int buf_size;\r
 977             get_buf_size(sz, levels.cols, &buf_size);\r
 978             \r
 979             ensureSizeIsEnough(1, buf_size, CV_8U, buffer);\r
 980 \r
 981             NppStreamHandler h(stream);\r
 982 \r
 983             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
 984 \r
 985             if (stream == 0)\r
 986                 cudaSafeCall( cudaDeviceSynchronize() );\r
 987         }\r
 988     };\r
 989     template<int SDEPTH, typename NppHistogramRangeFuncC4<SDEPTH>::func_ptr func, get_buf_size_c4_t get_buf_size>\r
 990     struct NppHistogramRangeC4\r
 991     {\r
 992         typedef typename NppHistogramRangeFuncC4<SDEPTH>::src_t src_t;\r
 993         typedef typename NppHistogramRangeFuncC1<SDEPTH>::level_t level_t;\r
 994         enum {LEVEL_TYPE_CODE=NppHistogramRangeFuncC1<SDEPTH>::LEVEL_TYPE_CODE};\r
 995 \r
 996         static void hist(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4], GpuMat& buffer, cudaStream_t stream)\r
 997         {\r
 998             CV_Assert(levels[0].type() == LEVEL_TYPE_CODE && levels[0].rows == 1);\r
 999             CV_Assert(levels[1].type() == LEVEL_TYPE_CODE && levels[1].rows == 1);\r
1000             CV_Assert(levels[2].type() == LEVEL_TYPE_CODE && levels[2].rows == 1);\r
1001             CV_Assert(levels[3].type() == LEVEL_TYPE_CODE && levels[3].rows == 1);\r
1002 \r
1003             hist[0].create(1, levels[0].cols - 1, CV_32S);\r
1004             hist[1].create(1, levels[1].cols - 1, CV_32S);\r
1005             hist[2].create(1, levels[2].cols - 1, CV_32S);\r
1006             hist[3].create(1, levels[3].cols - 1, CV_32S);\r
1007 \r
1008             Npp32s* pHist[] = {hist[0].ptr<Npp32s>(), hist[1].ptr<Npp32s>(), hist[2].ptr<Npp32s>(), hist[3].ptr<Npp32s>()};\r
1009             int nLevels[] = {levels[0].cols, levels[1].cols, levels[2].cols, levels[3].cols};\r
1010             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
1011 \r
1012             NppiSize sz;\r
1013             sz.width = src.cols;\r
1014             sz.height = src.rows;\r
1015 \r
1016             int buf_size;\r
1017             get_buf_size(sz, nLevels, &buf_size);\r
1018 \r
1019             ensureSizeIsEnough(1, buf_size, CV_8U, buffer);\r
1020 \r
1021             NppStreamHandler h(stream);\r
1022 \r
1023             nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), sz, pHist, pLevels, nLevels, buffer.ptr<Npp8u>()) );\r
1024 \r
1025             if (stream == 0)\r
1026                 cudaSafeCall( cudaDeviceSynchronize() );\r
1027         }\r
1028     };\r
1029 }\r
1030 \r
1031 void cv::gpu::evenLevels(GpuMat& levels, int nLevels, int lowerLevel, int upperLevel)\r
1032 {\r
1033     Mat host_levels(1, nLevels, CV_32SC1);\r
1034     nppSafeCall( nppiEvenLevelsHost_32s(host_levels.ptr<Npp32s>(), nLevels, lowerLevel, upperLevel) );\r
1035     levels.upload(host_levels);\r
1036 }\r
1037 \r
1038 void cv::gpu::histEven(const GpuMat& src, GpuMat& hist, int histSize, int lowerLevel, int upperLevel, Stream& stream)\r
1039 {\r
1040     GpuMat buf;\r
1041     histEven(src, hist, buf, histSize, lowerLevel, upperLevel, stream);\r
1042 }\r
1043 \r
1044 void cv::gpu::histEven(const GpuMat& src, GpuMat& hist, GpuMat& buf, int histSize, int lowerLevel, int upperLevel, Stream& stream)\r
1045 {\r
1046     CV_Assert(src.type() == CV_8UC1 || src.type() == CV_16UC1 || src.type() == CV_16SC1 );\r
1047 \r
1048     typedef void (*hist_t)(const GpuMat& src, GpuMat& hist, GpuMat& buf, int levels, int lowerLevel, int upperLevel, cudaStream_t stream);\r
1049     static const hist_t hist_callers[] =\r
1050     {\r
1051         NppHistogramEvenC1<CV_8U , nppiHistogramEven_8u_C1R , nppiHistogramEvenGetBufferSize_8u_C1R >::hist,\r
1052         0,\r
1053         NppHistogramEvenC1<CV_16U, nppiHistogramEven_16u_C1R, nppiHistogramEvenGetBufferSize_16u_C1R>::hist,\r
1054         NppHistogramEvenC1<CV_16S, nppiHistogramEven_16s_C1R, nppiHistogramEvenGetBufferSize_16s_C1R>::hist\r
1055     };\r
1056 \r
1057     hist_callers[src.depth()](src, hist, buf, histSize, lowerLevel, upperLevel, StreamAccessor::getStream(stream));\r
1058 }\r
1059 \r
1060 void cv::gpu::histEven(const GpuMat& src, GpuMat hist[4], int histSize[4], int lowerLevel[4], int upperLevel[4], Stream& stream)\r
1061 {\r
1062     GpuMat buf;\r
1063     histEven(src, hist, buf, histSize, lowerLevel, upperLevel, stream);\r
1064 }\r
1065 \r
1066 void cv::gpu::histEven(const GpuMat& src, GpuMat hist[4], GpuMat& buf, int histSize[4], int lowerLevel[4], int upperLevel[4], Stream& stream)\r
1067 {\r
1068     CV_Assert(src.type() == CV_8UC4 || src.type() == CV_16UC4 || src.type() == CV_16SC4 );\r
1069 \r
1070     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
1071     static const hist_t hist_callers[] =\r
1072     {\r
1073         NppHistogramEvenC4<CV_8U , nppiHistogramEven_8u_C4R , nppiHistogramEvenGetBufferSize_8u_C4R >::hist,\r
1074         0,\r
1075         NppHistogramEvenC4<CV_16U, nppiHistogramEven_16u_C4R, nppiHistogramEvenGetBufferSize_16u_C4R>::hist,\r
1076         NppHistogramEvenC4<CV_16S, nppiHistogramEven_16s_C4R, nppiHistogramEvenGetBufferSize_16s_C4R>::hist\r
1077     };\r
1078 \r
1079     hist_callers[src.depth()](src, hist, buf, histSize, lowerLevel, upperLevel, StreamAccessor::getStream(stream));\r
1080 }\r
1081 \r
1082 void cv::gpu::histRange(const GpuMat& src, GpuMat& hist, const GpuMat& levels, Stream& stream)\r
1083 {\r
1084     GpuMat buf;\r
1085     histRange(src, hist, levels, buf, stream);\r
1086 }\r
1087 \r
1088 \r
1089 void cv::gpu::histRange(const GpuMat& src, GpuMat& hist, const GpuMat& levels, GpuMat& buf, Stream& stream)\r
1090 {\r
1091     CV_Assert(src.type() == CV_8UC1 || src.type() == CV_16UC1 || src.type() == CV_16SC1 || src.type() == CV_32FC1);\r
1092 \r
1093     typedef void (*hist_t)(const GpuMat& src, GpuMat& hist, const GpuMat& levels, GpuMat& buf, cudaStream_t stream);\r
1094     static const hist_t hist_callers[] =\r
1095     {\r
1096         NppHistogramRangeC1<CV_8U , nppiHistogramRange_8u_C1R , nppiHistogramRangeGetBufferSize_8u_C1R >::hist,\r
1097         0,\r
1098         NppHistogramRangeC1<CV_16U, nppiHistogramRange_16u_C1R, nppiHistogramRangeGetBufferSize_16u_C1R>::hist,\r
1099         NppHistogramRangeC1<CV_16S, nppiHistogramRange_16s_C1R, nppiHistogramRangeGetBufferSize_16s_C1R>::hist,\r
1100         0,\r
1101         NppHistogramRangeC1<CV_32F, nppiHistogramRange_32f_C1R, nppiHistogramRangeGetBufferSize_32f_C1R>::hist\r
1102     };\r
1103 \r
1104     hist_callers[src.depth()](src, hist, levels, buf, StreamAccessor::getStream(stream));\r
1105 }\r
1106 \r
1107 void cv::gpu::histRange(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4], Stream& stream)\r
1108 {\r
1109     GpuMat buf;\r
1110     histRange(src, hist, levels, buf, stream);\r
1111 }\r
1112 \r
1113 void cv::gpu::histRange(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4], GpuMat& buf, Stream& stream)\r
1114 {\r
1115     CV_Assert(src.type() == CV_8UC4 || src.type() == CV_16UC4 || src.type() == CV_16SC4 || src.type() == CV_32FC4);\r
1116 \r
1117     typedef void (*hist_t)(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4], GpuMat& buf, cudaStream_t stream);\r
1118     static const hist_t hist_callers[] =\r
1119     {\r
1120         NppHistogramRangeC4<CV_8U , nppiHistogramRange_8u_C4R , nppiHistogramRangeGetBufferSize_8u_C4R >::hist,\r
1121         0,\r
1122         NppHistogramRangeC4<CV_16U, nppiHistogramRange_16u_C4R, nppiHistogramRangeGetBufferSize_16u_C4R>::hist,\r
1123         NppHistogramRangeC4<CV_16S, nppiHistogramRange_16s_C4R, nppiHistogramRangeGetBufferSize_16s_C4R>::hist,\r
1124         0,\r
1125         NppHistogramRangeC4<CV_32F, nppiHistogramRange_32f_C4R, nppiHistogramRangeGetBufferSize_32f_C4R>::hist\r
1126     };\r
1127 \r
1128     hist_callers[src.depth()](src, hist, levels, buf, StreamAccessor::getStream(stream));\r
1129 }\r
1130 \r
1131 namespace cv { namespace gpu { namespace histograms\r
1132 {\r
1133     void histogram256_gpu(DevMem2D src, int* hist, unsigned int* buf, cudaStream_t stream);\r
1134 \r
1135     const int PARTIAL_HISTOGRAM256_COUNT = 240;\r
1136     const int HISTOGRAM256_BIN_COUNT     = 256;\r
1137 }}}\r
1138 \r
1139 void cv::gpu::calcHist(const GpuMat& src, GpuMat& hist, Stream& stream)\r
1140 {\r
1141     GpuMat buf;\r
1142     calcHist(src, hist, buf, stream);\r
1143 }\r
1144 \r
1145 void cv::gpu::calcHist(const GpuMat& src, GpuMat& hist, GpuMat& buf, Stream& stream)\r
1146 {\r
1147     using namespace cv::gpu::histograms;\r
1148 \r
1149     CV_Assert(src.type() == CV_8UC1);\r
1150 \r
1151     hist.create(1, 256, CV_32SC1);\r
1152 \r
1153     ensureSizeIsEnough(1, PARTIAL_HISTOGRAM256_COUNT * HISTOGRAM256_BIN_COUNT, CV_32SC1, buf);\r
1154 \r
1155     histogram256_gpu(src, hist.ptr<int>(), buf.ptr<unsigned int>(), StreamAccessor::getStream(stream));\r
1156 }\r
1157 \r
1158 void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, Stream& stream)\r
1159 {\r
1160     GpuMat hist;\r
1161     GpuMat buf;\r
1162     equalizeHist(src, dst, hist, buf, stream);\r
1163 }\r
1164 \r
1165 void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, Stream& stream)\r
1166 {\r
1167     GpuMat buf;\r
1168     equalizeHist(src, dst, hist, buf, stream);\r
1169 }\r
1170 \r
1171 namespace cv { namespace gpu { namespace histograms\r
1172 {\r
1173     void equalizeHist_gpu(DevMem2D src, DevMem2D dst, const int* lut, cudaStream_t stream);\r
1174 }}}\r
1175 \r
1176 void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, GpuMat& buf, Stream& s)\r
1177 {\r
1178     using namespace cv::gpu::histograms;\r
1179 \r
1180     CV_Assert(src.type() == CV_8UC1);\r
1181 \r
1182     dst.create(src.size(), src.type());\r
1183 \r
1184     int intBufSize;\r
1185     nppSafeCall( nppsIntegralGetBufferSize_32s(256, &intBufSize) );\r
1186 \r
1187     int bufSize = static_cast<int>(std::max(256 * 240 * sizeof(int), intBufSize + 256 * sizeof(int)));\r
1188 \r
1189     ensureSizeIsEnough(1, bufSize, CV_8UC1, buf);\r
1190 \r
1191     GpuMat histBuf(1, 256 * 240, CV_32SC1, buf.ptr());\r
1192     GpuMat intBuf(1, intBufSize, CV_8UC1, buf.ptr());\r
1193     GpuMat lut(1, 256, CV_32S, buf.ptr() + intBufSize);\r
1194 \r
1195     calcHist(src, hist, histBuf, s);\r
1196 \r
1197     cudaStream_t stream = StreamAccessor::getStream(s);\r
1198 \r
1199     NppStreamHandler h(stream);\r
1200 \r
1201     nppSafeCall( nppsIntegral_32s(hist.ptr<Npp32s>(), lut.ptr<Npp32s>(), 256, intBuf.ptr<Npp8u>()) );\r
1202     \r
1203     if (stream == 0)\r
1204         cudaSafeCall( cudaDeviceSynchronize() );\r
1205 \r
1206     equalizeHist_gpu(src, dst, lut.ptr<int>(), stream);\r
1207 }\r
1208 \r
1209 ////////////////////////////////////////////////////////////////////////\r
1210 // cornerHarris & minEgenVal\r
1211 \r
1212 namespace cv { namespace gpu { namespace imgproc {\r
1213 \r
1214     void extractCovData_caller(const DevMem2Df Dx, const DevMem2Df Dy, PtrStepf dst);\r
1215     void cornerHarris_caller(const int block_size, const float k, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst, int border_type);\r
1216     void cornerMinEigenVal_caller(const int block_size, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst, int border_type);\r
1217 \r
1218 }}}\r
1219 \r
1220 namespace \r
1221 {\r
1222     template <typename T>\r
1223     void extractCovData(const GpuMat& src, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, int borderType)\r
1224     {   \r
1225         double scale = (double)(1 << ((ksize > 0 ? ksize : 3) - 1)) * blockSize;\r
1226         if (ksize < 0) \r
1227             scale *= 2.;\r
1228         if (src.depth() == CV_8U)\r
1229             scale *= 255.;\r
1230         scale = 1./scale;\r
1231 \r
1232         Dx.create(src.size(), CV_32F);\r
1233         Dy.create(src.size(), CV_32F);\r
1234 \r
1235         if (ksize > 0)\r
1236         {\r
1237             Sobel(src, Dx, CV_32F, 1, 0, ksize, scale, borderType);\r
1238             Sobel(src, Dy, CV_32F, 0, 1, ksize, scale, borderType);\r
1239         }\r
1240         else\r
1241         {\r
1242             Scharr(src, Dx, CV_32F, 1, 0, scale, borderType);\r
1243             Scharr(src, Dy, CV_32F, 0, 1, scale, borderType);\r
1244         }\r
1245     }\r
1246 \r
1247     void extractCovData(const GpuMat& src, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, int borderType)\r
1248     {\r
1249         switch (src.type())\r
1250         {\r
1251         case CV_8U:\r
1252             extractCovData<unsigned char>(src, Dx, Dy, blockSize, ksize, borderType);\r
1253             break;\r
1254         case CV_32F:\r
1255             extractCovData<float>(src, Dx, Dy, blockSize, ksize, borderType);\r
1256             break;\r
1257         default:\r
1258             CV_Error(CV_StsBadArg, "extractCovData: unsupported type of the source matrix");\r
1259         }\r
1260     }\r
1261 \r
1262 } // Anonymous namespace\r
1263 \r
1264 \r
1265 bool cv::gpu::tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType)\r
1266 {\r
1267     switch (cpuBorderType)\r
1268     {\r
1269     case cv::BORDER_REFLECT101:\r
1270         gpuBorderType = cv::gpu::BORDER_REFLECT101_GPU;\r
1271         return true;\r
1272     case cv::BORDER_REPLICATE:\r
1273         gpuBorderType = cv::gpu::BORDER_REPLICATE_GPU;\r
1274         return true;\r
1275     case cv::BORDER_CONSTANT:\r
1276         gpuBorderType = cv::gpu::BORDER_CONSTANT_GPU;\r
1277         return true;\r
1278     case cv::BORDER_REFLECT:\r
1279         gpuBorderType = cv::gpu::BORDER_REFLECT_GPU;\r
1280         return true;\r
1281     case cv::BORDER_WRAP:\r
1282         gpuBorderType = cv::gpu::BORDER_WRAP_GPU;\r
1283         return true;\r
1284     default:\r
1285         return false;\r
1286     };\r
1287     return false;\r
1288 }\r
1289 \r
1290 void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, double k, int borderType)\r
1291 {\r
1292     GpuMat Dx, Dy;\r
1293     cornerHarris(src, dst, Dx, Dy, blockSize, ksize, k, borderType);\r
1294 }\r
1295 \r
1296 void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, double k, int borderType)\r
1297 {\r
1298     CV_Assert(borderType == cv::BORDER_REFLECT101 ||\r
1299               borderType == cv::BORDER_REPLICATE);\r
1300 \r
1301     int gpuBorderType;\r
1302     CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
1303 \r
1304     extractCovData(src, Dx, Dy, blockSize, ksize, borderType);\r
1305     dst.create(src.size(), CV_32F);\r
1306     imgproc::cornerHarris_caller(blockSize, (float)k, Dx, Dy, dst, gpuBorderType);\r
1307 }\r
1308 \r
1309 void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, int borderType)\r
1310 {  \r
1311     GpuMat Dx, Dy;\r
1312     cornerMinEigenVal(src, dst, Dx, Dy, blockSize, ksize, borderType);\r
1313 }\r
1314 \r
1315 void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, int borderType)\r
1316 {  \r
1317     CV_Assert(borderType == cv::BORDER_REFLECT101 ||\r
1318               borderType == cv::BORDER_REPLICATE);\r
1319 \r
1320     int gpuBorderType;\r
1321     CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
1322 \r
1323     extractCovData(src, Dx, Dy, blockSize, ksize, borderType);    \r
1324     dst.create(src.size(), CV_32F);\r
1325     imgproc::cornerMinEigenVal_caller(blockSize, Dx, Dy, dst, gpuBorderType);\r
1326 }\r
1327 \r
1328 //////////////////////////////////////////////////////////////////////////////\r
1329 // mulSpectrums\r
1330 \r
1331 namespace cv { namespace gpu { namespace imgproc \r
1332 {\r
1333     void mulSpectrums(const PtrStep_<cufftComplex> a, const PtrStep_<cufftComplex> b, \r
1334                       DevMem2D_<cufftComplex> c);\r
1335 \r
1336     void mulSpectrums_CONJ(const PtrStep_<cufftComplex> a, const PtrStep_<cufftComplex> b, \r
1337                            DevMem2D_<cufftComplex> c);\r
1338 }}}\r
1339 \r
1340 \r
1341 void cv::gpu::mulSpectrums(const GpuMat& a, const GpuMat& b, GpuMat& c, \r
1342                            int flags, bool conjB) \r
1343 {\r
1344     typedef void (*Caller)(const PtrStep_<cufftComplex>, const PtrStep_<cufftComplex>, \r
1345                            DevMem2D_<cufftComplex>);\r
1346     static Caller callers[] = { imgproc::mulSpectrums, \r
1347                                 imgproc::mulSpectrums_CONJ };\r
1348 \r
1349     CV_Assert(a.type() == b.type() && a.type() == CV_32FC2);\r
1350     CV_Assert(a.size() == b.size());\r
1351 \r
1352     c.create(a.size(), CV_32FC2);\r
1353 \r
1354     Caller caller = callers[(int)conjB];\r
1355     caller(a, b, c);\r
1356 }\r
1357 \r
1358 //////////////////////////////////////////////////////////////////////////////\r
1359 // mulAndScaleSpectrums\r
1360 \r
1361 namespace cv { namespace gpu { namespace imgproc \r
1362 {\r
1363     void mulAndScaleSpectrums(const PtrStep_<cufftComplex> a, const PtrStep_<cufftComplex> b,\r
1364                              float scale, DevMem2D_<cufftComplex> c);\r
1365 \r
1366     void mulAndScaleSpectrums_CONJ(const PtrStep_<cufftComplex> a, const PtrStep_<cufftComplex> b,\r
1367                                   float scale, DevMem2D_<cufftComplex> c);\r
1368 }}}\r
1369 \r
1370 \r
1371 void cv::gpu::mulAndScaleSpectrums(const GpuMat& a, const GpuMat& b, GpuMat& c,\r
1372                                   int flags, float scale, bool conjB) \r
1373 {\r
1374     typedef void (*Caller)(const PtrStep_<cufftComplex>, const PtrStep_<cufftComplex>,\r
1375                            float scale, DevMem2D_<cufftComplex>);\r
1376     static Caller callers[] = { imgproc::mulAndScaleSpectrums, \r
1377                                 imgproc::mulAndScaleSpectrums_CONJ };\r
1378 \r
1379     CV_Assert(a.type() == b.type() && a.type() == CV_32FC2);\r
1380     CV_Assert(a.size() == b.size());\r
1381 \r
1382     c.create(a.size(), CV_32FC2);\r
1383 \r
1384     Caller caller = callers[(int)conjB];\r
1385     caller(a, b, scale, c);\r
1386 }\r
1387 \r
1388 //////////////////////////////////////////////////////////////////////////////\r
1389 // dft\r
1390 \r
1391 void cv::gpu::dft(const GpuMat& src, GpuMat& dst, Size dft_size, int flags)\r
1392 {\r
1393     CV_Assert(src.type() == CV_32F || src.type() == CV_32FC2);\r
1394 \r
1395     // We don't support unpacked output (in the case of real input)\r
1396     CV_Assert(!(flags & DFT_COMPLEX_OUTPUT));\r
1397 \r
1398     bool is_1d_input = (dft_size.height == 1) || (dft_size.width == 1);\r
1399     int is_row_dft = flags & DFT_ROWS;\r
1400     int is_scaled_dft = flags & DFT_SCALE;\r
1401     int is_inverse = flags & DFT_INVERSE;\r
1402     bool is_complex_input = src.channels() == 2;\r
1403     bool is_complex_output = !(flags & DFT_REAL_OUTPUT);\r
1404 \r
1405     // We don't support real-to-real transform\r
1406     CV_Assert(is_complex_input || is_complex_output);\r
1407 \r
1408     GpuMat src_data;\r
1409 \r
1410     // Make sure here we work with the continuous input, \r
1411     // as CUFFT can't handle gaps\r
1412     src_data = src;\r
1413     createContinuous(src.rows, src.cols, src.type(), src_data);\r
1414     if (src_data.data != src.data)\r
1415         src.copyTo(src_data);\r
1416 \r
1417     Size dft_size_opt = dft_size;\r
1418     if (is_1d_input && !is_row_dft)\r
1419     {\r
1420         // If the source matrix is single column handle it as single row\r
1421         dft_size_opt.width = std::max(dft_size.width, dft_size.height);\r
1422         dft_size_opt.height = std::min(dft_size.width, dft_size.height);\r
1423     }\r
1424 \r
1425     cufftType dft_type = CUFFT_R2C;\r
1426     if (is_complex_input) \r
1427         dft_type = is_complex_output ? CUFFT_C2C : CUFFT_C2R;\r
1428 \r
1429     CV_Assert(dft_size_opt.width > 1);\r
1430 \r
1431     cufftHandle plan;\r
1432     if (is_1d_input || is_row_dft)\r
1433         cufftPlan1d(&plan, dft_size_opt.width, dft_type, dft_size_opt.height);\r
1434     else\r
1435         cufftPlan2d(&plan, dft_size_opt.height, dft_size_opt.width, dft_type);\r
1436 \r
1437     if (is_complex_input)\r
1438     {\r
1439         if (is_complex_output)\r
1440         {\r
1441             createContinuous(dft_size, CV_32FC2, dst);\r
1442             cufftSafeCall(cufftExecC2C(\r
1443                     plan, src_data.ptr<cufftComplex>(), dst.ptr<cufftComplex>(),\r
1444                     is_inverse ? CUFFT_INVERSE : CUFFT_FORWARD));\r
1445         }\r
1446         else\r
1447         {\r
1448             createContinuous(dft_size, CV_32F, dst);\r
1449             cufftSafeCall(cufftExecC2R(\r
1450                     plan, src_data.ptr<cufftComplex>(), dst.ptr<cufftReal>()));\r
1451         }\r
1452     }\r
1453     else\r
1454     {\r
1455         // We could swap dft_size for efficiency. Here we must reflect it\r
1456         if (dft_size == dft_size_opt)\r
1457             createContinuous(Size(dft_size.width / 2 + 1, dft_size.height), CV_32FC2, dst);\r
1458         else\r
1459             createContinuous(Size(dft_size.width, dft_size.height / 2 + 1), CV_32FC2, dst);\r
1460 \r
1461         cufftSafeCall(cufftExecR2C(\r
1462                 plan, src_data.ptr<cufftReal>(), dst.ptr<cufftComplex>()));\r
1463     }\r
1464 \r
1465     cufftSafeCall(cufftDestroy(plan));\r
1466 \r
1467     if (is_scaled_dft)\r
1468         multiply(dst, Scalar::all(1. / dft_size.area()), dst);\r
1469 }\r
1470 \r
1471 //////////////////////////////////////////////////////////////////////////////\r
1472 // convolve\r
1473 \r
1474 \r
1475 void cv::gpu::ConvolveBuf::create(Size image_size, Size templ_size)\r
1476 {\r
1477     result_size = Size(image_size.width - templ_size.width + 1,\r
1478                        image_size.height - templ_size.height + 1);\r
1479     block_size = estimateBlockSize(result_size, templ_size);\r
1480 \r
1481     dft_size.width = getOptimalDFTSize(block_size.width + templ_size.width - 1);\r
1482     dft_size.height = getOptimalDFTSize(block_size.width + templ_size.height - 1);\r
1483     createContinuous(dft_size, CV_32F, image_block);\r
1484     createContinuous(dft_size, CV_32F, templ_block);\r
1485     createContinuous(dft_size, CV_32F, result_data);\r
1486 \r
1487     spect_len = dft_size.height * (dft_size.width / 2 + 1);\r
1488     createContinuous(1, spect_len, CV_32FC2, image_spect);\r
1489     createContinuous(1, spect_len, CV_32FC2, templ_spect);\r
1490     createContinuous(1, spect_len, CV_32FC2, result_spect);\r
1491 \r
1492     block_size.width = std::min(dft_size.width - templ_size.width + 1, result_size.width);\r
1493     block_size.height = std::min(dft_size.height - templ_size.height + 1, result_size.height);\r
1494 }\r
1495 \r
1496 \r
1497 Size cv::gpu::ConvolveBuf::estimateBlockSize(Size result_size, Size templ_size)\r
1498 {\r
1499     int scale = 40;\r
1500     Size bsize_min(1024, 1024);\r
1501 \r
1502     // Check whether we use Fermi generation or newer GPU\r
1503     if (DeviceInfo().majorVersion() >= 2)\r
1504     {\r
1505         bsize_min.width = 2048;\r
1506         bsize_min.height = 2048;\r
1507     }\r
1508 \r
1509     Size bsize(std::max(templ_size.width * scale, bsize_min.width),\r
1510                std::max(templ_size.height * scale, bsize_min.height));\r
1511 \r
1512     bsize.width = std::min(bsize.width, result_size.width);\r
1513     bsize.height = std::min(bsize.height, result_size.height);\r
1514     return bsize;\r
1515 }\r
1516 \r
1517 \r
1518 void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, \r
1519                        bool ccorr)\r
1520 {\r
1521     ConvolveBuf buf;\r
1522     convolve(image, templ, result, ccorr, buf);\r
1523 }\r
1524 \r
1525 \r
1526 void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, \r
1527                        bool ccorr, ConvolveBuf& buf)\r
1528 {\r
1529     StaticAssert<sizeof(float) == sizeof(cufftReal)>::check();\r
1530     StaticAssert<sizeof(float) * 2 == sizeof(cufftComplex)>::check();\r
1531 \r
1532     CV_Assert(image.type() == CV_32F);\r
1533     CV_Assert(templ.type() == CV_32F);\r
1534 \r
1535     buf.create(image.size(), templ.size());\r
1536     result.create(buf.result_size, CV_32F);\r
1537 \r
1538     Size& block_size = buf.block_size;\r
1539     Size& dft_size = buf.dft_size;\r
1540 \r
1541     GpuMat& image_block = buf.image_block;\r
1542     GpuMat& templ_block = buf.templ_block;\r
1543     GpuMat& result_data = buf.result_data;\r
1544 \r
1545     GpuMat& image_spect = buf.image_spect;\r
1546     GpuMat& templ_spect = buf.templ_spect;\r
1547     GpuMat& result_spect = buf.result_spect;\r
1548 \r
1549     cufftHandle planR2C, planC2R;\r
1550     cufftSafeCall(cufftPlan2d(&planC2R, dft_size.height, dft_size.width, CUFFT_C2R));\r
1551     cufftSafeCall(cufftPlan2d(&planR2C, dft_size.height, dft_size.width, CUFFT_R2C));\r
1552 \r
1553     GpuMat templ_roi(templ.size(), CV_32F, templ.data, templ.step);\r
1554     copyMakeBorder(templ_roi, templ_block, 0, templ_block.rows - templ_roi.rows, 0, \r
1555                    templ_block.cols - templ_roi.cols, 0);\r
1556 \r
1557     cufftSafeCall(cufftExecR2C(planR2C, templ_block.ptr<cufftReal>(), \r
1558                                templ_spect.ptr<cufftComplex>()));\r
1559 \r
1560     // Process all blocks of the result matrix\r
1561     for (int y = 0; y < result.rows; y += block_size.height)\r
1562     {\r
1563         for (int x = 0; x < result.cols; x += block_size.width)\r
1564         {\r
1565             Size image_roi_size(std::min(x + dft_size.width, image.cols) - x,\r
1566                                 std::min(y + dft_size.height, image.rows) - y);\r
1567             GpuMat image_roi(image_roi_size, CV_32F, (void*)(image.ptr<float>(y) + x), \r
1568                              image.step);\r
1569             copyMakeBorder(image_roi, image_block, 0, image_block.rows - image_roi.rows,\r
1570                            0, image_block.cols - image_roi.cols, 0);\r
1571 \r
1572             cufftSafeCall(cufftExecR2C(planR2C, image_block.ptr<cufftReal>(), \r
1573                                        image_spect.ptr<cufftComplex>()));\r
1574             mulAndScaleSpectrums(image_spect, templ_spect, result_spect, 0,\r
1575                                  1.f / dft_size.area(), ccorr);\r
1576             cufftSafeCall(cufftExecC2R(planC2R, result_spect.ptr<cufftComplex>(), \r
1577                                        result_data.ptr<cufftReal>()));\r
1578 \r
1579             Size result_roi_size(std::min(x + block_size.width, result.cols) - x,\r
1580                                  std::min(y + block_size.height, result.rows) - y);\r
1581             GpuMat result_roi(result_roi_size, result.type(), \r
1582                               (void*)(result.ptr<float>(y) + x), result.step);\r
1583             GpuMat result_block(result_roi_size, result_data.type(), \r
1584                                 result_data.ptr(), result_data.step);\r
1585             result_block.copyTo(result_roi);\r
1586         }\r
1587     }\r
1588 \r
1589     cufftSafeCall(cufftDestroy(planR2C));\r
1590     cufftSafeCall(cufftDestroy(planC2R));\r
1591 }\r
1592 \r
1593 //////////////////////////////////////////////////////////////////////////////\r
1594 // pyrDown\r
1595 \r
1596 namespace cv { namespace gpu { namespace imgproc\r
1597 {\r
1598     template <typename T, int cn> void pyrDown_gpu(const DevMem2D& src, const DevMem2D& dst, int borderType, cudaStream_t stream);\r
1599 }}}\r
1600 \r
1601 void cv::gpu::pyrDown(const GpuMat& src, GpuMat& dst, int borderType, Stream& stream)\r
1602 {\r
1603     using namespace cv::gpu::imgproc;\r
1604 \r
1605     typedef void (*func_t)(const DevMem2D& src, const DevMem2D& dst, int borderType, cudaStream_t stream);\r
1606 \r
1607     static const func_t funcs[6][4] = \r
1608     {\r
1609         {pyrDown_gpu<uchar, 1>, pyrDown_gpu<uchar, 2>, pyrDown_gpu<uchar, 3>, pyrDown_gpu<uchar, 4>},\r
1610         {pyrDown_gpu<schar, 1>, pyrDown_gpu<schar, 2>, pyrDown_gpu<schar, 3>, pyrDown_gpu<schar, 4>},\r
1611         {pyrDown_gpu<ushort, 1>, pyrDown_gpu<ushort, 2>, pyrDown_gpu<ushort, 3>, pyrDown_gpu<ushort, 4>},\r
1612         {pyrDown_gpu<short, 1>, pyrDown_gpu<short, 2>, pyrDown_gpu<short, 3>, pyrDown_gpu<short, 4>},\r
1613         {pyrDown_gpu<int, 1>, pyrDown_gpu<int, 2>, pyrDown_gpu<int, 3>, pyrDown_gpu<int, 4>},\r
1614         {pyrDown_gpu<float, 1>, pyrDown_gpu<float, 2>, pyrDown_gpu<float, 3>, pyrDown_gpu<float, 4>},\r
1615     };\r
1616 \r
1617     CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);\r
1618 \r
1619     CV_Assert(borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP);\r
1620     int gpuBorderType;\r
1621     CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
1622 \r
1623     dst.create((src.rows + 1) / 2, (src.cols + 1) / 2, src.type());\r
1624 \r
1625     funcs[src.depth()][src.channels() - 1](src, dst, gpuBorderType, StreamAccessor::getStream(stream));\r
1626 }\r
1627 \r
1628 \r
1629 //////////////////////////////////////////////////////////////////////////////\r
1630 // pyrUp\r
1631 \r
1632 namespace cv { namespace gpu { namespace imgproc\r
1633 {\r
1634     template <typename T, int cn> void pyrUp_gpu(const DevMem2D& src, const DevMem2D& dst, int borderType, cudaStream_t stream);\r
1635 }}}\r
1636 \r
1637 void cv::gpu::pyrUp(const GpuMat& src, GpuMat& dst, int borderType, Stream& stream)\r
1638 {\r
1639     using namespace cv::gpu::imgproc;\r
1640 \r
1641     typedef void (*func_t)(const DevMem2D& src, const DevMem2D& dst, int borderType, cudaStream_t stream);\r
1642 \r
1643     static const func_t funcs[6][4] = \r
1644     {\r
1645         {pyrUp_gpu<uchar, 1>, pyrUp_gpu<uchar, 2>, pyrUp_gpu<uchar, 3>, pyrUp_gpu<uchar, 4>},\r
1646         {pyrUp_gpu<schar, 1>, pyrUp_gpu<schar, 2>, pyrUp_gpu<schar, 3>, pyrUp_gpu<schar, 4>},\r
1647         {pyrUp_gpu<ushort, 1>, pyrUp_gpu<ushort, 2>, pyrUp_gpu<ushort, 3>, pyrUp_gpu<ushort, 4>},\r
1648         {pyrUp_gpu<short, 1>, pyrUp_gpu<short, 2>, pyrUp_gpu<short, 3>, pyrUp_gpu<short, 4>},\r
1649         {pyrUp_gpu<int, 1>, pyrUp_gpu<int, 2>, pyrUp_gpu<int, 3>, pyrUp_gpu<int, 4>},\r
1650         {pyrUp_gpu<float, 1>, pyrUp_gpu<float, 2>, pyrUp_gpu<float, 3>, pyrUp_gpu<float, 4>},\r
1651     };\r
1652 \r
1653     CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);\r
1654 \r
1655     CV_Assert(borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP);\r
1656     int gpuBorderType;\r
1657     CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
1658 \r
1659     dst.create(src.rows*2, src.cols*2, src.type());\r
1660 \r
1661     funcs[src.depth()][src.channels() - 1](src, dst, gpuBorderType, StreamAccessor::getStream(stream));\r
1662 }\r
1663 \r
1664 \r
1665 //////////////////////////////////////////////////////////////////////////////\r
1666 // Canny\r
1667 \r
1668 cv::gpu::CannyBuf::CannyBuf(const GpuMat& dx_, const GpuMat& dy_) : dx(dx_), dy(dy_)\r
1669 {\r
1670     CV_Assert(dx_.type() == CV_32SC1 && dy_.type() == CV_32SC1 && dx_.size() == dy_.size());\r
1671 \r
1672     create(dx_.size(), -1);\r
1673 }\r
1674 \r
1675 void cv::gpu::CannyBuf::create(const Size& image_size, int apperture_size)\r
1676 {\r
1677     ensureSizeIsEnough(image_size, CV_32SC1, dx);\r
1678     ensureSizeIsEnough(image_size, CV_32SC1, dy);\r
1679 \r
1680     if (apperture_size == 3)\r
1681     {\r
1682         ensureSizeIsEnough(image_size, CV_32SC1, dx_buf);\r
1683         ensureSizeIsEnough(image_size, CV_32SC1, dy_buf);\r
1684     }\r
1685     else if(apperture_size > 0)\r
1686     {\r
1687         if (!filterDX)\r
1688             filterDX = createDerivFilter_GPU(CV_8UC1, CV_32S, 1, 0, apperture_size, BORDER_REPLICATE);\r
1689         if (!filterDY)\r
1690             filterDY = createDerivFilter_GPU(CV_8UC1, CV_32S, 0, 1, apperture_size, BORDER_REPLICATE);\r
1691     }\r
1692 \r
1693     ensureSizeIsEnough(image_size.height + 2, image_size.width + 2, CV_32FC1, edgeBuf);\r
1694 \r
1695     ensureSizeIsEnough(1, image_size.width * image_size.height, CV_16UC2, trackBuf1);\r
1696     ensureSizeIsEnough(1, image_size.width * image_size.height, CV_16UC2, trackBuf2);\r
1697 }\r
1698 \r
1699 void cv::gpu::CannyBuf::release()\r
1700 {\r
1701     dx.release();\r
1702     dy.release();\r
1703     dx_buf.release();\r
1704     dy_buf.release();\r
1705     edgeBuf.release();\r
1706     trackBuf1.release();\r
1707     trackBuf2.release();\r
1708 }\r
1709 \r
1710 namespace cv { namespace gpu { namespace canny\r
1711 {    \r
1712     void calcSobelRowPass_gpu(PtrStep src, PtrStepi dx_buf, PtrStepi dy_buf, int rows, int cols);\r
1713 \r
1714     void calcMagnitude_gpu(PtrStepi dx_buf, PtrStepi dy_buf, PtrStepi dx, PtrStepi dy, PtrStepf mag, int rows, int cols, bool L2Grad);\r
1715     void calcMagnitude_gpu(PtrStepi dx, PtrStepi dy, PtrStepf mag, int rows, int cols, bool L2Grad);\r
1716 \r
1717     void calcMap_gpu(PtrStepi dx, PtrStepi dy, PtrStepf mag, PtrStepi map, int rows, int cols, float low_thresh, float high_thresh);\r
1718     \r
1719     void edgesHysteresisLocal_gpu(PtrStepi map, ushort2* st1, int rows, int cols);\r
1720 \r
1721     void edgesHysteresisGlobal_gpu(PtrStepi map, ushort2* st1, ushort2* st2, int rows, int cols);\r
1722 \r
1723     void getEdges_gpu(PtrStepi map, PtrStep dst, int rows, int cols);\r
1724 }}}\r
1725 \r
1726 namespace\r
1727 {\r
1728     void CannyCaller(CannyBuf& buf, GpuMat& dst, float low_thresh, float high_thresh)\r
1729     {\r
1730         using namespace cv::gpu::canny;\r
1731 \r
1732         calcMap_gpu(buf.dx, buf.dy, buf.edgeBuf, buf.edgeBuf, dst.rows, dst.cols, low_thresh, high_thresh);\r
1733         \r
1734         edgesHysteresisLocal_gpu(buf.edgeBuf, buf.trackBuf1.ptr<ushort2>(), dst.rows, dst.cols);\r
1735         \r
1736         edgesHysteresisGlobal_gpu(buf.edgeBuf, buf.trackBuf1.ptr<ushort2>(), buf.trackBuf2.ptr<ushort2>(), dst.rows, dst.cols);\r
1737         \r
1738         getEdges_gpu(buf.edgeBuf, dst, dst.rows, dst.cols);\r
1739     }\r
1740 }\r
1741 \r
1742 void cv::gpu::Canny(const GpuMat& src, GpuMat& dst, double low_thresh, double high_thresh, int apperture_size, bool L2gradient)\r
1743 {\r
1744     CannyBuf buf(src.size(), apperture_size);\r
1745     Canny(src, buf, dst, low_thresh, high_thresh, apperture_size, L2gradient);\r
1746 }\r
1747 \r
1748 void cv::gpu::Canny(const GpuMat& src, CannyBuf& buf, GpuMat& dst, double low_thresh, double high_thresh, int apperture_size, bool L2gradient)\r
1749 {\r
1750     using namespace cv::gpu::canny;\r
1751 \r
1752     CV_Assert(TargetArchs::builtWith(SHARED_ATOMICS) && DeviceInfo().supports(SHARED_ATOMICS));\r
1753     CV_Assert(src.type() == CV_8UC1);\r
1754 \r
1755     if( low_thresh > high_thresh )\r
1756         std::swap( low_thresh, high_thresh);\r
1757 \r
1758     dst.create(src.size(), CV_8U);\r
1759     dst.setTo(Scalar::all(0));\r
1760     \r
1761     buf.create(src.size(), apperture_size);\r
1762     buf.edgeBuf.setTo(Scalar::all(0));\r
1763 \r
1764     if (apperture_size == 3)\r
1765     {\r
1766         calcSobelRowPass_gpu(src, buf.dx_buf, buf.dy_buf, src.rows, src.cols);\r
1767 \r
1768         calcMagnitude_gpu(buf.dx_buf, buf.dy_buf, buf.dx, buf.dy, buf.edgeBuf, src.rows, src.cols, L2gradient);\r
1769     }\r
1770     else\r
1771     {\r
1772         buf.filterDX->apply(src, buf.dx, Rect(0, 0, src.cols, src.rows));\r
1773         buf.filterDY->apply(src, buf.dy, Rect(0, 0, src.cols, src.rows));\r
1774 \r
1775         calcMagnitude_gpu(buf.dx, buf.dy, buf.edgeBuf, src.rows, src.cols, L2gradient);\r
1776     }\r
1777 \r
1778     CannyCaller(buf, dst, static_cast<float>(low_thresh), static_cast<float>(high_thresh));\r
1779 }\r
1780 \r
1781 void cv::gpu::Canny(const GpuMat& dx, const GpuMat& dy, GpuMat& dst, double low_thresh, double high_thresh, bool L2gradient)\r
1782 {\r
1783     CannyBuf buf(dx, dy);\r
1784     Canny(dx, dy, buf, dst, low_thresh, high_thresh, L2gradient);\r
1785 }\r
1786 \r
1787 void cv::gpu::Canny(const GpuMat& dx, const GpuMat& dy, CannyBuf& buf, GpuMat& dst, double low_thresh, double high_thresh, bool L2gradient)\r
1788 {\r
1789     using namespace cv::gpu::canny;\r
1790 \r
1791     CV_Assert(TargetArchs::builtWith(SHARED_ATOMICS) && DeviceInfo().supports(SHARED_ATOMICS));\r
1792     CV_Assert(dx.type() == CV_32SC1 && dy.type() == CV_32SC1 && dx.size() == dy.size());\r
1793 \r
1794     if( low_thresh > high_thresh )\r
1795         std::swap( low_thresh, high_thresh);\r
1796 \r
1797     dst.create(dx.size(), CV_8U);\r
1798     dst.setTo(Scalar::all(0));\r
1799     \r
1800     buf.dx = dx; buf.dy = dy;\r
1801     buf.create(dx.size(), -1);\r
1802     buf.edgeBuf.setTo(Scalar::all(0));\r
1803 \r
1804     calcMagnitude_gpu(dx, dy, buf.edgeBuf, dx.rows, dx.cols, L2gradient);\r
1805 \r
1806     CannyCaller(buf, dst, static_cast<float>(low_thresh), static_cast<float>(high_thresh));\r
1807 }\r
1808 \r
1809 #endif /* !defined (HAVE_CUDA) */\r
1810 \r
1811 \r