1 /*M///////////////////////////////////////////////////////////////////////////////////////
2 //
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4 //
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install,
7 // copy or use the software.
8 //
9 //
10 // License Agreement
11 // For Open Source Computer Vision Library
12 //
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
15 // Third party copyrights are property of their respective owners.
16 //
17 // Redistribution and use in source and binary forms, with or without modification,
18 // are permitted provided that the following conditions are met:
19 //
20 // * Redistribution's of source code must retain the above copyright notice,
21 // this list of conditions and the following disclaimer.
22 //
23 // * Redistribution's in binary form must reproduce the above copyright notice,
24 // this list of conditions and the following disclaimer in the documentation
25 // and/or other materials provided with the distribution.
26 //
27 // * The name of the copyright holders may not be used to endorse or promote products
28 // derived from this software without specific prior written permission.
29 //
30 // This software is provided by the copyright holders and contributors "as is" and
31 // any express or implied warranties, including, but not limited to, the implied
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
34 // indirect, incidental, special, exemplary, or consequential damages
35 // (including, but not limited to, procurement of substitute goods or services;
36 // loss of use, data, or profits; or business interruption) however caused
37 // and on any theory of liability, whether in contract, strict liability,
38 // or tort (including negligence or otherwise) arising in any way out of
39 // the use of this software, even if advised of the possibility of such damage.
40 //
41 //M*/
42
43 #include "precomp.hpp"
44
45 using namespace cv;
46 using namespace cv::cuda;
47
48 #if !defined HAVE_CUDA || defined(CUDA_DISABLER)
49
warpAffine(InputArray,OutputArray,InputArray,Size,int,int,Scalar,Stream &)50 void cv::cuda::warpAffine(InputArray, OutputArray, InputArray, Size, int, int, Scalar, Stream&) { throw_no_cuda(); }
buildWarpAffineMaps(InputArray,bool,Size,OutputArray,OutputArray,Stream &)51 void cv::cuda::buildWarpAffineMaps(InputArray, bool, Size, OutputArray, OutputArray, Stream&) { throw_no_cuda(); }
52
warpPerspective(InputArray,OutputArray,InputArray,Size,int,int,Scalar,Stream &)53 void cv::cuda::warpPerspective(InputArray, OutputArray, InputArray, Size, int, int, Scalar, Stream&) { throw_no_cuda(); }
buildWarpPerspectiveMaps(InputArray,bool,Size,OutputArray,OutputArray,Stream &)54 void cv::cuda::buildWarpPerspectiveMaps(InputArray, bool, Size, OutputArray, OutputArray, Stream&) { throw_no_cuda(); }
55
rotate(InputArray,OutputArray,Size,double,double,double,int,Stream &)56 void cv::cuda::rotate(InputArray, OutputArray, Size, double, double, double, int, Stream&) { throw_no_cuda(); }
57
58 #else // HAVE_CUDA
59
60 namespace cv { namespace cuda { namespace device
61 {
62 namespace imgproc
63 {
64 void buildWarpAffineMaps_gpu(float coeffs[2 * 3], PtrStepSzf xmap, PtrStepSzf ymap, cudaStream_t stream);
65
66 template <typename T>
67 void warpAffine_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[2 * 3], PtrStepSzb dst, int interpolation,
68 int borderMode, const float* borderValue, cudaStream_t stream, bool cc20);
69
70 void buildWarpPerspectiveMaps_gpu(float coeffs[3 * 3], PtrStepSzf xmap, PtrStepSzf ymap, cudaStream_t stream);
71
72 template <typename T>
73 void warpPerspective_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[3 * 3], PtrStepSzb dst, int interpolation,
74 int borderMode, const float* borderValue, cudaStream_t stream, bool cc20);
75 }
76 }}}
77
buildWarpAffineMaps(InputArray _M,bool inverse,Size dsize,OutputArray _xmap,OutputArray _ymap,Stream & stream)78 void cv::cuda::buildWarpAffineMaps(InputArray _M, bool inverse, Size dsize, OutputArray _xmap, OutputArray _ymap, Stream& stream)
79 {
80 using namespace cv::cuda::device::imgproc;
81
82 Mat M = _M.getMat();
83
84 CV_Assert( M.rows == 2 && M.cols == 3 );
85
86 _xmap.create(dsize, CV_32FC1);
87 _ymap.create(dsize, CV_32FC1);
88
89 GpuMat xmap = _xmap.getGpuMat();
90 GpuMat ymap = _ymap.getGpuMat();
91
92 float coeffs[2 * 3];
93 Mat coeffsMat(2, 3, CV_32F, (void*)coeffs);
94
95 if (inverse)
96 M.convertTo(coeffsMat, coeffsMat.type());
97 else
98 {
99 cv::Mat iM;
100 invertAffineTransform(M, iM);
101 iM.convertTo(coeffsMat, coeffsMat.type());
102 }
103
104 buildWarpAffineMaps_gpu(coeffs, xmap, ymap, StreamAccessor::getStream(stream));
105 }
106
buildWarpPerspectiveMaps(InputArray _M,bool inverse,Size dsize,OutputArray _xmap,OutputArray _ymap,Stream & stream)107 void cv::cuda::buildWarpPerspectiveMaps(InputArray _M, bool inverse, Size dsize, OutputArray _xmap, OutputArray _ymap, Stream& stream)
108 {
109 using namespace cv::cuda::device::imgproc;
110
111 Mat M = _M.getMat();
112
113 CV_Assert( M.rows == 3 && M.cols == 3 );
114
115 _xmap.create(dsize, CV_32FC1);
116 _ymap.create(dsize, CV_32FC1);
117
118 GpuMat xmap = _xmap.getGpuMat();
119 GpuMat ymap = _ymap.getGpuMat();
120
121 float coeffs[3 * 3];
122 Mat coeffsMat(3, 3, CV_32F, (void*)coeffs);
123
124 if (inverse)
125 M.convertTo(coeffsMat, coeffsMat.type());
126 else
127 {
128 cv::Mat iM;
129 invert(M, iM);
130 iM.convertTo(coeffsMat, coeffsMat.type());
131 }
132
133 buildWarpPerspectiveMaps_gpu(coeffs, xmap, ymap, StreamAccessor::getStream(stream));
134 }
135
136 namespace
137 {
138 template <int DEPTH> struct NppWarpFunc
139 {
140 typedef typename NPPTypeTraits<DEPTH>::npp_type npp_type;
141
142 typedef NppStatus (*func_t)(const npp_type* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, npp_type* pDst,
143 int dstStep, NppiRect dstRoi, const double coeffs[][3],
144 int interpolation);
145 };
146
147 template <int DEPTH, typename NppWarpFunc<DEPTH>::func_t func> struct NppWarp
148 {
149 typedef typename NppWarpFunc<DEPTH>::npp_type npp_type;
150
call__anon78384e2c0111::NppWarp151 static void call(const cv::cuda::GpuMat& src, cv::cuda::GpuMat& dst, double coeffs[][3], int interpolation, cudaStream_t stream)
152 {
153 static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};
154
155 NppiSize srcsz;
156 srcsz.height = src.rows;
157 srcsz.width = src.cols;
158
159 NppiRect srcroi;
160 srcroi.x = 0;
161 srcroi.y = 0;
162 srcroi.height = src.rows;
163 srcroi.width = src.cols;
164
165 NppiRect dstroi;
166 dstroi.x = 0;
167 dstroi.y = 0;
168 dstroi.height = dst.rows;
169 dstroi.width = dst.cols;
170
171 cv::cuda::NppStreamHandler h(stream);
172
173 nppSafeCall( func(src.ptr<npp_type>(), srcsz, static_cast<int>(src.step), srcroi,
174 dst.ptr<npp_type>(), static_cast<int>(dst.step), dstroi,
175 coeffs, npp_inter[interpolation]) );
176
177 if (stream == 0)
178 cudaSafeCall( cudaDeviceSynchronize() );
179 }
180 };
181 }
182
warpAffine(InputArray _src,OutputArray _dst,InputArray _M,Size dsize,int flags,int borderMode,Scalar borderValue,Stream & stream)183 void cv::cuda::warpAffine(InputArray _src, OutputArray _dst, InputArray _M, Size dsize, int flags, int borderMode, Scalar borderValue, Stream& stream)
184 {
185 GpuMat src = _src.getGpuMat();
186 Mat M = _M.getMat();
187
188 CV_Assert( M.rows == 2 && M.cols == 3 );
189
190 const int interpolation = flags & INTER_MAX;
191
192 CV_Assert( src.depth() <= CV_32F && src.channels() <= 4 );
193 CV_Assert( interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC );
194 CV_Assert( borderMode == BORDER_REFLECT101 || borderMode == BORDER_REPLICATE || borderMode == BORDER_CONSTANT || borderMode == BORDER_REFLECT || borderMode == BORDER_WRAP );
195
196 _dst.create(dsize, src.type());
197 GpuMat dst = _dst.getGpuMat();
198
199 Size wholeSize;
200 Point ofs;
201 src.locateROI(wholeSize, ofs);
202
203 static const bool useNppTab[6][4][3] =
204 {
205 {
206 {false, false, true},
207 {false, false, false},
208 {false, true, true},
209 {false, false, false}
210 },
211 {
212 {false, false, false},
213 {false, false, false},
214 {false, false, false},
215 {false, false, false}
216 },
217 {
218 {false, true, true},
219 {false, false, false},
220 {false, true, true},
221 {false, false, false}
222 },
223 {
224 {false, false, false},
225 {false, false, false},
226 {false, false, false},
227 {false, false, false}
228 },
229 {
230 {false, true, true},
231 {false, false, false},
232 {false, true, true},
233 {false, false, true}
234 },
235 {
236 {false, true, true},
237 {false, false, false},
238 {false, true, true},
239 {false, false, true}
240 }
241 };
242
243 bool useNpp = borderMode == BORDER_CONSTANT && ofs.x == 0 && ofs.y == 0 && useNppTab[src.depth()][src.channels() - 1][interpolation];
244 // NPP bug on float data
245 useNpp = useNpp && src.depth() != CV_32F;
246
247 if (useNpp)
248 {
249 typedef void (*func_t)(const cv::cuda::GpuMat& src, cv::cuda::GpuMat& dst, double coeffs[][3], int flags, cudaStream_t stream);
250
251 static const func_t funcs[2][6][4] =
252 {
253 {
254 {NppWarp<CV_8U, nppiWarpAffine_8u_C1R>::call, 0, NppWarp<CV_8U, nppiWarpAffine_8u_C3R>::call, NppWarp<CV_8U, nppiWarpAffine_8u_C4R>::call},
255 {0, 0, 0, 0},
256 {NppWarp<CV_16U, nppiWarpAffine_16u_C1R>::call, 0, NppWarp<CV_16U, nppiWarpAffine_16u_C3R>::call, NppWarp<CV_16U, nppiWarpAffine_16u_C4R>::call},
257 {0, 0, 0, 0},
258 {NppWarp<CV_32S, nppiWarpAffine_32s_C1R>::call, 0, NppWarp<CV_32S, nppiWarpAffine_32s_C3R>::call, NppWarp<CV_32S, nppiWarpAffine_32s_C4R>::call},
259 {NppWarp<CV_32F, nppiWarpAffine_32f_C1R>::call, 0, NppWarp<CV_32F, nppiWarpAffine_32f_C3R>::call, NppWarp<CV_32F, nppiWarpAffine_32f_C4R>::call}
260 },
261 {
262 {NppWarp<CV_8U, nppiWarpAffineBack_8u_C1R>::call, 0, NppWarp<CV_8U, nppiWarpAffineBack_8u_C3R>::call, NppWarp<CV_8U, nppiWarpAffineBack_8u_C4R>::call},
263 {0, 0, 0, 0},
264 {NppWarp<CV_16U, nppiWarpAffineBack_16u_C1R>::call, 0, NppWarp<CV_16U, nppiWarpAffineBack_16u_C3R>::call, NppWarp<CV_16U, nppiWarpAffineBack_16u_C4R>::call},
265 {0, 0, 0, 0},
266 {NppWarp<CV_32S, nppiWarpAffineBack_32s_C1R>::call, 0, NppWarp<CV_32S, nppiWarpAffineBack_32s_C3R>::call, NppWarp<CV_32S, nppiWarpAffineBack_32s_C4R>::call},
267 {NppWarp<CV_32F, nppiWarpAffineBack_32f_C1R>::call, 0, NppWarp<CV_32F, nppiWarpAffineBack_32f_C3R>::call, NppWarp<CV_32F, nppiWarpAffineBack_32f_C4R>::call}
268 }
269 };
270
271 dst.setTo(borderValue, stream);
272
273 double coeffs[2][3];
274 Mat coeffsMat(2, 3, CV_64F, (void*)coeffs);
275 M.convertTo(coeffsMat, coeffsMat.type());
276
277 const func_t func = funcs[(flags & WARP_INVERSE_MAP) != 0][src.depth()][src.channels() - 1];
278 CV_Assert(func != 0);
279
280 func(src, dst, coeffs, interpolation, StreamAccessor::getStream(stream));
281 }
282 else
283 {
284 using namespace cv::cuda::device::imgproc;
285
286 typedef void (*func_t)(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[2 * 3], PtrStepSzb dst, int interpolation,
287 int borderMode, const float* borderValue, cudaStream_t stream, bool cc20);
288
289 static const func_t funcs[6][4] =
290 {
291 {warpAffine_gpu<uchar> , 0 /*warpAffine_gpu<uchar2>*/ , warpAffine_gpu<uchar3> , warpAffine_gpu<uchar4> },
292 {0 /*warpAffine_gpu<schar>*/, 0 /*warpAffine_gpu<char2>*/ , 0 /*warpAffine_gpu<char3>*/, 0 /*warpAffine_gpu<char4>*/},
293 {warpAffine_gpu<ushort> , 0 /*warpAffine_gpu<ushort2>*/, warpAffine_gpu<ushort3> , warpAffine_gpu<ushort4> },
294 {warpAffine_gpu<short> , 0 /*warpAffine_gpu<short2>*/ , warpAffine_gpu<short3> , warpAffine_gpu<short4> },
295 {0 /*warpAffine_gpu<int>*/ , 0 /*warpAffine_gpu<int2>*/ , 0 /*warpAffine_gpu<int3>*/ , 0 /*warpAffine_gpu<int4>*/ },
296 {warpAffine_gpu<float> , 0 /*warpAffine_gpu<float2>*/ , warpAffine_gpu<float3> , warpAffine_gpu<float4> }
297 };
298
299 const func_t func = funcs[src.depth()][src.channels() - 1];
300 CV_Assert(func != 0);
301
302 float coeffs[2 * 3];
303 Mat coeffsMat(2, 3, CV_32F, (void*)coeffs);
304
305 if (flags & WARP_INVERSE_MAP)
306 M.convertTo(coeffsMat, coeffsMat.type());
307 else
308 {
309 cv::Mat iM;
310 invertAffineTransform(M, iM);
311 iM.convertTo(coeffsMat, coeffsMat.type());
312 }
313
314 Scalar_<float> borderValueFloat;
315 borderValueFloat = borderValue;
316
317 func(src, PtrStepSzb(wholeSize.height, wholeSize.width, src.datastart, src.step), ofs.x, ofs.y, coeffs,
318 dst, interpolation, borderMode, borderValueFloat.val, StreamAccessor::getStream(stream), deviceSupports(FEATURE_SET_COMPUTE_20));
319 }
320 }
321
warpPerspective(InputArray _src,OutputArray _dst,InputArray _M,Size dsize,int flags,int borderMode,Scalar borderValue,Stream & stream)322 void cv::cuda::warpPerspective(InputArray _src, OutputArray _dst, InputArray _M, Size dsize, int flags, int borderMode, Scalar borderValue, Stream& stream)
323 {
324 GpuMat src = _src.getGpuMat();
325 Mat M = _M.getMat();
326
327 CV_Assert( M.rows == 3 && M.cols == 3 );
328
329 const int interpolation = flags & INTER_MAX;
330
331 CV_Assert( src.depth() <= CV_32F && src.channels() <= 4 );
332 CV_Assert( interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC );
333 CV_Assert( borderMode == BORDER_REFLECT101 || borderMode == BORDER_REPLICATE || borderMode == BORDER_CONSTANT || borderMode == BORDER_REFLECT || borderMode == BORDER_WRAP) ;
334
335 _dst.create(dsize, src.type());
336 GpuMat dst = _dst.getGpuMat();
337
338 Size wholeSize;
339 Point ofs;
340 src.locateROI(wholeSize, ofs);
341
342 static const bool useNppTab[6][4][3] =
343 {
344 {
345 {false, false, true},
346 {false, false, false},
347 {false, true, true},
348 {false, false, false}
349 },
350 {
351 {false, false, false},
352 {false, false, false},
353 {false, false, false},
354 {false, false, false}
355 },
356 {
357 {false, true, true},
358 {false, false, false},
359 {false, true, true},
360 {false, false, false}
361 },
362 {
363 {false, false, false},
364 {false, false, false},
365 {false, false, false},
366 {false, false, false}
367 },
368 {
369 {false, true, true},
370 {false, false, false},
371 {false, true, true},
372 {false, false, true}
373 },
374 {
375 {false, true, true},
376 {false, false, false},
377 {false, true, true},
378 {false, false, true}
379 }
380 };
381
382 bool useNpp = borderMode == BORDER_CONSTANT && ofs.x == 0 && ofs.y == 0 && useNppTab[src.depth()][src.channels() - 1][interpolation];
383 // NPP bug on float data
384 useNpp = useNpp && src.depth() != CV_32F;
385
386 if (useNpp)
387 {
388 typedef void (*func_t)(const cv::cuda::GpuMat& src, cv::cuda::GpuMat& dst, double coeffs[][3], int flags, cudaStream_t stream);
389
390 static const func_t funcs[2][6][4] =
391 {
392 {
393 {NppWarp<CV_8U, nppiWarpPerspective_8u_C1R>::call, 0, NppWarp<CV_8U, nppiWarpPerspective_8u_C3R>::call, NppWarp<CV_8U, nppiWarpPerspective_8u_C4R>::call},
394 {0, 0, 0, 0},
395 {NppWarp<CV_16U, nppiWarpPerspective_16u_C1R>::call, 0, NppWarp<CV_16U, nppiWarpPerspective_16u_C3R>::call, NppWarp<CV_16U, nppiWarpPerspective_16u_C4R>::call},
396 {0, 0, 0, 0},
397 {NppWarp<CV_32S, nppiWarpPerspective_32s_C1R>::call, 0, NppWarp<CV_32S, nppiWarpPerspective_32s_C3R>::call, NppWarp<CV_32S, nppiWarpPerspective_32s_C4R>::call},
398 {NppWarp<CV_32F, nppiWarpPerspective_32f_C1R>::call, 0, NppWarp<CV_32F, nppiWarpPerspective_32f_C3R>::call, NppWarp<CV_32F, nppiWarpPerspective_32f_C4R>::call}
399 },
400 {
401 {NppWarp<CV_8U, nppiWarpPerspectiveBack_8u_C1R>::call, 0, NppWarp<CV_8U, nppiWarpPerspectiveBack_8u_C3R>::call, NppWarp<CV_8U, nppiWarpPerspectiveBack_8u_C4R>::call},
402 {0, 0, 0, 0},
403 {NppWarp<CV_16U, nppiWarpPerspectiveBack_16u_C1R>::call, 0, NppWarp<CV_16U, nppiWarpPerspectiveBack_16u_C3R>::call, NppWarp<CV_16U, nppiWarpPerspectiveBack_16u_C4R>::call},
404 {0, 0, 0, 0},
405 {NppWarp<CV_32S, nppiWarpPerspectiveBack_32s_C1R>::call, 0, NppWarp<CV_32S, nppiWarpPerspectiveBack_32s_C3R>::call, NppWarp<CV_32S, nppiWarpPerspectiveBack_32s_C4R>::call},
406 {NppWarp<CV_32F, nppiWarpPerspectiveBack_32f_C1R>::call, 0, NppWarp<CV_32F, nppiWarpPerspectiveBack_32f_C3R>::call, NppWarp<CV_32F, nppiWarpPerspectiveBack_32f_C4R>::call}
407 }
408 };
409
410 dst.setTo(borderValue, stream);
411
412 double coeffs[3][3];
413 Mat coeffsMat(3, 3, CV_64F, (void*)coeffs);
414 M.convertTo(coeffsMat, coeffsMat.type());
415
416 const func_t func = funcs[(flags & WARP_INVERSE_MAP) != 0][src.depth()][src.channels() - 1];
417 CV_Assert(func != 0);
418
419 func(src, dst, coeffs, interpolation, StreamAccessor::getStream(stream));
420 }
421 else
422 {
423 using namespace cv::cuda::device::imgproc;
424
425 typedef void (*func_t)(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[2 * 3], PtrStepSzb dst, int interpolation,
426 int borderMode, const float* borderValue, cudaStream_t stream, bool cc20);
427
428 static const func_t funcs[6][4] =
429 {
430 {warpPerspective_gpu<uchar> , 0 /*warpPerspective_gpu<uchar2>*/ , warpPerspective_gpu<uchar3> , warpPerspective_gpu<uchar4> },
431 {0 /*warpPerspective_gpu<schar>*/, 0 /*warpPerspective_gpu<char2>*/ , 0 /*warpPerspective_gpu<char3>*/, 0 /*warpPerspective_gpu<char4>*/},
432 {warpPerspective_gpu<ushort> , 0 /*warpPerspective_gpu<ushort2>*/, warpPerspective_gpu<ushort3> , warpPerspective_gpu<ushort4> },
433 {warpPerspective_gpu<short> , 0 /*warpPerspective_gpu<short2>*/ , warpPerspective_gpu<short3> , warpPerspective_gpu<short4> },
434 {0 /*warpPerspective_gpu<int>*/ , 0 /*warpPerspective_gpu<int2>*/ , 0 /*warpPerspective_gpu<int3>*/ , 0 /*warpPerspective_gpu<int4>*/ },
435 {warpPerspective_gpu<float> , 0 /*warpPerspective_gpu<float2>*/ , warpPerspective_gpu<float3> , warpPerspective_gpu<float4> }
436 };
437
438 const func_t func = funcs[src.depth()][src.channels() - 1];
439 CV_Assert(func != 0);
440
441 float coeffs[3 * 3];
442 Mat coeffsMat(3, 3, CV_32F, (void*)coeffs);
443
444 if (flags & WARP_INVERSE_MAP)
445 M.convertTo(coeffsMat, coeffsMat.type());
446 else
447 {
448 cv::Mat iM;
449 invert(M, iM);
450 iM.convertTo(coeffsMat, coeffsMat.type());
451 }
452
453 Scalar_<float> borderValueFloat;
454 borderValueFloat = borderValue;
455
456 func(src, PtrStepSzb(wholeSize.height, wholeSize.width, src.datastart, src.step), ofs.x, ofs.y, coeffs,
457 dst, interpolation, borderMode, borderValueFloat.val, StreamAccessor::getStream(stream), deviceSupports(FEATURE_SET_COMPUTE_20));
458 }
459 }
460
461 ////////////////////////////////////////////////////////////////////////
462 // rotate
463
464 namespace
465 {
466 template <int DEPTH> struct NppRotateFunc
467 {
468 typedef typename NPPTypeTraits<DEPTH>::npp_type npp_type;
469
470 typedef NppStatus (*func_t)(const npp_type* pSrc, NppiSize oSrcSize, int nSrcStep, NppiRect oSrcROI,
471 npp_type* pDst, int nDstStep, NppiRect oDstROI,
472 double nAngle, double nShiftX, double nShiftY, int eInterpolation);
473 };
474
475 template <int DEPTH, typename NppRotateFunc<DEPTH>::func_t func> struct NppRotate
476 {
477 typedef typename NppRotateFunc<DEPTH>::npp_type npp_type;
478
call__anon78384e2c0211::NppRotate479 static void call(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, cudaStream_t stream)
480 {
481 (void)dsize;
482 static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};
483
484 NppStreamHandler h(stream);
485
486 NppiSize srcsz;
487 srcsz.height = src.rows;
488 srcsz.width = src.cols;
489 NppiRect srcroi;
490 srcroi.x = srcroi.y = 0;
491 srcroi.height = src.rows;
492 srcroi.width = src.cols;
493 NppiRect dstroi;
494 dstroi.x = dstroi.y = 0;
495 dstroi.height = dst.rows;
496 dstroi.width = dst.cols;
497
498 nppSafeCall( func(src.ptr<npp_type>(), srcsz, static_cast<int>(src.step), srcroi,
499 dst.ptr<npp_type>(), static_cast<int>(dst.step), dstroi, angle, xShift, yShift, npp_inter[interpolation]) );
500
501 if (stream == 0)
502 cudaSafeCall( cudaDeviceSynchronize() );
503 }
504 };
505 }
506
rotate(InputArray _src,OutputArray _dst,Size dsize,double angle,double xShift,double yShift,int interpolation,Stream & stream)507 void cv::cuda::rotate(InputArray _src, OutputArray _dst, Size dsize, double angle, double xShift, double yShift, int interpolation, Stream& stream)
508 {
509 typedef void (*func_t)(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, cudaStream_t stream);
510 static const func_t funcs[6][4] =
511 {
512 {NppRotate<CV_8U, nppiRotate_8u_C1R>::call, 0, NppRotate<CV_8U, nppiRotate_8u_C3R>::call, NppRotate<CV_8U, nppiRotate_8u_C4R>::call},
513 {0,0,0,0},
514 {NppRotate<CV_16U, nppiRotate_16u_C1R>::call, 0, NppRotate<CV_16U, nppiRotate_16u_C3R>::call, NppRotate<CV_16U, nppiRotate_16u_C4R>::call},
515 {0,0,0,0},
516 {0,0,0,0},
517 {NppRotate<CV_32F, nppiRotate_32f_C1R>::call, 0, NppRotate<CV_32F, nppiRotate_32f_C3R>::call, NppRotate<CV_32F, nppiRotate_32f_C4R>::call}
518 };
519
520 GpuMat src = _src.getGpuMat();
521
522 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
523 CV_Assert( src.channels() == 1 || src.channels() == 3 || src.channels() == 4 );
524 CV_Assert( interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC );
525
526 _dst.create(dsize, src.type());
527 GpuMat dst = _dst.getGpuMat();
528
529 dst.setTo(Scalar::all(0), stream);
530
531 funcs[src.depth()][src.channels() - 1](src, dst, dsize, angle, xShift, yShift, interpolation, StreamAccessor::getStream(stream));
532 }
533
534 #endif // HAVE_CUDA
535