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 "opencv2/opencv_modules.hpp" 44 45 #ifndef HAVE_OPENCV_CUDEV 46 47 #error "opencv_cudev is required" 48 49 #else 50 51 #include "opencv2/core/cuda.hpp" 52 #include "opencv2/cudev.hpp" 53 54 using namespace cv; 55 using namespace cv::cuda; 56 using namespace cv::cudev; 57 58 namespace 59 { 60 class DefaultAllocator : public GpuMat::Allocator 61 { 62 public: 63 bool allocate(GpuMat* mat, int rows, int cols, size_t elemSize); 64 void free(GpuMat* mat); 65 }; 66 67 bool DefaultAllocator::allocate(GpuMat* mat, int rows, int cols, size_t elemSize) 68 { 69 if (rows > 1 && cols > 1) 70 { 71 CV_CUDEV_SAFE_CALL( cudaMallocPitch(&mat->data, &mat->step, elemSize * cols, rows) ); 72 } 73 else 74 { 75 // Single row or single column must be continuous 76 CV_CUDEV_SAFE_CALL( cudaMalloc(&mat->data, elemSize * cols * rows) ); 77 mat->step = elemSize * cols; 78 } 79 80 mat->refcount = (int*) fastMalloc(sizeof(int)); 81 82 return true; 83 } 84 85 void DefaultAllocator::free(GpuMat* mat) 86 { 87 cudaFree(mat->datastart); 88 fastFree(mat->refcount); 89 } 90 91 DefaultAllocator cudaDefaultAllocator; 92 GpuMat::Allocator* g_defaultAllocator = &cudaDefaultAllocator; 93 } 94 95 GpuMat::Allocator* cv::cuda::GpuMat::defaultAllocator() 96 { 97 return g_defaultAllocator; 98 } 99 100 void cv::cuda::GpuMat::setDefaultAllocator(Allocator* allocator) 101 { 102 CV_Assert( allocator != 0 ); 103 g_defaultAllocator = allocator; 104 } 105 106 ///////////////////////////////////////////////////// 107 /// create 108 109 void cv::cuda::GpuMat::create(int _rows, int _cols, int _type) 110 { 111 CV_DbgAssert( _rows >= 0 && _cols >= 0 ); 112 113 _type &= Mat::TYPE_MASK; 114 115 if (rows == _rows && cols == _cols && type() == _type && data) 116 return; 117 118 if (data) 119 release(); 120 121 if (_rows > 0 && _cols > 0) 122 { 123 flags = Mat::MAGIC_VAL + _type; 124 rows = _rows; 125 cols = _cols; 126 127 const size_t esz = elemSize(); 128 129 bool allocSuccess = allocator->allocate(this, rows, cols, esz); 130 131 if (!allocSuccess) 132 { 133 // custom allocator fails, try default allocator 134 allocator = defaultAllocator(); 135 allocSuccess = allocator->allocate(this, rows, cols, esz); 136 CV_Assert( allocSuccess ); 137 } 138 139 if (esz * cols == step) 140 flags |= Mat::CONTINUOUS_FLAG; 141 142 int64 _nettosize = static_cast<int64>(step) * rows; 143 size_t nettosize = static_cast<size_t>(_nettosize); 144 145 datastart = data; 146 dataend = data + nettosize; 147 148 if (refcount) 149 *refcount = 1; 150 } 151 } 152 153 ///////////////////////////////////////////////////// 154 /// release 155 156 void cv::cuda::GpuMat::release() 157 { 158 CV_DbgAssert( allocator != 0 ); 159 160 if (refcount && CV_XADD(refcount, -1) == 1) 161 allocator->free(this); 162 163 dataend = data = datastart = 0; 164 step = rows = cols = 0; 165 refcount = 0; 166 } 167 168 ///////////////////////////////////////////////////// 169 /// upload 170 171 void cv::cuda::GpuMat::upload(InputArray arr) 172 { 173 Mat mat = arr.getMat(); 174 175 CV_DbgAssert( !mat.empty() ); 176 177 create(mat.size(), mat.type()); 178 179 CV_CUDEV_SAFE_CALL( cudaMemcpy2D(data, step, mat.data, mat.step, cols * elemSize(), rows, cudaMemcpyHostToDevice) ); 180 } 181 182 void cv::cuda::GpuMat::upload(InputArray arr, Stream& _stream) 183 { 184 Mat mat = arr.getMat(); 185 186 CV_DbgAssert( !mat.empty() ); 187 188 create(mat.size(), mat.type()); 189 190 cudaStream_t stream = StreamAccessor::getStream(_stream); 191 CV_CUDEV_SAFE_CALL( cudaMemcpy2DAsync(data, step, mat.data, mat.step, cols * elemSize(), rows, cudaMemcpyHostToDevice, stream) ); 192 } 193 194 ///////////////////////////////////////////////////// 195 /// download 196 197 void cv::cuda::GpuMat::download(OutputArray _dst) const 198 { 199 CV_DbgAssert( !empty() ); 200 201 _dst.create(size(), type()); 202 Mat dst = _dst.getMat(); 203 204 CV_CUDEV_SAFE_CALL( cudaMemcpy2D(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToHost) ); 205 } 206 207 void cv::cuda::GpuMat::download(OutputArray _dst, Stream& _stream) const 208 { 209 CV_DbgAssert( !empty() ); 210 211 _dst.create(size(), type()); 212 Mat dst = _dst.getMat(); 213 214 cudaStream_t stream = StreamAccessor::getStream(_stream); 215 CV_CUDEV_SAFE_CALL( cudaMemcpy2DAsync(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToHost, stream) ); 216 } 217 218 ///////////////////////////////////////////////////// 219 /// copyTo 220 221 void cv::cuda::GpuMat::copyTo(OutputArray _dst) const 222 { 223 CV_DbgAssert( !empty() ); 224 225 _dst.create(size(), type()); 226 GpuMat dst = _dst.getGpuMat(); 227 228 CV_CUDEV_SAFE_CALL( cudaMemcpy2D(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice) ); 229 } 230 231 void cv::cuda::GpuMat::copyTo(OutputArray _dst, Stream& _stream) const 232 { 233 CV_DbgAssert( !empty() ); 234 235 _dst.create(size(), type()); 236 GpuMat dst = _dst.getGpuMat(); 237 238 cudaStream_t stream = StreamAccessor::getStream(_stream); 239 CV_CUDEV_SAFE_CALL( cudaMemcpy2DAsync(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice, stream) ); 240 } 241 242 namespace 243 { 244 template <size_t size> struct CopyToPolicy : DefaultTransformPolicy 245 { 246 }; 247 template <> struct CopyToPolicy<4> : DefaultTransformPolicy 248 { 249 enum { 250 shift = 2 251 }; 252 }; 253 template <> struct CopyToPolicy<8> : DefaultTransformPolicy 254 { 255 enum { 256 shift = 1 257 }; 258 }; 259 260 template <typename T> 261 void copyWithMask(const GpuMat& src, const GpuMat& dst, const GpuMat& mask, Stream& stream) 262 { 263 gridTransformUnary_< CopyToPolicy<sizeof(typename VecTraits<T>::elem_type)> >(globPtr<T>(src), globPtr<T>(dst), identity<T>(), globPtr<uchar>(mask), stream); 264 } 265 } 266 267 void cv::cuda::GpuMat::copyTo(OutputArray _dst, InputArray _mask, Stream& stream) const 268 { 269 CV_DbgAssert( !empty() ); 270 CV_DbgAssert( depth() <= CV_64F && channels() <= 4 ); 271 272 GpuMat mask = _mask.getGpuMat(); 273 CV_DbgAssert( size() == mask.size() && mask.depth() == CV_8U && (mask.channels() == 1 || mask.channels() == channels()) ); 274 275 uchar* data0 = _dst.getGpuMat().data; 276 277 _dst.create(size(), type()); 278 GpuMat dst = _dst.getGpuMat(); 279 280 // do not leave dst uninitialized 281 if (dst.data != data0) 282 dst.setTo(Scalar::all(0), stream); 283 284 typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, const GpuMat& mask, Stream& stream); 285 static const func_t funcs[9][4] = 286 { 287 {0,0,0,0}, 288 {copyWithMask<uchar>, copyWithMask<uchar2>, copyWithMask<uchar3>, copyWithMask<uchar4>}, 289 {copyWithMask<ushort>, copyWithMask<ushort2>, copyWithMask<ushort3>, copyWithMask<ushort4>}, 290 {0,0,0,0}, 291 {copyWithMask<int>, copyWithMask<int2>, copyWithMask<int3>, copyWithMask<int4>}, 292 {0,0,0,0}, 293 {0,0,0,0}, 294 {0,0,0,0}, 295 {copyWithMask<double>, copyWithMask<double2>, copyWithMask<double3>, copyWithMask<double4>} 296 }; 297 298 if (mask.channels() == channels()) 299 { 300 const func_t func = funcs[elemSize1()][0]; 301 CV_DbgAssert( func != 0 ); 302 func(reshape(1), dst.reshape(1), mask.reshape(1), stream); 303 } 304 else 305 { 306 const func_t func = funcs[elemSize1()][channels() - 1]; 307 CV_DbgAssert( func != 0 ); 308 func(*this, dst, mask, stream); 309 } 310 } 311 312 ///////////////////////////////////////////////////// 313 /// setTo 314 315 namespace 316 { 317 template <typename T> 318 void setToWithOutMask(const GpuMat& mat, Scalar _scalar, Stream& stream) 319 { 320 Scalar_<typename VecTraits<T>::elem_type> scalar = _scalar; 321 gridTransformUnary(constantPtr(VecTraits<T>::make(scalar.val), mat.rows, mat.cols), globPtr<T>(mat), identity<T>(), stream); 322 } 323 324 template <typename T> 325 void setToWithMask(const GpuMat& mat, const GpuMat& mask, Scalar _scalar, Stream& stream) 326 { 327 Scalar_<typename VecTraits<T>::elem_type> scalar = _scalar; 328 gridTransformUnary(constantPtr(VecTraits<T>::make(scalar.val), mat.rows, mat.cols), globPtr<T>(mat), identity<T>(), globPtr<uchar>(mask), stream); 329 } 330 } 331 332 GpuMat& cv::cuda::GpuMat::setTo(Scalar value, Stream& stream) 333 { 334 CV_DbgAssert( !empty() ); 335 CV_DbgAssert( depth() <= CV_64F && channels() <= 4 ); 336 337 if (value[0] == 0.0 && value[1] == 0.0 && value[2] == 0.0 && value[3] == 0.0) 338 { 339 // Zero fill 340 341 if (stream) 342 CV_CUDEV_SAFE_CALL( cudaMemset2DAsync(data, step, 0, cols * elemSize(), rows, StreamAccessor::getStream(stream)) ); 343 else 344 CV_CUDEV_SAFE_CALL( cudaMemset2D(data, step, 0, cols * elemSize(), rows) ); 345 346 return *this; 347 } 348 349 if (depth() == CV_8U) 350 { 351 const int cn = channels(); 352 353 if (cn == 1 354 || (cn == 2 && value[0] == value[1]) 355 || (cn == 3 && value[0] == value[1] && value[0] == value[2]) 356 || (cn == 4 && value[0] == value[1] && value[0] == value[2] && value[0] == value[3])) 357 { 358 const int val = cv::saturate_cast<uchar>(value[0]); 359 360 if (stream) 361 CV_CUDEV_SAFE_CALL( cudaMemset2DAsync(data, step, val, cols * elemSize(), rows, StreamAccessor::getStream(stream)) ); 362 else 363 CV_CUDEV_SAFE_CALL( cudaMemset2D(data, step, val, cols * elemSize(), rows) ); 364 365 return *this; 366 } 367 } 368 369 typedef void (*func_t)(const GpuMat& mat, Scalar scalar, Stream& stream); 370 static const func_t funcs[7][4] = 371 { 372 {setToWithOutMask<uchar>,setToWithOutMask<uchar2>,setToWithOutMask<uchar3>,setToWithOutMask<uchar4>}, 373 {setToWithOutMask<schar>,setToWithOutMask<char2>,setToWithOutMask<char3>,setToWithOutMask<char4>}, 374 {setToWithOutMask<ushort>,setToWithOutMask<ushort2>,setToWithOutMask<ushort3>,setToWithOutMask<ushort4>}, 375 {setToWithOutMask<short>,setToWithOutMask<short2>,setToWithOutMask<short3>,setToWithOutMask<short4>}, 376 {setToWithOutMask<int>,setToWithOutMask<int2>,setToWithOutMask<int3>,setToWithOutMask<int4>}, 377 {setToWithOutMask<float>,setToWithOutMask<float2>,setToWithOutMask<float3>,setToWithOutMask<float4>}, 378 {setToWithOutMask<double>,setToWithOutMask<double2>,setToWithOutMask<double3>,setToWithOutMask<double4>} 379 }; 380 381 funcs[depth()][channels() - 1](*this, value, stream); 382 383 return *this; 384 } 385 386 GpuMat& cv::cuda::GpuMat::setTo(Scalar value, InputArray _mask, Stream& stream) 387 { 388 CV_DbgAssert( !empty() ); 389 CV_DbgAssert( depth() <= CV_64F && channels() <= 4 ); 390 391 GpuMat mask = _mask.getGpuMat(); 392 393 if (mask.empty()) 394 { 395 return setTo(value, stream); 396 } 397 398 CV_DbgAssert( size() == mask.size() && mask.type() == CV_8UC1 ); 399 400 typedef void (*func_t)(const GpuMat& mat, const GpuMat& mask, Scalar scalar, Stream& stream); 401 static const func_t funcs[7][4] = 402 { 403 {setToWithMask<uchar>,setToWithMask<uchar2>,setToWithMask<uchar3>,setToWithMask<uchar4>}, 404 {setToWithMask<schar>,setToWithMask<char2>,setToWithMask<char3>,setToWithMask<char4>}, 405 {setToWithMask<ushort>,setToWithMask<ushort2>,setToWithMask<ushort3>,setToWithMask<ushort4>}, 406 {setToWithMask<short>,setToWithMask<short2>,setToWithMask<short3>,setToWithMask<short4>}, 407 {setToWithMask<int>,setToWithMask<int2>,setToWithMask<int3>,setToWithMask<int4>}, 408 {setToWithMask<float>,setToWithMask<float2>,setToWithMask<float3>,setToWithMask<float4>}, 409 {setToWithMask<double>,setToWithMask<double2>,setToWithMask<double3>,setToWithMask<double4>} 410 }; 411 412 funcs[depth()][channels() - 1](*this, mask, value, stream); 413 414 return *this; 415 } 416 417 ///////////////////////////////////////////////////// 418 /// convertTo 419 420 namespace 421 { 422 template <typename T> struct ConvertToPolicy : DefaultTransformPolicy 423 { 424 }; 425 template <> struct ConvertToPolicy<double> : DefaultTransformPolicy 426 { 427 enum { 428 shift = 1 429 }; 430 }; 431 432 template <typename T, typename D> 433 void convertToNoScale(const GpuMat& src, const GpuMat& dst, Stream& stream) 434 { 435 typedef typename VecTraits<T>::elem_type src_elem_type; 436 typedef typename VecTraits<D>::elem_type dst_elem_type; 437 typedef typename LargerType<src_elem_type, float>::type larger_elem_type; 438 typedef typename LargerType<float, dst_elem_type>::type scalar_type; 439 440 gridTransformUnary_< ConvertToPolicy<scalar_type> >(globPtr<T>(src), globPtr<D>(dst), saturate_cast_func<T, D>(), stream); 441 } 442 443 template <typename T, typename D, typename S> struct Convertor : unary_function<T, D> 444 { 445 S alpha; 446 S beta; 447 448 __device__ __forceinline__ D operator ()(typename TypeTraits<T>::parameter_type src) const 449 { 450 return cudev::saturate_cast<D>(alpha * src + beta); 451 } 452 }; 453 454 template <typename T, typename D> 455 void convertToScale(const GpuMat& src, const GpuMat& dst, double alpha, double beta, Stream& stream) 456 { 457 typedef typename VecTraits<T>::elem_type src_elem_type; 458 typedef typename VecTraits<D>::elem_type dst_elem_type; 459 typedef typename LargerType<src_elem_type, float>::type larger_elem_type; 460 typedef typename LargerType<float, dst_elem_type>::type scalar_type; 461 462 Convertor<T, D, scalar_type> op; 463 op.alpha = cv::saturate_cast<scalar_type>(alpha); 464 op.beta = cv::saturate_cast<scalar_type>(beta); 465 466 gridTransformUnary_< ConvertToPolicy<scalar_type> >(globPtr<T>(src), globPtr<D>(dst), op, stream); 467 } 468 } 469 470 void cv::cuda::GpuMat::convertTo(OutputArray _dst, int rtype, Stream& stream) const 471 { 472 if (rtype < 0) 473 rtype = type(); 474 else 475 rtype = CV_MAKE_TYPE(CV_MAT_DEPTH(rtype), channels()); 476 477 const int sdepth = depth(); 478 const int ddepth = CV_MAT_DEPTH(rtype); 479 if (sdepth == ddepth) 480 { 481 if (stream) 482 copyTo(_dst, stream); 483 else 484 copyTo(_dst); 485 486 return; 487 } 488 489 CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F ); 490 491 GpuMat src = *this; 492 493 _dst.create(size(), rtype); 494 GpuMat dst = _dst.getGpuMat(); 495 496 typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream); 497 static const func_t funcs[7][7] = 498 { 499 {0, convertToNoScale<uchar, schar>, convertToNoScale<uchar, ushort>, convertToNoScale<uchar, short>, convertToNoScale<uchar, int>, convertToNoScale<uchar, float>, convertToNoScale<uchar, double>}, 500 {convertToNoScale<schar, uchar>, 0, convertToNoScale<schar, ushort>, convertToNoScale<schar, short>, convertToNoScale<schar, int>, convertToNoScale<schar, float>, convertToNoScale<schar, double>}, 501 {convertToNoScale<ushort, uchar>, convertToNoScale<ushort, schar>, 0, convertToNoScale<ushort, short>, convertToNoScale<ushort, int>, convertToNoScale<ushort, float>, convertToNoScale<ushort, double>}, 502 {convertToNoScale<short, uchar>, convertToNoScale<short, schar>, convertToNoScale<short, ushort>, 0, convertToNoScale<short, int>, convertToNoScale<short, float>, convertToNoScale<short, double>}, 503 {convertToNoScale<int, uchar>, convertToNoScale<int, schar>, convertToNoScale<int, ushort>, convertToNoScale<int, short>, 0, convertToNoScale<int, float>, convertToNoScale<int, double>}, 504 {convertToNoScale<float, uchar>, convertToNoScale<float, schar>, convertToNoScale<float, ushort>, convertToNoScale<float, short>, convertToNoScale<float, int>, 0, convertToNoScale<float, double>}, 505 {convertToNoScale<double, uchar>, convertToNoScale<double, schar>, convertToNoScale<double, ushort>, convertToNoScale<double, short>, convertToNoScale<double, int>, convertToNoScale<double, float>, 0} 506 }; 507 508 funcs[sdepth][ddepth](reshape(1), dst.reshape(1), stream); 509 } 510 511 void cv::cuda::GpuMat::convertTo(OutputArray _dst, int rtype, double alpha, double beta, Stream& stream) const 512 { 513 if (rtype < 0) 514 rtype = type(); 515 else 516 rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), channels()); 517 518 const int sdepth = depth(); 519 const int ddepth = CV_MAT_DEPTH(rtype); 520 521 GpuMat src = *this; 522 523 _dst.create(size(), rtype); 524 GpuMat dst = _dst.getGpuMat(); 525 526 typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, double alpha, double beta, Stream& stream); 527 static const func_t funcs[7][7] = 528 { 529 {convertToScale<uchar, uchar>, convertToScale<uchar, schar>, convertToScale<uchar, ushort>, convertToScale<uchar, short>, convertToScale<uchar, int>, convertToScale<uchar, float>, convertToScale<uchar, double>}, 530 {convertToScale<schar, uchar>, convertToScale<schar, schar>, convertToScale<schar, ushort>, convertToScale<schar, short>, convertToScale<schar, int>, convertToScale<schar, float>, convertToScale<schar, double>}, 531 {convertToScale<ushort, uchar>, convertToScale<ushort, schar>, convertToScale<ushort, ushort>, convertToScale<ushort, short>, convertToScale<ushort, int>, convertToScale<ushort, float>, convertToScale<ushort, double>}, 532 {convertToScale<short, uchar>, convertToScale<short, schar>, convertToScale<short, ushort>, convertToScale<short, short>, convertToScale<short, int>, convertToScale<short, float>, convertToScale<short, double>}, 533 {convertToScale<int, uchar>, convertToScale<int, schar>, convertToScale<int, ushort>, convertToScale<int, short>, convertToScale<int, int>, convertToScale<int, float>, convertToScale<int, double>}, 534 {convertToScale<float, uchar>, convertToScale<float, schar>, convertToScale<float, ushort>, convertToScale<float, short>, convertToScale<float, int>, convertToScale<float, float>, convertToScale<float, double>}, 535 {convertToScale<double, uchar>, convertToScale<double, schar>, convertToScale<double, ushort>, convertToScale<double, short>, convertToScale<double, int>, convertToScale<double, float>, convertToScale<double, double>} 536 }; 537 538 funcs[sdepth][ddepth](reshape(1), dst.reshape(1), alpha, beta, stream); 539 } 540 541 #endif 542