diff --git a/modules/core/include/opencv2/core/cuda/common.hpp b/modules/core/include/opencv2/core/cuda/common.hpp index 680ec497c..203efae1e 100644 --- a/modules/core/include/opencv2/core/cuda/common.hpp +++ b/modules/core/include/opencv2/core/cuda/common.hpp @@ -45,10 +45,8 @@ #include #include "opencv2/core/cuda_devptrs.hpp" - -#ifndef CV_PI - #define CV_PI 3.1415926535897932384626433832795 -#endif +#include "opencv2/core/cvdef.h" +#include "opencv2/core/base.hpp" #ifndef CV_PI_F #ifndef CV_PI @@ -58,16 +56,22 @@ #endif #endif +namespace cv { namespace gpu { namespace cuda { + static inline void checkError(cudaError_t err, const char* file, const int line, const char* func) + { + if (cudaSuccess != err) + cv::error(cv::Error::GpuApiCallError, cudaGetErrorString(err), func, file, line); + } +}}} + #if defined(__GNUC__) - #define cudaSafeCall(expr) ___cudaSafeCall(expr, __FILE__, __LINE__, __func__) + #define cvCudaSafeCall(expr) cv::gpu::cuda::checkError((expr), __FILE__, __LINE__, __func__) #else /* defined(__CUDACC__) || defined(__MSVC__) */ - #define cudaSafeCall(expr) ___cudaSafeCall(expr, __FILE__, __LINE__) + #define cvCudaSafeCall(expr) cv::gpu::cuda::checkError((expr), __FILE__, __LINE__, "") #endif namespace cv { namespace gpu { - void error(const char *error_string, const char *file, const int line, const char *func); - template static inline bool isAligned(const T* ptr, size_t size) { return reinterpret_cast(ptr) % size == 0; @@ -79,38 +83,32 @@ namespace cv { namespace gpu } }} -static inline void ___cudaSafeCall(cudaError_t err, const char *file, const int line, const char *func = "") -{ - if (cudaSuccess != err) - cv::gpu::error(cudaGetErrorString(err), file, line, func); -} - namespace cv { namespace gpu { - __host__ __device__ __forceinline__ int divUp(int total, int grain) + enum { - return (total + grain - 1) / grain; - } - - namespace cuda - { - using cv::gpu::divUp; + BORDER_REFLECT101_GPU = 0, + BORDER_REPLICATE_GPU, + BORDER_CONSTANT_GPU, + BORDER_REFLECT_GPU, + BORDER_WRAP_GPU + }; #ifdef __CUDACC__ - typedef unsigned char uchar; - typedef unsigned short ushort; - typedef signed char schar; - #if defined (_WIN32) || defined (__APPLE__) - typedef unsigned int uint; - #endif + namespace cuda + { + __host__ __device__ __forceinline__ int divUp(int total, int grain) + { + return (total + grain - 1) / grain; + } template inline void bindTexture(const textureReference* tex, const PtrStepSz& img) { cudaChannelFormatDesc desc = cudaCreateChannelDesc(); - cudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) ); + cvCudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) ); } -#endif // __CUDACC__ } +#endif // __CUDACC__ }} diff --git a/modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp b/modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp index 0ed8d64db..731aa15ce 100644 --- a/modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp +++ b/modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp @@ -317,10 +317,10 @@ namespace cv { namespace gpu { namespace cuda const dim3 grid(divUp(src.cols, threads.x), divUp(src.rows, threads.y), 1); transformSimple<<>>(src, dst, mask, op); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -332,10 +332,10 @@ namespace cv { namespace gpu { namespace cuda const dim3 grid(divUp(src1.cols, threads.x), divUp(src1.rows, threads.y), 1); transformSimple<<>>(src1, src2, dst, mask, op); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; template<> struct TransformDispatcher @@ -345,7 +345,7 @@ namespace cv { namespace gpu { namespace cuda { typedef TransformFunctorTraits ft; - StaticAssert::check(); + CV_StaticAssert(ft::smart_shift != 1, ""); if (!isAligned(src.data, ft::smart_shift * sizeof(T)) || !isAligned(src.step, ft::smart_shift * sizeof(T)) || !isAligned(dst.data, ft::smart_shift * sizeof(D)) || !isAligned(dst.step, ft::smart_shift * sizeof(D))) @@ -358,10 +358,10 @@ namespace cv { namespace gpu { namespace cuda const dim3 grid(divUp(src.cols, threads.x * ft::smart_shift), divUp(src.rows, threads.y), 1); transformSmart<<>>(src, dst, mask, op); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -369,7 +369,7 @@ namespace cv { namespace gpu { namespace cuda { typedef TransformFunctorTraits ft; - StaticAssert::check(); + CV_StaticAssert(ft::smart_shift != 1, ""); if (!isAligned(src1.data, ft::smart_shift * sizeof(T1)) || !isAligned(src1.step, ft::smart_shift * sizeof(T1)) || !isAligned(src2.data, ft::smart_shift * sizeof(T2)) || !isAligned(src2.step, ft::smart_shift * sizeof(T2)) || @@ -383,10 +383,10 @@ namespace cv { namespace gpu { namespace cuda const dim3 grid(divUp(src1.cols, threads.x * ft::smart_shift), divUp(src1.rows, threads.y), 1); transformSmart<<>>(src1, src2, dst, mask, op); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; } // namespace transform_detail diff --git a/modules/core/include/opencv2/core/cuda_devptrs.hpp b/modules/core/include/opencv2/core/cuda_devptrs.hpp index 9e0ba11da..c82ce61b3 100644 --- a/modules/core/include/opencv2/core/cuda_devptrs.hpp +++ b/modules/core/include/opencv2/core/cuda_devptrs.hpp @@ -58,9 +58,6 @@ namespace cv // Simple lightweight structures that encapsulates information about an image on device. // It is intended to pass to nvcc-compiled code. GpuMat depends on headers that nvcc can't compile - template struct StaticAssert; - template <> struct StaticAssert {static __CV_GPU_HOST_DEVICE__ void check(){}}; - template struct DevPtr { typedef T elem_type; diff --git a/modules/core/include/opencv2/core/gpu_private.hpp b/modules/core/include/opencv2/core/gpu_private.hpp new file mode 100644 index 000000000..31ea5e6e9 --- /dev/null +++ b/modules/core/include/opencv2/core/gpu_private.hpp @@ -0,0 +1,134 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Copyright (C) 2013, OpenCV Foundation, all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __OPENCV_CORE_GPU_PRIVATE_HPP__ +#define __OPENCV_CORE_GPU_PRIVATE_HPP__ + +#ifndef __OPENCV_BUILD +# error this is a private header which should not be used from outside of the OpenCV library +#endif + +#include "cvconfig.h" + +#include "opencv2/core/cvdef.h" +#include "opencv2/core/base.hpp" + +#ifdef HAVE_CUDA +# include +# include +# include +# include "opencv2/core/stream_accessor.hpp" +# include "opencv2/core/cuda/common.hpp" + +# define CUDART_MINIMUM_REQUIRED_VERSION 4020 + +# if (CUDART_VERSION < CUDART_MINIMUM_REQUIRED_VERSION) +# error "Insufficient Cuda Runtime library version, please update it." +# endif + +# if defined(CUDA_ARCH_BIN_OR_PTX_10) +# error "OpenCV GPU module doesn't support NVIDIA compute capability 1.0" +# endif +#endif + +namespace cv { namespace gpu { + CV_EXPORTS cv::String getNppErrorMessage(int code); + + static inline void checkNppError(int code, const char* file, const int line, const char* func) + { + if (code < 0) + cv::error(cv::Error::GpuApiCallError, getNppErrorMessage(code), func, file, line); + } + + // Converts CPU border extrapolation mode into GPU internal analogue. + // Returns true if the GPU analogue exists, false otherwise. + CV_EXPORTS bool tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType); +}} + +#ifndef HAVE_CUDA + +static inline void throw_no_cuda() { CV_Error(cv::Error::GpuNotSupported, "The library is compiled without GPU support"); } + +#else // HAVE_CUDA + +static inline void throw_no_cuda() { CV_Error(cv::Error::StsNotImplemented, "The called functionality is disabled for current build or platform"); } + +#if defined(__GNUC__) + #define nppSafeCall(expr) cv::gpu::checkNppError(expr, __FILE__, __LINE__, __func__) +#else /* defined(__CUDACC__) || defined(__MSVC__) */ + #define nppSafeCall(expr) cv::gpu::checkNppError(expr, __FILE__, __LINE__, "") +#endif + +namespace cv { namespace gpu +{ + template struct NPPTypeTraits; + template<> struct NPPTypeTraits { typedef Npp8u npp_type; }; + template<> struct NPPTypeTraits { typedef Npp8s npp_type; }; + template<> struct NPPTypeTraits { typedef Npp16u npp_type; }; + template<> struct NPPTypeTraits { typedef Npp16s npp_type; }; + template<> struct NPPTypeTraits { typedef Npp32s npp_type; }; + template<> struct NPPTypeTraits { typedef Npp32f npp_type; }; + template<> struct NPPTypeTraits { typedef Npp64f npp_type; }; + + class NppStreamHandler + { + public: + inline explicit NppStreamHandler(cudaStream_t newStream) + { + oldStream = nppGetStream(); + nppSetStream(newStream); + } + + inline ~NppStreamHandler() + { + nppSetStream(oldStream); + } + + private: + cudaStream_t oldStream; + }; +}} + +#endif // HAVE_CUDA + +#endif // __OPENCV_CORE_GPU_PRIVATE_HPP__ diff --git a/modules/core/include/opencv2/core/gpumat.hpp b/modules/core/include/opencv2/core/gpumat.hpp index 60d37e500..52b87b000 100644 --- a/modules/core/include/opencv2/core/gpumat.hpp +++ b/modules/core/include/opencv2/core/gpumat.hpp @@ -454,11 +454,6 @@ CV_EXPORTS void ensureSizeIsEnough(int rows, int cols, int type, GpuMat& m); CV_EXPORTS GpuMat allocMatFromBuf(int rows, int cols, int type, GpuMat &mat); -//////////////////////////////////////////////////////////////////////// -// Error handling - -CV_EXPORTS void error(const char* error_string, const char* file, const int line, const char* func = ""); - //////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////// diff --git a/modules/core/include/opencv2/core/stream_accessor.hpp b/modules/core/include/opencv2/core/stream_accessor.hpp index 9b87470e7..3f98eb0a3 100644 --- a/modules/core/include/opencv2/core/stream_accessor.hpp +++ b/modules/core/include/opencv2/core/stream_accessor.hpp @@ -43,17 +43,20 @@ #ifndef __OPENCV_CUDA_STREAM_ACCESSOR_HPP__ #define __OPENCV_CUDA_STREAM_ACCESSOR_HPP__ -#include "opencv2/core/gpumat.hpp" -#include "cuda_runtime_api.h" +#include +#include "opencv2/core/cvdef.h" + +// This is only header file that depends on Cuda. All other headers are independent. +// So if you use OpenCV binaries you do noot need to install Cuda Toolkit. +// But of you wanna use GPU by yourself, may get cuda stream instance using the class below. +// In this case you have to install Cuda Toolkit. namespace cv { namespace gpu { - // This is only header file that depends on Cuda. All other headers are independent. - // So if you use OpenCV binaries you do noot need to install Cuda Toolkit. - // But of you wanna use GPU by yourself, may get cuda stream instance using the class below. - // In this case you have to install Cuda Toolkit. + class Stream; + struct StreamAccessor { CV_EXPORTS static cudaStream_t getStream(const Stream& stream); diff --git a/modules/core/src/cuda/matrix_operations.cu b/modules/core/src/cuda/matrix_operations.cu index 2cb82184c..725321286 100644 --- a/modules/core/src/cuda/matrix_operations.cu +++ b/modules/core/src/cuda/matrix_operations.cu @@ -124,31 +124,31 @@ namespace cv { namespace gpu { namespace cuda void writeScalar(const uchar* vals) { - cudaSafeCall( cudaMemcpyToSymbol(scalar_8u, vals, sizeof(uchar) * 4) ); + cvCudaSafeCall( cudaMemcpyToSymbol(scalar_8u, vals, sizeof(uchar) * 4) ); } void writeScalar(const schar* vals) { - cudaSafeCall( cudaMemcpyToSymbol(scalar_8s, vals, sizeof(schar) * 4) ); + cvCudaSafeCall( cudaMemcpyToSymbol(scalar_8s, vals, sizeof(schar) * 4) ); } void writeScalar(const ushort* vals) { - cudaSafeCall( cudaMemcpyToSymbol(scalar_16u, vals, sizeof(ushort) * 4) ); + cvCudaSafeCall( cudaMemcpyToSymbol(scalar_16u, vals, sizeof(ushort) * 4) ); } void writeScalar(const short* vals) { - cudaSafeCall( cudaMemcpyToSymbol(scalar_16s, vals, sizeof(short) * 4) ); + cvCudaSafeCall( cudaMemcpyToSymbol(scalar_16s, vals, sizeof(short) * 4) ); } void writeScalar(const int* vals) { - cudaSafeCall( cudaMemcpyToSymbol(scalar_32s, vals, sizeof(int) * 4) ); + cvCudaSafeCall( cudaMemcpyToSymbol(scalar_32s, vals, sizeof(int) * 4) ); } void writeScalar(const float* vals) { - cudaSafeCall( cudaMemcpyToSymbol(scalar_32f, vals, sizeof(float) * 4) ); + cvCudaSafeCall( cudaMemcpyToSymbol(scalar_32f, vals, sizeof(float) * 4) ); } void writeScalar(const double* vals) { - cudaSafeCall( cudaMemcpyToSymbol(scalar_64f, vals, sizeof(double) * 4) ); + cvCudaSafeCall( cudaMemcpyToSymbol(scalar_64f, vals, sizeof(double) * 4) ); } template @@ -186,10 +186,10 @@ namespace cv { namespace gpu { namespace cuda dim3 numBlocks (mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1); set_to_with_mask<<>>((T*)mat.data, (uchar*)mask.data, mat.cols, mat.rows, mat.step, channels, mask.step); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall ( cudaDeviceSynchronize() ); + cvCudaSafeCall ( cudaDeviceSynchronize() ); } template void set_to_gpu(PtrStepSzb mat, const uchar* scalar, PtrStepSzb mask, int channels, cudaStream_t stream); @@ -209,10 +209,10 @@ namespace cv { namespace gpu { namespace cuda dim3 numBlocks (mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1); set_to_without_mask<<>>((T*)mat.data, mat.cols, mat.rows, mat.step, channels); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall ( cudaDeviceSynchronize() ); + cvCudaSafeCall ( cudaDeviceSynchronize() ); } template void set_to_gpu(PtrStepSzb mat, const uchar* scalar, int channels, cudaStream_t stream); @@ -290,8 +290,8 @@ namespace cv { namespace gpu { namespace cuda template void cvt_(PtrStepSzb src, PtrStepSzb dst, double alpha, double beta, cudaStream_t stream) { - cudaSafeCall( cudaSetDoubleForDevice(&alpha) ); - cudaSafeCall( cudaSetDoubleForDevice(&beta) ); + cvCudaSafeCall( cudaSetDoubleForDevice(&alpha) ); + cvCudaSafeCall( cudaSetDoubleForDevice(&beta) ); Convertor op(static_cast(alpha), static_cast(beta)); cv::gpu::cuda::transform((PtrStepSz)src, (PtrStepSz)dst, op, WithOutMask(), stream); } diff --git a/modules/core/src/cudastream.cpp b/modules/core/src/cudastream.cpp index 270865a9d..a6d1a41e6 100644 --- a/modules/core/src/cudastream.cpp +++ b/modules/core/src/cudastream.cpp @@ -46,33 +46,30 @@ using namespace cv; using namespace cv::gpu; #if !defined (HAVE_CUDA) -#define throw_nogpu() CV_Error(CV_GpuNotSupported, "The library is compiled without CUDA support") -cv::gpu::Stream::Stream() { throw_nogpu(); } +cv::gpu::Stream::Stream() { throw_no_cuda(); } cv::gpu::Stream::~Stream() {} -cv::gpu::Stream::Stream(const Stream&) { throw_nogpu(); } -Stream& cv::gpu::Stream::operator=(const Stream&) { throw_nogpu(); return *this; } -bool cv::gpu::Stream::queryIfComplete() { throw_nogpu(); return false; } -void cv::gpu::Stream::waitForCompletion() { throw_nogpu(); } -void cv::gpu::Stream::enqueueDownload(const GpuMat&, Mat&) { throw_nogpu(); } -void cv::gpu::Stream::enqueueDownload(const GpuMat&, CudaMem&) { throw_nogpu(); } -void cv::gpu::Stream::enqueueUpload(const CudaMem&, GpuMat&) { throw_nogpu(); } -void cv::gpu::Stream::enqueueUpload(const Mat&, GpuMat&) { throw_nogpu(); } -void cv::gpu::Stream::enqueueCopy(const GpuMat&, GpuMat&) { throw_nogpu(); } -void cv::gpu::Stream::enqueueMemSet(GpuMat&, Scalar) { throw_nogpu(); } -void cv::gpu::Stream::enqueueMemSet(GpuMat&, Scalar, const GpuMat&) { throw_nogpu(); } -void cv::gpu::Stream::enqueueConvert(const GpuMat&, GpuMat&, int, double, double) { throw_nogpu(); } -void cv::gpu::Stream::enqueueHostCallback(StreamCallback, void*) { throw_nogpu(); } -Stream& cv::gpu::Stream::Null() { throw_nogpu(); static Stream s; return s; } -cv::gpu::Stream::operator bool() const { throw_nogpu(); return false; } -cv::gpu::Stream::Stream(Impl*) { throw_nogpu(); } -void cv::gpu::Stream::create() { throw_nogpu(); } -void cv::gpu::Stream::release() { throw_nogpu(); } +cv::gpu::Stream::Stream(const Stream&) { throw_no_cuda(); } +Stream& cv::gpu::Stream::operator=(const Stream&) { throw_no_cuda(); return *this; } +bool cv::gpu::Stream::queryIfComplete() { throw_no_cuda(); return false; } +void cv::gpu::Stream::waitForCompletion() { throw_no_cuda(); } +void cv::gpu::Stream::enqueueDownload(const GpuMat&, Mat&) { throw_no_cuda(); } +void cv::gpu::Stream::enqueueDownload(const GpuMat&, CudaMem&) { throw_no_cuda(); } +void cv::gpu::Stream::enqueueUpload(const CudaMem&, GpuMat&) { throw_no_cuda(); } +void cv::gpu::Stream::enqueueUpload(const Mat&, GpuMat&) { throw_no_cuda(); } +void cv::gpu::Stream::enqueueCopy(const GpuMat&, GpuMat&) { throw_no_cuda(); } +void cv::gpu::Stream::enqueueMemSet(GpuMat&, Scalar) { throw_no_cuda(); } +void cv::gpu::Stream::enqueueMemSet(GpuMat&, Scalar, const GpuMat&) { throw_no_cuda(); } +void cv::gpu::Stream::enqueueConvert(const GpuMat&, GpuMat&, int, double, double) { throw_no_cuda(); } +void cv::gpu::Stream::enqueueHostCallback(StreamCallback, void*) { throw_no_cuda(); } +Stream& cv::gpu::Stream::Null() { throw_no_cuda(); static Stream s; return s; } +cv::gpu::Stream::operator bool() const { throw_no_cuda(); return false; } +cv::gpu::Stream::Stream(Impl*) { throw_no_cuda(); } +void cv::gpu::Stream::create() { throw_no_cuda(); } +void cv::gpu::Stream::release() { throw_no_cuda(); } #else /* !defined (HAVE_CUDA) */ -#include "opencv2/core/stream_accessor.hpp" - namespace cv { namespace gpu { void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream); @@ -134,14 +131,14 @@ bool cv::gpu::Stream::queryIfComplete() if (err == cudaErrorNotReady || err == cudaSuccess) return err == cudaSuccess; - cudaSafeCall(err); + cvCudaSafeCall(err); return false; } void cv::gpu::Stream::waitForCompletion() { cudaStream_t stream = Impl::getStream(impl); - cudaSafeCall( cudaStreamSynchronize(stream) ); + cvCudaSafeCall( cudaStreamSynchronize(stream) ); } void cv::gpu::Stream::enqueueDownload(const GpuMat& src, Mat& dst) @@ -151,7 +148,7 @@ void cv::gpu::Stream::enqueueDownload(const GpuMat& src, Mat& dst) cudaStream_t stream = Impl::getStream(impl); size_t bwidth = src.cols * src.elemSize(); - cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) ); + cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) ); } void cv::gpu::Stream::enqueueDownload(const GpuMat& src, CudaMem& dst) @@ -160,7 +157,7 @@ void cv::gpu::Stream::enqueueDownload(const GpuMat& src, CudaMem& dst) cudaStream_t stream = Impl::getStream(impl); size_t bwidth = src.cols * src.elemSize(); - cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) ); + cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) ); } void cv::gpu::Stream::enqueueUpload(const CudaMem& src, GpuMat& dst) @@ -169,7 +166,7 @@ void cv::gpu::Stream::enqueueUpload(const CudaMem& src, GpuMat& dst) cudaStream_t stream = Impl::getStream(impl); size_t bwidth = src.cols * src.elemSize(); - cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) ); + cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) ); } void cv::gpu::Stream::enqueueUpload(const Mat& src, GpuMat& dst) @@ -178,7 +175,7 @@ void cv::gpu::Stream::enqueueUpload(const Mat& src, GpuMat& dst) cudaStream_t stream = Impl::getStream(impl); size_t bwidth = src.cols * src.elemSize(); - cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) ); + cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) ); } void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst) @@ -187,7 +184,7 @@ void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst) cudaStream_t stream = Impl::getStream(impl); size_t bwidth = src.cols * src.elemSize(); - cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToDevice, stream) ); + cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToDevice, stream) ); } void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val) @@ -204,7 +201,7 @@ void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val) if (val[0] == 0.0 && val[1] == 0.0 && val[2] == 0.0 && val[3] == 0.0) { - cudaSafeCall( cudaMemset2DAsync(src.data, src.step, 0, src.cols * src.elemSize(), src.rows, stream) ); + cvCudaSafeCall( cudaMemset2DAsync(src.data, src.step, 0, src.cols * src.elemSize(), src.rows, stream) ); return; } @@ -215,7 +212,7 @@ void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val) if (cn == 1 || (cn == 2 && val[0] == val[1]) || (cn == 3 && val[0] == val[1] && val[0] == val[2]) || (cn == 4 && val[0] == val[1] && val[0] == val[2] && val[0] == val[3])) { int ival = saturate_cast(val[0]); - cudaSafeCall( cudaMemset2DAsync(src.data, src.step, ival, src.cols * src.elemSize(), src.rows, stream) ); + cvCudaSafeCall( cudaMemset2DAsync(src.data, src.step, ival, src.cols * src.elemSize(), src.rows, stream) ); return; } } @@ -302,7 +299,7 @@ void cv::gpu::Stream::enqueueHostCallback(StreamCallback callback, void* userDat cudaStream_t stream = Impl::getStream(impl); - cudaSafeCall( cudaStreamAddCallback(stream, cudaStreamCallback, data, 0) ); + cvCudaSafeCall( cudaStreamAddCallback(stream, cudaStreamCallback, data, 0) ); #else (void) callback; (void) userData; @@ -331,7 +328,7 @@ void cv::gpu::Stream::create() release(); cudaStream_t stream; - cudaSafeCall( cudaStreamCreate( &stream ) ); + cvCudaSafeCall( cudaStreamCreate( &stream ) ); impl = (Stream::Impl*) fastMalloc(sizeof(Stream::Impl)); @@ -343,7 +340,7 @@ void cv::gpu::Stream::release() { if (impl && CV_XADD(&impl->ref_counter, -1) == 1) { - cudaSafeCall( cudaStreamDestroy(impl->stream) ); + cvCudaSafeCall( cudaStreamDestroy(impl->stream) ); cv::fastFree(impl); } } diff --git a/modules/core/src/gpumat.cpp b/modules/core/src/gpumat.cpp index 12252b538..3e1f3fe0b 100644 --- a/modules/core/src/gpumat.cpp +++ b/modules/core/src/gpumat.cpp @@ -45,64 +45,38 @@ using namespace cv; using namespace cv::gpu; -#ifndef HAVE_CUDA - -#define throw_nogpu CV_Error(CV_GpuNotSupported, "The library is compiled without CUDA support") - -#else // HAVE_CUDA - -namespace -{ -#if defined(__GNUC__) - #define nppSafeCall(expr) ___nppSafeCall(expr, __FILE__, __LINE__, __func__) -#else /* defined(__CUDACC__) || defined(__MSVC__) */ - #define nppSafeCall(expr) ___nppSafeCall(expr, __FILE__, __LINE__) -#endif - - inline void ___nppSafeCall(int err, const char *file, const int line, const char *func = "") - { - if (err < 0) - { - String msg = cv::format("NPP API Call Error: %d", err); - cv::gpu::error(msg.c_str(), file, line, func); - } - } -} - -#endif // HAVE_CUDA - //////////////////////////////// Initialization & Info //////////////////////// #ifndef HAVE_CUDA int cv::gpu::getCudaEnabledDeviceCount() { return 0; } -void cv::gpu::setDevice(int) { throw_nogpu; } -int cv::gpu::getDevice() { throw_nogpu; return 0; } +void cv::gpu::setDevice(int) { throw_no_cuda(); } +int cv::gpu::getDevice() { throw_no_cuda(); return 0; } -void cv::gpu::resetDevice() { throw_nogpu; } +void cv::gpu::resetDevice() { throw_no_cuda(); } -bool cv::gpu::deviceSupports(FeatureSet) { throw_nogpu; return false; } +bool cv::gpu::deviceSupports(FeatureSet) { throw_no_cuda(); return false; } -bool cv::gpu::TargetArchs::builtWith(FeatureSet) { throw_nogpu; return false; } -bool cv::gpu::TargetArchs::has(int, int) { throw_nogpu; return false; } -bool cv::gpu::TargetArchs::hasPtx(int, int) { throw_nogpu; return false; } -bool cv::gpu::TargetArchs::hasBin(int, int) { throw_nogpu; return false; } -bool cv::gpu::TargetArchs::hasEqualOrLessPtx(int, int) { throw_nogpu; return false; } -bool cv::gpu::TargetArchs::hasEqualOrGreater(int, int) { throw_nogpu; return false; } -bool cv::gpu::TargetArchs::hasEqualOrGreaterPtx(int, int) { throw_nogpu; return false; } -bool cv::gpu::TargetArchs::hasEqualOrGreaterBin(int, int) { throw_nogpu; return false; } +bool cv::gpu::TargetArchs::builtWith(FeatureSet) { throw_no_cuda(); return false; } +bool cv::gpu::TargetArchs::has(int, int) { throw_no_cuda(); return false; } +bool cv::gpu::TargetArchs::hasPtx(int, int) { throw_no_cuda(); return false; } +bool cv::gpu::TargetArchs::hasBin(int, int) { throw_no_cuda(); return false; } +bool cv::gpu::TargetArchs::hasEqualOrLessPtx(int, int) { throw_no_cuda(); return false; } +bool cv::gpu::TargetArchs::hasEqualOrGreater(int, int) { throw_no_cuda(); return false; } +bool cv::gpu::TargetArchs::hasEqualOrGreaterPtx(int, int) { throw_no_cuda(); return false; } +bool cv::gpu::TargetArchs::hasEqualOrGreaterBin(int, int) { throw_no_cuda(); return false; } -size_t cv::gpu::DeviceInfo::sharedMemPerBlock() const { throw_nogpu; return 0; } -void cv::gpu::DeviceInfo::queryMemory(size_t&, size_t&) const { throw_nogpu; } -size_t cv::gpu::DeviceInfo::freeMemory() const { throw_nogpu; return 0; } -size_t cv::gpu::DeviceInfo::totalMemory() const { throw_nogpu; return 0; } -bool cv::gpu::DeviceInfo::supports(FeatureSet) const { throw_nogpu; return false; } -bool cv::gpu::DeviceInfo::isCompatible() const { throw_nogpu; return false; } -void cv::gpu::DeviceInfo::query() { throw_nogpu; } +size_t cv::gpu::DeviceInfo::sharedMemPerBlock() const { throw_no_cuda(); return 0; } +void cv::gpu::DeviceInfo::queryMemory(size_t&, size_t&) const { throw_no_cuda(); } +size_t cv::gpu::DeviceInfo::freeMemory() const { throw_no_cuda(); return 0; } +size_t cv::gpu::DeviceInfo::totalMemory() const { throw_no_cuda(); return 0; } +bool cv::gpu::DeviceInfo::supports(FeatureSet) const { throw_no_cuda(); return false; } +bool cv::gpu::DeviceInfo::isCompatible() const { throw_no_cuda(); return false; } +void cv::gpu::DeviceInfo::query() { throw_no_cuda(); } -void cv::gpu::printCudaDeviceInfo(int) { throw_nogpu; } -void cv::gpu::printShortCudaDeviceInfo(int) { throw_nogpu; } +void cv::gpu::printCudaDeviceInfo(int) { throw_no_cuda(); } +void cv::gpu::printShortCudaDeviceInfo(int) { throw_no_cuda(); } #else // HAVE_CUDA @@ -117,25 +91,25 @@ int cv::gpu::getCudaEnabledDeviceCount() if (error == cudaErrorNoDevice) return 0; - cudaSafeCall( error ); + cvCudaSafeCall( error ); return count; } void cv::gpu::setDevice(int device) { - cudaSafeCall( cudaSetDevice( device ) ); + cvCudaSafeCall( cudaSetDevice( device ) ); } int cv::gpu::getDevice() { int device; - cudaSafeCall( cudaGetDevice( &device ) ); + cvCudaSafeCall( cudaGetDevice( &device ) ); return device; } void cv::gpu::resetDevice() { - cudaSafeCall( cudaDeviceReset() ); + cvCudaSafeCall( cudaDeviceReset() ); } namespace @@ -328,7 +302,7 @@ namespace if (!props_[devID]) { props_[devID] = new cudaDeviceProp; - cudaSafeCall( cudaGetDeviceProperties(props_[devID], devID) ); + cvCudaSafeCall( cudaGetDeviceProperties(props_[devID], devID) ); } return props_[devID]; @@ -348,7 +322,7 @@ void cv::gpu::DeviceInfo::queryMemory(size_t& _totalMemory, size_t& _freeMemory) if (prevDeviceID != device_id_) setDevice(device_id_); - cudaSafeCall( cudaMemGetInfo(&_freeMemory, &_totalMemory) ); + cvCudaSafeCall( cudaMemGetInfo(&_freeMemory, &_totalMemory) ); if (prevDeviceID != device_id_) setDevice(prevDeviceID); @@ -434,8 +408,8 @@ void cv::gpu::printCudaDeviceInfo(int device) printf("Device count: %d\n", count); int driverVersion = 0, runtimeVersion = 0; - cudaSafeCall( cudaDriverGetVersion(&driverVersion) ); - cudaSafeCall( cudaRuntimeGetVersion(&runtimeVersion) ); + cvCudaSafeCall( cudaDriverGetVersion(&driverVersion) ); + cvCudaSafeCall( cudaRuntimeGetVersion(&runtimeVersion) ); const char *computeMode[] = { "Default (multiple host threads can use ::cudaSetDevice() with device simultaneously)", @@ -449,7 +423,7 @@ void cv::gpu::printCudaDeviceInfo(int device) for(int dev = beg; dev < end; ++dev) { cudaDeviceProp prop; - cudaSafeCall( cudaGetDeviceProperties(&prop, dev) ); + cvCudaSafeCall( cudaGetDeviceProperties(&prop, dev) ); printf("\nDevice %d: \"%s\"\n", dev, prop.name); printf(" CUDA Driver Version / Runtime Version %d.%d / %d.%d\n", driverVersion/1000, driverVersion%100, runtimeVersion/1000, runtimeVersion%100); @@ -511,13 +485,13 @@ void cv::gpu::printShortCudaDeviceInfo(int device) int end = valid ? device+1 : count; int driverVersion = 0, runtimeVersion = 0; - cudaSafeCall( cudaDriverGetVersion(&driverVersion) ); - cudaSafeCall( cudaRuntimeGetVersion(&runtimeVersion) ); + cvCudaSafeCall( cudaDriverGetVersion(&driverVersion) ); + cvCudaSafeCall( cudaRuntimeGetVersion(&runtimeVersion) ); for(int dev = beg; dev < end; ++dev) { cudaDeviceProp prop; - cudaSafeCall( cudaGetDeviceProperties(&prop, dev) ); + cvCudaSafeCall( cudaGetDeviceProperties(&prop, dev) ); const char *arch_str = prop.major < 2 ? " (not Fermi)" : ""; printf("Device %d: \"%s\" %.0fMb", dev, prop.name, (float)prop.totalGlobalMem/1048576.0f); @@ -846,18 +820,18 @@ namespace class EmptyFuncTable : public GpuFuncTable { public: - void copy(const Mat&, GpuMat&) const { throw_nogpu; } - void copy(const GpuMat&, Mat&) const { throw_nogpu; } - void copy(const GpuMat&, GpuMat&) const { throw_nogpu; } + void copy(const Mat&, GpuMat&) const { throw_no_cuda(); } + void copy(const GpuMat&, Mat&) const { throw_no_cuda(); } + void copy(const GpuMat&, GpuMat&) const { throw_no_cuda(); } - void copyWithMask(const GpuMat&, GpuMat&, const GpuMat&) const { throw_nogpu; } + void copyWithMask(const GpuMat&, GpuMat&, const GpuMat&) const { throw_no_cuda(); } - void convert(const GpuMat&, GpuMat&) const { throw_nogpu; } - void convert(const GpuMat&, GpuMat&, double, double) const { throw_nogpu; } + void convert(const GpuMat&, GpuMat&) const { throw_no_cuda(); } + void convert(const GpuMat&, GpuMat&, double, double) const { throw_no_cuda(); } - void setTo(GpuMat&, Scalar, const GpuMat&) const { throw_nogpu; } + void setTo(GpuMat&, Scalar, const GpuMat&) const { throw_no_cuda(); } - void mallocPitch(void**, size_t*, size_t, size_t) const { throw_nogpu; } + void mallocPitch(void**, size_t*, size_t, size_t) const { throw_no_cuda(); } void free(void*) const {} }; @@ -1009,7 +983,7 @@ namespace nppSafeCall( func(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz) ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; template::func_ptr func> struct NppCvt @@ -1024,7 +998,7 @@ namespace nppSafeCall( func(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz, NPP_RND_NEAR) ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -1066,7 +1040,7 @@ namespace nppSafeCall( func(nppS.val, src.ptr(), static_cast(src.step), sz) ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; template::func_ptr func> struct NppSet @@ -1083,7 +1057,7 @@ namespace nppSafeCall( func(nppS[0], src.ptr(), static_cast(src.step), sz) ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -1114,7 +1088,7 @@ namespace nppSafeCall( func(nppS.val, src.ptr(), static_cast(src.step), sz, mask.ptr(), static_cast(mask.step)) ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; template::func_ptr func> struct NppSetMask @@ -1131,7 +1105,7 @@ namespace nppSafeCall( func(nppS[0], src.ptr(), static_cast(src.step), sz, mask.ptr(), static_cast(mask.step)) ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -1157,7 +1131,7 @@ namespace nppSafeCall( func(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz, mask.ptr(), static_cast(mask.step)) ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -1174,15 +1148,15 @@ namespace public: void copy(const Mat& src, GpuMat& dst) const { - cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyHostToDevice) ); + cvCudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyHostToDevice) ); } void copy(const GpuMat& src, Mat& dst) const { - cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToHost) ); + cvCudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToHost) ); } void copy(const GpuMat& src, GpuMat& dst) const { - cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToDevice) ); + cvCudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToDevice) ); } void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask) const @@ -1327,7 +1301,7 @@ namespace { if (s[0] == 0.0 && s[1] == 0.0 && s[2] == 0.0 && s[3] == 0.0) { - cudaSafeCall( cudaMemset2D(m.data, m.step, 0, m.cols * m.elemSize(), m.rows) ); + cvCudaSafeCall( cudaMemset2D(m.data, m.step, 0, m.cols * m.elemSize(), m.rows) ); return; } @@ -1338,7 +1312,7 @@ namespace if (cn == 1 || (cn == 2 && s[0] == s[1]) || (cn == 3 && s[0] == s[1] && s[0] == s[2]) || (cn == 4 && s[0] == s[1] && s[0] == s[2] && s[0] == s[3])) { int val = saturate_cast(s[0]); - cudaSafeCall( cudaMemset2D(m.data, m.step, val, m.cols * m.elemSize(), m.rows) ); + cvCudaSafeCall( cudaMemset2D(m.data, m.step, val, m.cols * m.elemSize(), m.rows) ); return; } } @@ -1393,7 +1367,7 @@ namespace void mallocPitch(void** devPtr, size_t* step, size_t width, size_t height) const { - cudaSafeCall( cudaMallocPitch(devPtr, step, width, height) ); + cvCudaSafeCall( cudaMallocPitch(devPtr, step, width, height) ); } void free(void* devPtr) const @@ -1551,18 +1525,117 @@ void cv::gpu::GpuMat::release() //////////////////////////////////////////////////////////////////////// // Error handling -void cv::gpu::error(const char *error_string, const char *file, const int line, const char *func) +#ifdef HAVE_CUDA + +namespace { - int code = CV_GpuApiCallError; + #define error_entry(entry) { entry, #entry } - if (std::uncaught_exception()) + struct ErrorEntry { - const char* errorStr = cvErrorStr(code); - const char* function = func ? func : "unknown function"; + int code; + const char* str; + }; - fprintf(stderr, "OpenCV Error: %s(%s) in %s, file %s, line %d", errorStr, error_string, function, file, line); - fflush(stderr); - } - else - cv::error( cv::Exception(code, error_string, func, file, line) ); + struct ErrorEntryComparer + { + int code; + ErrorEntryComparer(int code_) : code(code_) {} + bool operator()(const ErrorEntry& e) const { return e.code == code; } + }; + + const ErrorEntry npp_errors [] = + { + error_entry( NPP_NOT_SUPPORTED_MODE_ERROR ), + error_entry( NPP_ROUND_MODE_NOT_SUPPORTED_ERROR ), + error_entry( NPP_RESIZE_NO_OPERATION_ERROR ), + +#if defined (_MSC_VER) + error_entry( NPP_NOT_SUFFICIENT_COMPUTE_CAPABILITY ), +#endif + + error_entry( NPP_BAD_ARG_ERROR ), + error_entry( NPP_LUT_NUMBER_OF_LEVELS_ERROR ), + error_entry( NPP_TEXTURE_BIND_ERROR ), + error_entry( NPP_COEFF_ERROR ), + error_entry( NPP_RECT_ERROR ), + error_entry( NPP_QUAD_ERROR ), + error_entry( NPP_WRONG_INTERSECTION_ROI_ERROR ), + error_entry( NPP_NOT_EVEN_STEP_ERROR ), + error_entry( NPP_INTERPOLATION_ERROR ), + error_entry( NPP_RESIZE_FACTOR_ERROR ), + error_entry( NPP_HAAR_CLASSIFIER_PIXEL_MATCH_ERROR ), + error_entry( NPP_MEMFREE_ERR ), + error_entry( NPP_MEMSET_ERR ), + error_entry( NPP_MEMCPY_ERROR ), + error_entry( NPP_MEM_ALLOC_ERR ), + error_entry( NPP_HISTO_NUMBER_OF_LEVELS_ERROR ), + error_entry( NPP_MIRROR_FLIP_ERR ), + error_entry( NPP_INVALID_INPUT ), + error_entry( NPP_ALIGNMENT_ERROR ), + error_entry( NPP_STEP_ERROR ), + error_entry( NPP_SIZE_ERROR ), + error_entry( NPP_POINTER_ERROR ), + error_entry( NPP_NULL_POINTER_ERROR ), + error_entry( NPP_CUDA_KERNEL_EXECUTION_ERROR ), + error_entry( NPP_NOT_IMPLEMENTED_ERROR ), + error_entry( NPP_ERROR ), + error_entry( NPP_NO_ERROR ), + error_entry( NPP_SUCCESS ), + error_entry( NPP_WARNING ), + error_entry( NPP_WRONG_INTERSECTION_QUAD_WARNING ), + error_entry( NPP_MISALIGNED_DST_ROI_WARNING ), + error_entry( NPP_AFFINE_QUAD_INCORRECT_WARNING ), + error_entry( NPP_DOUBLE_SIZE_WARNING ), + error_entry( NPP_ODD_ROI_WARNING ) + }; + + const size_t npp_error_num = sizeof(npp_errors) / sizeof(npp_errors[0]); +} + +#endif + +String cv::gpu::getNppErrorMessage(int code) +{ +#ifndef HAVE_CUDA + (void) code; + return String(); +#else + size_t idx = std::find_if(npp_errors, npp_errors + npp_error_num, ErrorEntryComparer(code)) - npp_errors; + + const char* msg = (idx != npp_error_num) ? npp_errors[idx].str : "Unknown error code"; + String str = cv::format("%s [Code = %d]", msg, code); + + return str; +#endif +} + +bool cv::gpu::tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType) +{ +#ifndef HAVE_CUDA + (void) cpuBorderType; + (void) gpuBorderType; + return false; +#else + switch (cpuBorderType) + { + case IPL_BORDER_REFLECT_101: + gpuBorderType = cv::gpu::BORDER_REFLECT101_GPU; + return true; + case IPL_BORDER_REPLICATE: + gpuBorderType = cv::gpu::BORDER_REPLICATE_GPU; + return true; + case IPL_BORDER_CONSTANT: + gpuBorderType = cv::gpu::BORDER_CONSTANT_GPU; + return true; + case IPL_BORDER_REFLECT: + gpuBorderType = cv::gpu::BORDER_REFLECT_GPU; + return true; + case IPL_BORDER_WRAP: + gpuBorderType = cv::gpu::BORDER_WRAP_GPU; + return true; + default: + return false; + }; +#endif } diff --git a/modules/core/src/matrix_operations.cpp b/modules/core/src/matrix_operations.cpp index 53764a168..ef09ef62a 100644 --- a/modules/core/src/matrix_operations.cpp +++ b/modules/core/src/matrix_operations.cpp @@ -41,7 +41,6 @@ //M*/ #include "precomp.hpp" -#include "opencv2/core/gpumat.hpp" using namespace cv; using namespace cv::gpu; @@ -181,30 +180,29 @@ bool cv::gpu::CudaMem::empty() const #if !defined (HAVE_CUDA) -void cv::gpu::registerPageLocked(Mat&) { CV_Error(CV_GpuNotSupported, "The library is compiled without CUDA support"); } -void cv::gpu::unregisterPageLocked(Mat&) { CV_Error(CV_GpuNotSupported, "The library is compiled without CUDA support"); } -void cv::gpu::CudaMem::create(int /*_rows*/, int /*_cols*/, int /*_type*/, int /*type_alloc*/) -{ CV_Error(CV_GpuNotSupported, "The library is compiled without CUDA support"); } -bool cv::gpu::CudaMem::canMapHostMemory() { CV_Error(CV_GpuNotSupported, "The library is compiled without CUDA support"); return false; } -void cv::gpu::CudaMem::release() { CV_Error(CV_GpuNotSupported, "The library is compiled without CUDA support"); } -GpuMat cv::gpu::CudaMem::createGpuMatHeader () const { CV_Error(CV_GpuNotSupported, "The library is compiled without CUDA support"); return GpuMat(); } +void cv::gpu::registerPageLocked(Mat&) { throw_no_cuda(); } +void cv::gpu::unregisterPageLocked(Mat&) { throw_no_cuda(); } +void cv::gpu::CudaMem::create(int, int, int, int) { throw_no_cuda(); } +bool cv::gpu::CudaMem::canMapHostMemory() { throw_no_cuda(); return false; } +void cv::gpu::CudaMem::release() { throw_no_cuda(); } +GpuMat cv::gpu::CudaMem::createGpuMatHeader () const { throw_no_cuda(); return GpuMat(); } #else /* !defined (HAVE_CUDA) */ void cv::gpu::registerPageLocked(Mat& m) { - cudaSafeCall( cudaHostRegister(m.ptr(), m.step * m.rows, cudaHostRegisterPortable) ); + cvCudaSafeCall( cudaHostRegister(m.ptr(), m.step * m.rows, cudaHostRegisterPortable) ); } void cv::gpu::unregisterPageLocked(Mat& m) { - cudaSafeCall( cudaHostUnregister(m.ptr()) ); + cvCudaSafeCall( cudaHostUnregister(m.ptr()) ); } bool cv::gpu::CudaMem::canMapHostMemory() { cudaDeviceProp prop; - cudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) ); + cvCudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) ); return (prop.canMapHostMemory != 0) ? true : false; } @@ -222,7 +220,7 @@ namespace void cv::gpu::CudaMem::create(int _rows, int _cols, int _type, int _alloc_type) { if (_alloc_type == ALLOC_ZEROCOPY && !canMapHostMemory()) - cv::gpu::error("ZeroCopy is not supported by current device", __FILE__, __LINE__); + CV_Error(cv::Error::GpuApiCallError, "ZeroCopy is not supported by current device"); _type &= Mat::TYPE_MASK; if( rows == _rows && cols == _cols && type() == _type && data ) @@ -239,7 +237,7 @@ void cv::gpu::CudaMem::create(int _rows, int _cols, int _type, int _alloc_type) if (_alloc_type == ALLOC_ZEROCOPY) { cudaDeviceProp prop; - cudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) ); + cvCudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) ); step = alignUpStep(step, prop.textureAlignment); } int64 _nettosize = (int64)step*rows; @@ -254,10 +252,10 @@ void cv::gpu::CudaMem::create(int _rows, int _cols, int _type, int _alloc_type) switch (alloc_type) { - case ALLOC_PAGE_LOCKED: cudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocDefault) ); break; - case ALLOC_ZEROCOPY: cudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocMapped) ); break; - case ALLOC_WRITE_COMBINED: cudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocWriteCombined) ); break; - default: cv::gpu::error("Invalid alloc type", __FILE__, __LINE__); + case ALLOC_PAGE_LOCKED: cvCudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocDefault) ); break; + case ALLOC_ZEROCOPY: cvCudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocMapped) ); break; + case ALLOC_WRITE_COMBINED: cvCudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocWriteCombined) ); break; + default: CV_Error(cv::Error::StsBadFlag, "Invalid alloc type"); } datastart = data = (uchar*)ptr; @@ -270,15 +268,13 @@ void cv::gpu::CudaMem::create(int _rows, int _cols, int _type, int _alloc_type) GpuMat cv::gpu::CudaMem::createGpuMatHeader () const { + CV_Assert( alloc_type == ALLOC_ZEROCOPY ); + GpuMat res; - if (alloc_type == ALLOC_ZEROCOPY) - { - void *pdev; - cudaSafeCall( cudaHostGetDevicePointer( &pdev, data, 0 ) ); - res = GpuMat(rows, cols, type(), pdev, step); - } - else - cv::gpu::error("Zero-copy is not supported or memory was allocated without zero-copy flag", __FILE__, __LINE__); + + void *pdev; + cvCudaSafeCall( cudaHostGetDevicePointer( &pdev, data, 0 ) ); + res = GpuMat(rows, cols, type(), pdev, step); return res; } @@ -287,7 +283,7 @@ void cv::gpu::CudaMem::release() { if( refcount && CV_XADD(refcount, -1) == 1 ) { - cudaSafeCall( cudaFreeHost(datastart ) ); + cvCudaSafeCall( cudaFreeHost(datastart ) ); fastFree(refcount); } data = datastart = dataend = 0; diff --git a/modules/core/src/opengl_interop.cpp b/modules/core/src/opengl_interop.cpp index cfed9fcec..19eabfa7d 100644 --- a/modules/core/src/opengl_interop.cpp +++ b/modules/core/src/opengl_interop.cpp @@ -41,16 +41,12 @@ //M*/ #include "precomp.hpp" -#include "opencv2/core/opengl.hpp" -#include "opencv2/core/gpumat.hpp" #ifdef HAVE_OPENGL - #include "gl_core_3_1.hpp" - - #ifdef HAVE_CUDA - #include - #include - #endif +# include "gl_core_3_1.hpp" +# ifdef HAVE_CUDA +# include +# endif #endif using namespace cv; @@ -59,15 +55,9 @@ using namespace cv::gpu; namespace { #ifndef HAVE_OPENGL - void throw_nogl() { CV_Error(CV_OpenGlNotSupported, "The library is compiled without OpenGL support"); } + void throw_no_ogl() { CV_Error(CV_OpenGlNotSupported, "The library is compiled without OpenGL support"); } #else - void throw_nogl() { CV_Error(CV_OpenGlApiCallError, "OpenGL context doesn't exist"); } - - #ifndef HAVE_CUDA - void throw_nocuda() { CV_Error(CV_GpuNotSupported, "The library is compiled without GPU support"); } - #else - void throw_nocuda() { CV_Error(CV_StsNotImplemented, "The called functionality is disabled for current build or platform"); } - #endif + void throw_no_ogl() { CV_Error(CV_OpenGlApiCallError, "OpenGL context doesn't exist"); } #endif bool checkError(const char* file, const int line, const char* func = 0) @@ -137,13 +127,13 @@ void cv::gpu::setGlDevice(int device) { #ifndef HAVE_OPENGL (void) device; - throw_nogl(); + throw_no_ogl(); #else #if !defined(HAVE_CUDA) || defined(CUDA_DISABLER) (void) device; - throw_nocuda(); + throw_no_cuda(); #else - cudaSafeCall( cudaGLSetGLDevice(device) ); + cvCudaSafeCall( cudaGLSetGLDevice(device) ); #endif #endif } @@ -194,7 +184,7 @@ namespace return; cudaGraphicsResource_t resource; - cudaSafeCall( cudaGraphicsGLRegisterBuffer(&resource, buffer, cudaGraphicsMapFlagsNone) ); + cvCudaSafeCall( cudaGraphicsGLRegisterBuffer(&resource, buffer, cudaGraphicsMapFlagsNone) ); release(); @@ -227,7 +217,7 @@ namespace CudaResource::GraphicsMapHolder::GraphicsMapHolder(cudaGraphicsResource_t* resource, cudaStream_t stream) : resource_(resource), stream_(stream) { if (resource_) - cudaSafeCall( cudaGraphicsMapResources(1, resource_, stream_) ); + cvCudaSafeCall( cudaGraphicsMapResources(1, resource_, stream_) ); } CudaResource::GraphicsMapHolder::~GraphicsMapHolder() @@ -250,14 +240,14 @@ namespace void* dst; size_t size; - cudaSafeCall( cudaGraphicsResourceGetMappedPointer(&dst, &size, resource_) ); + cvCudaSafeCall( cudaGraphicsResourceGetMappedPointer(&dst, &size, resource_) ); CV_DbgAssert( width * height == size ); if (stream == 0) - cudaSafeCall( cudaMemcpy2D(dst, width, src, spitch, width, height, cudaMemcpyDeviceToDevice) ); + cvCudaSafeCall( cudaMemcpy2D(dst, width, src, spitch, width, height, cudaMemcpyDeviceToDevice) ); else - cudaSafeCall( cudaMemcpy2DAsync(dst, width, src, spitch, width, height, cudaMemcpyDeviceToDevice, stream) ); + cvCudaSafeCall( cudaMemcpy2DAsync(dst, width, src, spitch, width, height, cudaMemcpyDeviceToDevice, stream) ); } void CudaResource::copyTo(void* dst, size_t dpitch, size_t width, size_t height, cudaStream_t stream) @@ -269,14 +259,14 @@ namespace void* src; size_t size; - cudaSafeCall( cudaGraphicsResourceGetMappedPointer(&src, &size, resource_) ); + cvCudaSafeCall( cudaGraphicsResourceGetMappedPointer(&src, &size, resource_) ); CV_DbgAssert( width * height == size ); if (stream == 0) - cudaSafeCall( cudaMemcpy2D(dst, dpitch, src, width, width, height, cudaMemcpyDeviceToDevice) ); + cvCudaSafeCall( cudaMemcpy2D(dst, dpitch, src, width, width, height, cudaMemcpyDeviceToDevice) ); else - cudaSafeCall( cudaMemcpy2DAsync(dst, dpitch, src, width, width, height, cudaMemcpyDeviceToDevice, stream) ); + cvCudaSafeCall( cudaMemcpy2DAsync(dst, dpitch, src, width, width, height, cudaMemcpyDeviceToDevice, stream) ); } void* CudaResource::map(cudaStream_t stream) @@ -287,7 +277,7 @@ namespace void* ptr; size_t size; - cudaSafeCall( cudaGraphicsResourceGetMappedPointer(&ptr, &size, resource_) ); + cvCudaSafeCall( cudaGraphicsResourceGetMappedPointer(&ptr, &size, resource_) ); h.reset(); @@ -476,7 +466,7 @@ void cv::ogl::Buffer::Impl::unmapHost() cv::ogl::Buffer::Buffer() : rows_(0), cols_(0), type_(0) { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); #else impl_ = Impl::empty(); #endif @@ -490,7 +480,7 @@ cv::ogl::Buffer::Buffer(int arows, int acols, int atype, unsigned int abufId, bo (void) atype; (void) abufId; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else impl_ = new Impl(abufId, autoRelease); rows_ = arows; @@ -506,7 +496,7 @@ cv::ogl::Buffer::Buffer(Size asize, int atype, unsigned int abufId, bool autoRel (void) atype; (void) abufId; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else impl_ = new Impl(abufId, autoRelease); rows_ = asize.height; @@ -531,7 +521,7 @@ cv::ogl::Buffer::Buffer(InputArray arr, Target target, bool autoRelease) : rows_ (void) arr; (void) target; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else const int kind = arr.kind(); @@ -578,7 +568,7 @@ void cv::ogl::Buffer::create(int arows, int acols, int atype, Target target, boo (void) atype; (void) target; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else if (rows_ != arows || cols_ != acols || type_ != atype) { @@ -607,7 +597,7 @@ void cv::ogl::Buffer::setAutoRelease(bool flag) { #ifndef HAVE_OPENGL (void) flag; - throw_nogl(); + throw_no_ogl(); #else impl_->setAutoRelease(flag); #endif @@ -619,7 +609,7 @@ void cv::ogl::Buffer::copyFrom(InputArray arr, Target target, bool autoRelease) (void) arr; (void) target; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else const int kind = arr.kind(); @@ -647,7 +637,7 @@ void cv::ogl::Buffer::copyFrom(InputArray arr, Target target, bool autoRelease) case _InputArray::GPU_MAT: { #if !defined HAVE_CUDA || defined(CUDA_DISABLER) - throw_nocuda(); + throw_no_cuda(); #else GpuMat dmat = arr.getGpuMat(); impl_->copyFrom(dmat.data, dmat.step, dmat.cols * dmat.elemSize(), dmat.rows); @@ -672,7 +662,7 @@ void cv::ogl::Buffer::copyTo(OutputArray arr, Target target, bool autoRelease) c (void) arr; (void) target; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else const int kind = arr.kind(); @@ -693,7 +683,7 @@ void cv::ogl::Buffer::copyTo(OutputArray arr, Target target, bool autoRelease) c case _InputArray::GPU_MAT: { #if !defined HAVE_CUDA || defined(CUDA_DISABLER) - throw_nocuda(); + throw_no_cuda(); #else GpuMat& dmat = arr.getGpuMatRef(); dmat.create(rows_, cols_, type_); @@ -719,7 +709,7 @@ cv::ogl::Buffer cv::ogl::Buffer::clone(Target target, bool autoRelease) const #ifndef HAVE_OPENGL (void) target; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); return cv::ogl::Buffer(); #else ogl::Buffer buf; @@ -732,7 +722,7 @@ void cv::ogl::Buffer::bind(Target target) const { #ifndef HAVE_OPENGL (void) target; - throw_nogl(); + throw_no_ogl(); #else impl_->bind(target); #endif @@ -742,7 +732,7 @@ void cv::ogl::Buffer::unbind(Target target) { #ifndef HAVE_OPENGL (void) target; - throw_nogl(); + throw_no_ogl(); #else gl::BindBuffer(target, 0); CV_CheckGlError(); @@ -753,7 +743,7 @@ Mat cv::ogl::Buffer::mapHost(Access access) { #ifndef HAVE_OPENGL (void) access; - throw_nogl(); + throw_no_ogl(); return Mat(); #else return Mat(rows_, cols_, type_, impl_->mapHost(access)); @@ -763,7 +753,7 @@ Mat cv::ogl::Buffer::mapHost(Access access) void cv::ogl::Buffer::unmapHost() { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); #else return impl_->unmapHost(); #endif @@ -772,11 +762,11 @@ void cv::ogl::Buffer::unmapHost() GpuMat cv::ogl::Buffer::mapDevice() { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); return GpuMat(); #else #if !defined HAVE_CUDA || defined(CUDA_DISABLER) - throw_nocuda(); + throw_no_cuda(); return GpuMat(); #else return GpuMat(rows_, cols_, type_, impl_->mapDevice()); @@ -787,10 +777,10 @@ GpuMat cv::ogl::Buffer::mapDevice() void cv::ogl::Buffer::unmapDevice() { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); #else #if !defined HAVE_CUDA || defined(CUDA_DISABLER) - throw_nocuda(); + throw_no_cuda(); #else impl_->unmapDevice(); #endif @@ -800,7 +790,7 @@ void cv::ogl::Buffer::unmapDevice() unsigned int cv::ogl::Buffer::bufId() const { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); return 0; #else return impl_->bufId(); @@ -926,7 +916,7 @@ void cv::ogl::Texture2D::Impl::bind() const cv::ogl::Texture2D::Texture2D() : rows_(0), cols_(0), format_(NONE) { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); #else impl_ = Impl::empty(); #endif @@ -940,7 +930,7 @@ cv::ogl::Texture2D::Texture2D(int arows, int acols, Format aformat, unsigned int (void) aformat; (void) atexId; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else impl_ = new Impl(atexId, autoRelease); rows_ = arows; @@ -956,7 +946,7 @@ cv::ogl::Texture2D::Texture2D(Size asize, Format aformat, unsigned int atexId, b (void) aformat; (void) atexId; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else impl_ = new Impl(atexId, autoRelease); rows_ = asize.height; @@ -980,7 +970,7 @@ cv::ogl::Texture2D::Texture2D(InputArray arr, bool autoRelease) : rows_(0), cols #ifndef HAVE_OPENGL (void) arr; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else const int kind = arr.kind(); @@ -1016,7 +1006,7 @@ cv::ogl::Texture2D::Texture2D(InputArray arr, bool autoRelease) : rows_(0), cols case _InputArray::GPU_MAT: { #if !defined HAVE_CUDA || defined(CUDA_DISABLER) - throw_nocuda(); + throw_no_cuda(); #else GpuMat dmat = arr.getGpuMat(); ogl::Buffer buf(dmat, ogl::Buffer::PIXEL_UNPACK_BUFFER); @@ -1051,7 +1041,7 @@ void cv::ogl::Texture2D::create(int arows, int acols, Format aformat, bool autoR (void) acols; (void) aformat; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else if (rows_ != arows || cols_ != acols || format_ != aformat) { @@ -1080,7 +1070,7 @@ void cv::ogl::Texture2D::setAutoRelease(bool flag) { #ifndef HAVE_OPENGL (void) flag; - throw_nogl(); + throw_no_ogl(); #else impl_->setAutoRelease(flag); #endif @@ -1091,7 +1081,7 @@ void cv::ogl::Texture2D::copyFrom(InputArray arr, bool autoRelease) #ifndef HAVE_OPENGL (void) arr; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else const int kind = arr.kind(); @@ -1129,7 +1119,7 @@ void cv::ogl::Texture2D::copyFrom(InputArray arr, bool autoRelease) case _InputArray::GPU_MAT: { #if !defined HAVE_CUDA || defined(CUDA_DISABLER) - throw_nocuda(); + throw_no_cuda(); #else GpuMat dmat = arr.getGpuMat(); ogl::Buffer buf(dmat, ogl::Buffer::PIXEL_UNPACK_BUFFER); @@ -1158,7 +1148,7 @@ void cv::ogl::Texture2D::copyTo(OutputArray arr, int ddepth, bool autoRelease) c (void) arr; (void) ddepth; (void) autoRelease; - throw_nogl(); + throw_no_ogl(); #else const int kind = arr.kind(); @@ -1180,7 +1170,7 @@ void cv::ogl::Texture2D::copyTo(OutputArray arr, int ddepth, bool autoRelease) c case _InputArray::GPU_MAT: { #if !defined HAVE_CUDA || defined(CUDA_DISABLER) - throw_nocuda(); + throw_no_cuda(); #else ogl::Buffer buf(rows_, cols_, CV_MAKE_TYPE(ddepth, cn), ogl::Buffer::PIXEL_PACK_BUFFER); buf.bind(ogl::Buffer::PIXEL_PACK_BUFFER); @@ -1207,7 +1197,7 @@ void cv::ogl::Texture2D::copyTo(OutputArray arr, int ddepth, bool autoRelease) c void cv::ogl::Texture2D::bind() const { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); #else impl_->bind(); #endif @@ -1216,7 +1206,7 @@ void cv::ogl::Texture2D::bind() const unsigned int cv::ogl::Texture2D::texId() const { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); return 0; #else return impl_->texId(); @@ -1331,7 +1321,7 @@ void cv::ogl::Arrays::setAutoRelease(bool flag) void cv::ogl::Arrays::bind() const { #ifndef HAVE_OPENGL - throw_nogl(); + throw_no_ogl(); #else CV_Assert( texCoord_.empty() || texCoord_.size().area() == size_ ); CV_Assert( normal_.empty() || normal_.size().area() == size_ ); @@ -1416,7 +1406,7 @@ void cv::ogl::render(const ogl::Texture2D& tex, Rect_ wndRect, Rect_ #include @@ -64,37 +66,6 @@ #define GET_OPTIMIZED(func) (func) #endif -#ifdef HAVE_CUDA - -# include -# include - -# define CUDART_MINIMUM_REQUIRED_VERSION 4020 -# define NPP_MINIMUM_REQUIRED_VERSION 4200 - -# if (CUDART_VERSION < CUDART_MINIMUM_REQUIRED_VERSION) -# error "Insufficient Cuda Runtime library version, please update it." -# endif - -# if (NPP_VERSION_MAJOR * 1000 + NPP_VERSION_MINOR * 100 + NPP_VERSION_BUILD < NPP_MINIMUM_REQUIRED_VERSION) -# error "Insufficient NPP version, please update it." -# endif - -# if defined(__GNUC__) -# define cudaSafeCall(expr) ___cudaSafeCall(expr, __FILE__, __LINE__, __func__) -# else -# define cudaSafeCall(expr) ___cudaSafeCall(expr, __FILE__, __LINE__) -# endif - -static inline void ___cudaSafeCall(cudaError_t err, const char *file, const int line, const char *func = "") -{ - if (cudaSuccess != err) cv::gpu::error(cudaGetErrorString(err), file, line, func); -} - -#else -# define cudaSafeCall(expr) -#endif //HAVE_CUDA - namespace cv { diff --git a/modules/gpu/perf/perf_precomp.hpp b/modules/gpu/perf/perf_precomp.hpp index 56227223a..40930f7c0 100644 --- a/modules/gpu/perf/perf_precomp.hpp +++ b/modules/gpu/perf/perf_precomp.hpp @@ -54,10 +54,6 @@ #include #include -#ifdef HAVE_CUDA -#include -#endif - #include "opencv2/ts.hpp" #include "opencv2/ts/gpu_perf.hpp" @@ -70,7 +66,7 @@ #include "opencv2/legacy.hpp" #include "opencv2/photo.hpp" -#include "opencv2/core/private.hpp" +#include "opencv2/core/gpu_private.hpp" #ifdef GTEST_CREATE_SHARED_LIBRARY #error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index cefc29f20..8f9113991 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -47,19 +47,19 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) -void cv::gpu::gemm(const GpuMat&, const GpuMat&, double, const GpuMat&, double, GpuMat&, int, Stream&) { throw_nogpu(); } -void cv::gpu::transpose(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::flip(const GpuMat&, GpuMat&, int, Stream&) { throw_nogpu(); } -void cv::gpu::LUT(const GpuMat&, const Mat&, GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::magnitude(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::magnitudeSqr(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::magnitudeSqr(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::phase(const GpuMat&, const GpuMat&, GpuMat&, bool, Stream&) { throw_nogpu(); } -void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, Stream&) { throw_nogpu(); } -void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, Stream&) { throw_nogpu(); } -void cv::gpu::normalize(const GpuMat&, GpuMat&, double, double, int, int, const GpuMat&) { throw_nogpu(); } -void cv::gpu::normalize(const GpuMat&, GpuMat&, double, double, int, int, const GpuMat&, GpuMat&, GpuMat&) { throw_nogpu(); } +void cv::gpu::gemm(const GpuMat&, const GpuMat&, double, const GpuMat&, double, GpuMat&, int, Stream&) { throw_no_cuda(); } +void cv::gpu::transpose(const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::flip(const GpuMat&, GpuMat&, int, Stream&) { throw_no_cuda(); } +void cv::gpu::LUT(const GpuMat&, const Mat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::magnitude(const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::magnitudeSqr(const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::magnitudeSqr(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::phase(const GpuMat&, const GpuMat&, GpuMat&, bool, Stream&) { throw_no_cuda(); } +void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, Stream&) { throw_no_cuda(); } +void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, Stream&) { throw_no_cuda(); } +void cv::gpu::normalize(const GpuMat&, GpuMat&, double, double, int, int, const GpuMat&) { throw_no_cuda(); } +void cv::gpu::normalize(const GpuMat&, GpuMat&, double, double, int, int, const GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); } #else /* !defined (HAVE_CUDA) */ @@ -246,7 +246,7 @@ void cv::gpu::transpose(const GpuMat& src, GpuMat& dst, Stream& s) } if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////// @@ -287,7 +287,7 @@ namespace (flipCode == 0 ? NPP_HORIZONTAL_AXIS : (flipCode > 0 ? NPP_VERTICAL_AXIS : NPP_BOTH_AXIS))) ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; } @@ -402,7 +402,7 @@ void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst, Stream& s) } if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////// @@ -427,7 +427,7 @@ namespace nppSafeCall( func(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz) ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } } diff --git a/modules/gpu/src/bgfg_gmg.cpp b/modules/gpu/src/bgfg_gmg.cpp index b474823a5..79777fcb7 100644 --- a/modules/gpu/src/bgfg_gmg.cpp +++ b/modules/gpu/src/bgfg_gmg.cpp @@ -44,9 +44,9 @@ #if !defined HAVE_CUDA || defined(CUDA_DISABLER) -cv::gpu::GMG_GPU::GMG_GPU() { throw_nogpu(); } -void cv::gpu::GMG_GPU::initialize(cv::Size, float, float) { throw_nogpu(); } -void cv::gpu::GMG_GPU::operator ()(const cv::gpu::GpuMat&, cv::gpu::GpuMat&, float, cv::gpu::Stream&) { throw_nogpu(); } +cv::gpu::GMG_GPU::GMG_GPU() { throw_no_cuda(); } +void cv::gpu::GMG_GPU::initialize(cv::Size, float, float) { throw_no_cuda(); } +void cv::gpu::GMG_GPU::operator ()(const cv::gpu::GpuMat&, cv::gpu::GpuMat&, float, cv::gpu::Stream&) { throw_no_cuda(); } void cv::gpu::GMG_GPU::release() {} #else diff --git a/modules/gpu/src/bgfg_mog.cpp b/modules/gpu/src/bgfg_mog.cpp index 13db07911..78cd9a4a0 100644 --- a/modules/gpu/src/bgfg_mog.cpp +++ b/modules/gpu/src/bgfg_mog.cpp @@ -44,16 +44,16 @@ #if !defined HAVE_CUDA || defined(CUDA_DISABLER) -cv::gpu::MOG_GPU::MOG_GPU(int) { throw_nogpu(); } -void cv::gpu::MOG_GPU::initialize(cv::Size, int) { throw_nogpu(); } -void cv::gpu::MOG_GPU::operator()(const cv::gpu::GpuMat&, cv::gpu::GpuMat&, float, Stream&) { throw_nogpu(); } -void cv::gpu::MOG_GPU::getBackgroundImage(GpuMat&, Stream&) const { throw_nogpu(); } +cv::gpu::MOG_GPU::MOG_GPU(int) { throw_no_cuda(); } +void cv::gpu::MOG_GPU::initialize(cv::Size, int) { throw_no_cuda(); } +void cv::gpu::MOG_GPU::operator()(const cv::gpu::GpuMat&, cv::gpu::GpuMat&, float, Stream&) { throw_no_cuda(); } +void cv::gpu::MOG_GPU::getBackgroundImage(GpuMat&, Stream&) const { throw_no_cuda(); } void cv::gpu::MOG_GPU::release() {} -cv::gpu::MOG2_GPU::MOG2_GPU(int) { throw_nogpu(); } -void cv::gpu::MOG2_GPU::initialize(cv::Size, int) { throw_nogpu(); } -void cv::gpu::MOG2_GPU::operator()(const GpuMat&, GpuMat&, float, Stream&) { throw_nogpu(); } -void cv::gpu::MOG2_GPU::getBackgroundImage(GpuMat&, Stream&) const { throw_nogpu(); } +cv::gpu::MOG2_GPU::MOG2_GPU(int) { throw_no_cuda(); } +void cv::gpu::MOG2_GPU::initialize(cv::Size, int) { throw_no_cuda(); } +void cv::gpu::MOG2_GPU::operator()(const GpuMat&, GpuMat&, float, Stream&) { throw_no_cuda(); } +void cv::gpu::MOG2_GPU::getBackgroundImage(GpuMat&, Stream&) const { throw_no_cuda(); } void cv::gpu::MOG2_GPU::release() {} #else diff --git a/modules/gpu/src/bilateral_filter.cpp b/modules/gpu/src/bilateral_filter.cpp index 41640a556..41b6275f2 100644 --- a/modules/gpu/src/bilateral_filter.cpp +++ b/modules/gpu/src/bilateral_filter.cpp @@ -47,10 +47,10 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) -cv::gpu::DisparityBilateralFilter::DisparityBilateralFilter(int, int, int) { throw_nogpu(); } -cv::gpu::DisparityBilateralFilter::DisparityBilateralFilter(int, int, int, float, float, float) { throw_nogpu(); } +cv::gpu::DisparityBilateralFilter::DisparityBilateralFilter(int, int, int) { throw_no_cuda(); } +cv::gpu::DisparityBilateralFilter::DisparityBilateralFilter(int, int, int, float, float, float) { throw_no_cuda(); } -void cv::gpu::DisparityBilateralFilter::operator()(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::DisparityBilateralFilter::operator()(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } #else /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/blend.cpp b/modules/gpu/src/blend.cpp index 08193386b..97469f6b8 100644 --- a/modules/gpu/src/blend.cpp +++ b/modules/gpu/src/blend.cpp @@ -47,7 +47,7 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) -void cv::gpu::blendLinear(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::blendLinear(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } #else diff --git a/modules/gpu/src/brute_force_matcher.cpp b/modules/gpu/src/brute_force_matcher.cpp index 1f9c11cd0..8fd4e11fd 100644 --- a/modules/gpu/src/brute_force_matcher.cpp +++ b/modules/gpu/src/brute_force_matcher.cpp @@ -47,37 +47,37 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) -cv::gpu::BFMatcher_GPU::BFMatcher_GPU(int) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::add(const std::vector&) { throw_nogpu(); } -const std::vector& cv::gpu::BFMatcher_GPU::getTrainDescriptors() const { throw_nogpu(); return trainDescCollection; } -void cv::gpu::BFMatcher_GPU::clear() { throw_nogpu(); } -bool cv::gpu::BFMatcher_GPU::empty() const { throw_nogpu(); return true; } -bool cv::gpu::BFMatcher_GPU::isMaskSupported() const { throw_nogpu(); return true; } -void cv::gpu::BFMatcher_GPU::matchSingle(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, const GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::matchDownload(const GpuMat&, const GpuMat&, std::vector&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::matchConvert(const Mat&, const Mat&, std::vector&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::match(const GpuMat&, const GpuMat&, std::vector&, const GpuMat&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::makeGpuCollection(GpuMat&, GpuMat&, const std::vector&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::matchCollection(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, const GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::matchDownload(const GpuMat&, const GpuMat&, const GpuMat&, std::vector&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::matchConvert(const Mat&, const Mat&, const Mat&, std::vector&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::match(const GpuMat&, std::vector&, const std::vector&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::knnMatchSingle(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, int, const GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::knnMatchDownload(const GpuMat&, const GpuMat&, std::vector< std::vector >&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::knnMatchConvert(const Mat&, const Mat&, std::vector< std::vector >&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::knnMatch(const GpuMat&, const GpuMat&, std::vector< std::vector >&, int, const GpuMat&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::knnMatch2Collection(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, const GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::knnMatch2Download(const GpuMat&, const GpuMat&, const GpuMat&, std::vector< std::vector >&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::knnMatch2Convert(const Mat&, const Mat&, const Mat&, std::vector< std::vector >&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::knnMatch(const GpuMat&, std::vector< std::vector >&, int, const std::vector&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::radiusMatchSingle(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, float, const GpuMat&, Stream&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::radiusMatchDownload(const GpuMat&, const GpuMat&, const GpuMat&, std::vector< std::vector >&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::radiusMatchConvert(const Mat&, const Mat&, const Mat&, std::vector< std::vector >&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::radiusMatch(const GpuMat&, const GpuMat&, std::vector< std::vector >&, float, const GpuMat&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::radiusMatchCollection(const GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, float, const std::vector&, Stream&) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::radiusMatchDownload(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, std::vector< std::vector >&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::radiusMatchConvert(const Mat&, const Mat&, const Mat&, const Mat&, std::vector< std::vector >&, bool) { throw_nogpu(); } -void cv::gpu::BFMatcher_GPU::radiusMatch(const GpuMat&, std::vector< std::vector >&, float, const std::vector&, bool) { throw_nogpu(); } +cv::gpu::BFMatcher_GPU::BFMatcher_GPU(int) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::add(const std::vector&) { throw_no_cuda(); } +const std::vector& cv::gpu::BFMatcher_GPU::getTrainDescriptors() const { throw_no_cuda(); return trainDescCollection; } +void cv::gpu::BFMatcher_GPU::clear() { throw_no_cuda(); } +bool cv::gpu::BFMatcher_GPU::empty() const { throw_no_cuda(); return true; } +bool cv::gpu::BFMatcher_GPU::isMaskSupported() const { throw_no_cuda(); return true; } +void cv::gpu::BFMatcher_GPU::matchSingle(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, const GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::matchDownload(const GpuMat&, const GpuMat&, std::vector&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::matchConvert(const Mat&, const Mat&, std::vector&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::match(const GpuMat&, const GpuMat&, std::vector&, const GpuMat&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::makeGpuCollection(GpuMat&, GpuMat&, const std::vector&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::matchCollection(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, const GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::matchDownload(const GpuMat&, const GpuMat&, const GpuMat&, std::vector&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::matchConvert(const Mat&, const Mat&, const Mat&, std::vector&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::match(const GpuMat&, std::vector&, const std::vector&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::knnMatchSingle(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, int, const GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::knnMatchDownload(const GpuMat&, const GpuMat&, std::vector< std::vector >&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::knnMatchConvert(const Mat&, const Mat&, std::vector< std::vector >&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::knnMatch(const GpuMat&, const GpuMat&, std::vector< std::vector >&, int, const GpuMat&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::knnMatch2Collection(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, const GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::knnMatch2Download(const GpuMat&, const GpuMat&, const GpuMat&, std::vector< std::vector >&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::knnMatch2Convert(const Mat&, const Mat&, const Mat&, std::vector< std::vector >&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::knnMatch(const GpuMat&, std::vector< std::vector >&, int, const std::vector&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::radiusMatchSingle(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, float, const GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::radiusMatchDownload(const GpuMat&, const GpuMat&, const GpuMat&, std::vector< std::vector >&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::radiusMatchConvert(const Mat&, const Mat&, const Mat&, std::vector< std::vector >&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::radiusMatch(const GpuMat&, const GpuMat&, std::vector< std::vector >&, float, const GpuMat&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::radiusMatchCollection(const GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, float, const std::vector&, Stream&) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::radiusMatchDownload(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, std::vector< std::vector >&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::radiusMatchConvert(const Mat&, const Mat&, const Mat&, const Mat&, std::vector< std::vector >&, bool) { throw_no_cuda(); } +void cv::gpu::BFMatcher_GPU::radiusMatch(const GpuMat&, std::vector< std::vector >&, float, const std::vector&, bool) { throw_no_cuda(); } #else /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/calib3d.cpp b/modules/gpu/src/calib3d.cpp index 78408c73b..8c8f753b6 100644 --- a/modules/gpu/src/calib3d.cpp +++ b/modules/gpu/src/calib3d.cpp @@ -47,11 +47,11 @@ using namespace cv::gpu; #if !defined HAVE_CUDA || defined(CUDA_DISABLER) -void cv::gpu::transformPoints(const GpuMat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::transformPoints(const GpuMat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_no_cuda(); } -void cv::gpu::projectPoints(const GpuMat&, const Mat&, const Mat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::projectPoints(const GpuMat&, const Mat&, const Mat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_no_cuda(); } -void cv::gpu::solvePnPRansac(const Mat&, const Mat&, const Mat&, const Mat&, Mat&, Mat&, bool, int, float, int, std::vector*) { throw_nogpu(); } +void cv::gpu::solvePnPRansac(const Mat&, const Mat&, const Mat&, const Mat&, Mat&, Mat&, bool, int, float, int, std::vector*) { throw_no_cuda(); } #else diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp index 3ba8368d1..4ccc86a69 100644 --- a/modules/gpu/src/cascadeclassifier.cpp +++ b/modules/gpu/src/cascadeclassifier.cpp @@ -49,15 +49,15 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) -cv::gpu::CascadeClassifier_GPU::CascadeClassifier_GPU() { throw_nogpu(); } -cv::gpu::CascadeClassifier_GPU::CascadeClassifier_GPU(const String&) { throw_nogpu(); } -cv::gpu::CascadeClassifier_GPU::~CascadeClassifier_GPU() { throw_nogpu(); } -bool cv::gpu::CascadeClassifier_GPU::empty() const { throw_nogpu(); return true; } -bool cv::gpu::CascadeClassifier_GPU::load(const String&) { throw_nogpu(); return true; } -Size cv::gpu::CascadeClassifier_GPU::getClassifierSize() const { throw_nogpu(); return Size();} -void cv::gpu::CascadeClassifier_GPU::release() { throw_nogpu(); } -int cv::gpu::CascadeClassifier_GPU::detectMultiScale( const GpuMat&, GpuMat&, double, int, Size) {throw_nogpu(); return -1;} -int cv::gpu::CascadeClassifier_GPU::detectMultiScale( const GpuMat&, GpuMat&, Size, Size, double, int) {throw_nogpu(); return -1;} +cv::gpu::CascadeClassifier_GPU::CascadeClassifier_GPU() { throw_no_cuda(); } +cv::gpu::CascadeClassifier_GPU::CascadeClassifier_GPU(const String&) { throw_no_cuda(); } +cv::gpu::CascadeClassifier_GPU::~CascadeClassifier_GPU() { throw_no_cuda(); } +bool cv::gpu::CascadeClassifier_GPU::empty() const { throw_no_cuda(); return true; } +bool cv::gpu::CascadeClassifier_GPU::load(const String&) { throw_no_cuda(); return true; } +Size cv::gpu::CascadeClassifier_GPU::getClassifierSize() const { throw_no_cuda(); return Size();} +void cv::gpu::CascadeClassifier_GPU::release() { throw_no_cuda(); } +int cv::gpu::CascadeClassifier_GPU::detectMultiScale( const GpuMat&, GpuMat&, double, int, Size) {throw_no_cuda(); return -1;} +int cv::gpu::CascadeClassifier_GPU::detectMultiScale( const GpuMat&, GpuMat&, Size, Size, double, int) {throw_no_cuda(); return -1;} #else @@ -403,7 +403,7 @@ public: unsigned int classified = 0; GpuMat dclassified(1, 1, CV_32S); - cudaSafeCall( cudaMemcpy(dclassified.ptr(), &classified, sizeof(int), cudaMemcpyHostToDevice) ); + cvCudaSafeCall( cudaMemcpy(dclassified.ptr(), &classified, sizeof(int), cudaMemcpyHostToDevice) ); PyrLavel level(0, 1.0f, image.size(), NxM, minObjectSize); @@ -448,11 +448,11 @@ public: if (groupThreshold <= 0 || objects.empty()) return 0; - cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); + cvCudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); cuda::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified.ptr()); - cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); + cvCudaSafeCall( cudaDeviceSynchronize() ); return classified; } @@ -481,7 +481,7 @@ private: roiSize.height = frame.height; cudaDeviceProp prop; - cudaSafeCall( cudaGetDeviceProperties(&prop, cv::gpu::getDevice()) ); + cvCudaSafeCall( cudaGetDeviceProperties(&prop, cv::gpu::getDevice()) ); Ncv32u bufSize; ncvSafeCall( nppiStIntegralGetSize_8u32u(roiSize, &bufSize, prop) ); diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp index 9f6ca2598..33d57360c 100644 --- a/modules/gpu/src/color.cpp +++ b/modules/gpu/src/color.cpp @@ -47,10 +47,10 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) -void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, Stream&) { throw_nogpu(); } -void cv::gpu::demosaicing(const GpuMat&, GpuMat&, int, int, Stream&) { throw_nogpu(); } -void cv::gpu::swapChannels(GpuMat&, const int[], Stream&) { throw_nogpu(); } -void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_nogpu(); } +void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, Stream&) { throw_no_cuda(); } +void cv::gpu::demosaicing(const GpuMat&, GpuMat&, int, int, Stream&) { throw_no_cuda(); } +void cv::gpu::swapChannels(GpuMat&, const int[], Stream&) { throw_no_cuda(); } +void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_no_cuda(); } #else /* !defined (HAVE_CUDA) */ @@ -1600,7 +1600,7 @@ namespace nppSafeCall( nppiAlphaPremul_16u_AC4R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI) ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); #endif } @@ -1942,7 +1942,7 @@ void cv::gpu::swapChannels(GpuMat& image, const int dstOrder[4], Stream& s) nppSafeCall( nppiSwapChannels_8u_C4IR(image.ptr(), static_cast(image.step), sz, dstOrder) ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } void cv::gpu::gammaCorrection(const GpuMat& src, GpuMat& dst, bool forward, Stream& stream) diff --git a/modules/gpu/src/cuda/NV12ToARGB.cu b/modules/gpu/src/cuda/NV12ToARGB.cu index 4110dbc2a..cccf6eb46 100644 --- a/modules/gpu/src/cuda/NV12ToARGB.cu +++ b/modules/gpu/src/cuda/NV12ToARGB.cu @@ -60,7 +60,7 @@ namespace cv { namespace gpu { namespace cuda { void loadHueCSC(float hueCSC[9]) { - cudaSafeCall( cudaMemcpyToSymbol(constHueColorSpaceMat, hueCSC, 9 * sizeof(float)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(constHueColorSpaceMat, hueCSC, 9 * sizeof(float)) ); } __device__ void YUV2RGB(const uint* yuvi, float* red, float* green, float* blue) @@ -190,10 +190,10 @@ namespace cv { namespace gpu { namespace cuda { NV12ToARGB<<>>(decodedFrame.data, decodedFrame.step, interopFrame.data, interopFrame.step, interopFrame.cols, interopFrame.rows); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } } }}} diff --git a/modules/gpu/src/cuda/bf_knnmatch.cu b/modules/gpu/src/cuda/bf_knnmatch.cu index 629ac694a..f03b1c499 100644 --- a/modules/gpu/src/cuda/bf_knnmatch.cu +++ b/modules/gpu/src/cuda/bf_knnmatch.cu @@ -417,10 +417,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (BLOCK_SIZE * (MAX_DESC_LEN >= BLOCK_SIZE ? MAX_DESC_LEN : BLOCK_SIZE) + BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolledCached<<>>(query, train, mask, trainIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -478,10 +478,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (BLOCK_SIZE * (MAX_DESC_LEN >= 2 * BLOCK_SIZE ? MAX_DESC_LEN : 2 * BLOCK_SIZE) + BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolledCached<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -594,10 +594,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolled<<>>(query, train, mask, trainIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -653,10 +653,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolled<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -768,10 +768,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); match<<>>(query, train, mask, trainIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -827,10 +827,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); match<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -959,10 +959,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); calcDistanceUnrolled<<>>(query, train, mask, allDist); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -1022,10 +1022,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); calcDistance<<>>(query, train, mask, allDist); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -1115,11 +1115,11 @@ namespace cv { namespace gpu { namespace cuda for (int i = 0; i < k; ++i) { findBestMatch<<>>(allDist, i, trainIdx, distance); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); } if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } void findKnnMatchDispatcher(int k, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, cudaStream_t stream) diff --git a/modules/gpu/src/cuda/bf_match.cu b/modules/gpu/src/cuda/bf_match.cu index 1e7cf8ab8..2ecb48ead 100644 --- a/modules/gpu/src/cuda/bf_match.cu +++ b/modules/gpu/src/cuda/bf_match.cu @@ -177,10 +177,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (BLOCK_SIZE * (MAX_DESC_LEN >= BLOCK_SIZE ? MAX_DESC_LEN : BLOCK_SIZE) + BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolledCached<<>>(query, train, mask, trainIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -236,10 +236,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (BLOCK_SIZE * (MAX_DESC_LEN >= 2 * BLOCK_SIZE ? MAX_DESC_LEN : 2 * BLOCK_SIZE) + BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolledCached<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -335,10 +335,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolled<<>>(query, train, mask, trainIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -392,10 +392,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolled<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -490,10 +490,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); match<<>>(query, train, mask, trainIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -546,10 +546,10 @@ namespace cv { namespace gpu { namespace cuda const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); match<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// diff --git a/modules/gpu/src/cuda/bf_radius_match.cu b/modules/gpu/src/cuda/bf_radius_match.cu index 1c0857748..56b99faf5 100644 --- a/modules/gpu/src/cuda/bf_radius_match.cu +++ b/modules/gpu/src/cuda/bf_radius_match.cu @@ -122,10 +122,10 @@ namespace cv { namespace gpu { namespace cuda matchUnrolled<<>>(query, 0, train, maxDistance, mask, trainIdx, PtrStepi(), distance, nMatches.data, trainIdx.cols); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -153,11 +153,11 @@ namespace cv { namespace gpu { namespace cuda matchUnrolled<<>>(query, i, train, maxDistance, WithOutMask(), trainIdx, imgIdx, distance, nMatches.data, trainIdx.cols); } - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); } if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -230,10 +230,10 @@ namespace cv { namespace gpu { namespace cuda match<<>>(query, 0, train, maxDistance, mask, trainIdx, PtrStepi(), distance, nMatches.data, trainIdx.cols); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -261,11 +261,11 @@ namespace cv { namespace gpu { namespace cuda match<<>>(query, i, train, maxDistance, WithOutMask(), trainIdx, imgIdx, distance, nMatches.data, trainIdx.cols); } - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); } if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// diff --git a/modules/gpu/src/cuda/bgfg_gmg.cu b/modules/gpu/src/cuda/bgfg_gmg.cu index 0047fe53e..14e23df94 100644 --- a/modules/gpu/src/cuda/bgfg_gmg.cu +++ b/modules/gpu/src/cuda/bgfg_gmg.cu @@ -62,15 +62,15 @@ namespace cv { namespace gpu { namespace cuda { void loadConstants(int width, int height, float minVal, float maxVal, int quantizationLevels, float backgroundPrior, float decisionThreshold, int maxFeatures, int numInitializationFrames) { - cudaSafeCall( cudaMemcpyToSymbol(c_width, &width, sizeof(width)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_height, &height, sizeof(height)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_minVal, &minVal, sizeof(minVal)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_maxVal, &maxVal, sizeof(maxVal)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_quantizationLevels, &quantizationLevels, sizeof(quantizationLevels)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_backgroundPrior, &backgroundPrior, sizeof(backgroundPrior)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_decisionThreshold, &decisionThreshold, sizeof(decisionThreshold)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_maxFeatures, &maxFeatures, sizeof(maxFeatures)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_numInitializationFrames, &numInitializationFrames, sizeof(numInitializationFrames)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_width, &width, sizeof(width)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_height, &height, sizeof(height)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_minVal, &minVal, sizeof(minVal)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_maxVal, &maxVal, sizeof(maxVal)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_quantizationLevels, &quantizationLevels, sizeof(quantizationLevels)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_backgroundPrior, &backgroundPrior, sizeof(backgroundPrior)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_decisionThreshold, &decisionThreshold, sizeof(decisionThreshold)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_maxFeatures, &maxFeatures, sizeof(maxFeatures)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_numInitializationFrames, &numInitializationFrames, sizeof(numInitializationFrames)) ); } __device__ float findFeature(const int color, const PtrStepi& colors, const PtrStepf& weights, const int x, const int y, const int nfeatures) @@ -230,14 +230,14 @@ namespace cv { namespace gpu { namespace cuda { const dim3 block(32, 8); const dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); - cudaSafeCall( cudaFuncSetCacheConfig(update, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(update, cudaFuncCachePreferL1) ); update<<>>((PtrStepSz) frame, fgmask, colors, weights, nfeatures, frameNum, learningRate, updateBackgroundModel); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template void update_gpu(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); diff --git a/modules/gpu/src/cuda/bgfg_mog.cu b/modules/gpu/src/cuda/bgfg_mog.cu index 6a514f7d3..771f367ce 100644 --- a/modules/gpu/src/cuda/bgfg_mog.cu +++ b/modules/gpu/src/cuda/bgfg_mog.cu @@ -180,16 +180,16 @@ namespace cv { namespace gpu { namespace cuda dim3 block(32, 8); dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); - cudaSafeCall( cudaFuncSetCacheConfig(mog_withoutLearning, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(mog_withoutLearning, cudaFuncCachePreferL1) ); mog_withoutLearning<<>>((PtrStepSz) frame, fgmask, weight, (PtrStepSz) mean, (PtrStepSz) var, nmixtures, varThreshold, backgroundRatio); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////// @@ -333,16 +333,16 @@ namespace cv { namespace gpu { namespace cuda dim3 block(32, 8); dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); - cudaSafeCall( cudaFuncSetCacheConfig(mog_withLearning, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(mog_withLearning, cudaFuncCachePreferL1) ); mog_withLearning<<>>((PtrStepSz) frame, fgmask, weight, sortKey, (PtrStepSz) mean, (PtrStepSz) var, nmixtures, varThreshold, backgroundRatio, learningRate, minVar); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////// @@ -406,13 +406,13 @@ namespace cv { namespace gpu { namespace cuda dim3 block(32, 8); dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); - cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage, cudaFuncCachePreferL1) ); getBackgroundImage<<>>(weight, (PtrStepSz) mean, (PtrStepSz) dst, nmixtures, backgroundRatio); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } void getBackgroundImage_gpu(int cn, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, int nmixtures, float backgroundRatio, cudaStream_t stream) @@ -445,15 +445,15 @@ namespace cv { namespace gpu { namespace cuda varMin = ::fminf(varMin, varMax); varMax = ::fmaxf(varMin, varMax); - cudaSafeCall( cudaMemcpyToSymbol(c_nmixtures, &nmixtures, sizeof(int)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_Tb, &Tb, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_TB, &TB, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_Tg, &Tg, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_varInit, &varInit, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_varMin, &varMin, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_varMax, &varMax, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_tau, &tau, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(c_shadowVal, &shadowVal, sizeof(unsigned char)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_nmixtures, &nmixtures, sizeof(int)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_Tb, &Tb, sizeof(float)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_TB, &TB, sizeof(float)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_Tg, &Tg, sizeof(float)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_varInit, &varInit, sizeof(float)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_varMin, &varMin, sizeof(float)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_varMax, &varMax, sizeof(float)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_tau, &tau, sizeof(float)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(c_shadowVal, &shadowVal, sizeof(unsigned char)) ); } template @@ -665,7 +665,7 @@ namespace cv { namespace gpu { namespace cuda if (detectShadows) { - cudaSafeCall( cudaFuncSetCacheConfig(mog2, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(mog2, cudaFuncCachePreferL1) ); mog2<<>>((PtrStepSz) frame, fgmask, modesUsed, weight, variance, (PtrStepSz) mean, @@ -673,17 +673,17 @@ namespace cv { namespace gpu { namespace cuda } else { - cudaSafeCall( cudaFuncSetCacheConfig(mog2, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(mog2, cudaFuncCachePreferL1) ); mog2<<>>((PtrStepSz) frame, fgmask, modesUsed, weight, variance, (PtrStepSz) mean, alphaT, alpha1, prune); } - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } void mog2_gpu(PtrStepSzb frame, int cn, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean, @@ -737,13 +737,13 @@ namespace cv { namespace gpu { namespace cuda dim3 block(32, 8); dim3 grid(divUp(modesUsed.cols, block.x), divUp(modesUsed.rows, block.y)); - cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage2, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage2, cudaFuncCachePreferL1) ); getBackgroundImage2<<>>(modesUsed, weight, (PtrStepSz) mean, (PtrStepSz) dst); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } void getBackgroundImage2_gpu(int cn, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, cudaStream_t stream) diff --git a/modules/gpu/src/cuda/bilateral_filter.cu b/modules/gpu/src/cuda/bilateral_filter.cu index 17c7c1d8e..8a6f78111 100644 --- a/modules/gpu/src/cuda/bilateral_filter.cu +++ b/modules/gpu/src/cuda/bilateral_filter.cu @@ -135,12 +135,12 @@ namespace cv { namespace gpu { namespace cuda float sigma_spatial2_inv_half = -0.5f/(sigma_spatial * sigma_spatial); float sigma_color2_inv_half = -0.5f/(sigma_color * sigma_color); - cudaSafeCall( cudaFuncSetCacheConfig (bilateral_kernel >, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig (bilateral_kernel >, cudaFuncCachePreferL1) ); bilateral_kernel<<>>((PtrStepSz)src, (PtrStepSz)dst, b, kernel_size, sigma_spatial2_inv_half, sigma_color2_inv_half); - cudaSafeCall ( cudaGetLastError () ); + cvCudaSafeCall ( cudaGetLastError () ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template diff --git a/modules/gpu/src/cuda/blend.cu b/modules/gpu/src/cuda/blend.cu index b4ecfbb7f..f5f734cc5 100644 --- a/modules/gpu/src/cuda/blend.cu +++ b/modules/gpu/src/cuda/blend.cu @@ -73,10 +73,10 @@ namespace cv { namespace gpu { namespace cuda dim3 grid(divUp(cols * cn, threads.x), divUp(rows, threads.y)); blendLinearKernel<<>>(rows, cols * cn, cn, img1, img2, weights1, weights2, result); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall(cudaDeviceSynchronize()); + cvCudaSafeCall(cudaDeviceSynchronize()); } template void blendLinearCaller(int, int, int, PtrStep, PtrStep, PtrStepf, PtrStepf, PtrStep, cudaStream_t stream); @@ -109,10 +109,10 @@ namespace cv { namespace gpu { namespace cuda dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); blendLinearKernel8UC4<<>>(rows, cols, img1, img2, weights1, weights2, result); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall(cudaDeviceSynchronize()); + cvCudaSafeCall(cudaDeviceSynchronize()); } } // namespace blend }}} // namespace cv { namespace gpu { namespace cuda diff --git a/modules/gpu/src/cuda/calib3d.cu b/modules/gpu/src/cuda/calib3d.cu index 5d12405d5..f8882cfc2 100644 --- a/modules/gpu/src/cuda/calib3d.cu +++ b/modules/gpu/src/cuda/calib3d.cu @@ -75,10 +75,10 @@ namespace cv { namespace gpu { namespace cuda const float* transl, PtrStepSz dst, cudaStream_t stream) { - cudaSafeCall(cudaMemcpyToSymbol(crot0, rot, sizeof(float) * 3)); - cudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3)); - cudaSafeCall(cudaMemcpyToSymbol(crot2, rot + 6, sizeof(float) * 3)); - cudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(crot0, rot, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(crot2, rot + 6, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3)); cv::gpu::cuda::transform(src, dst, TransformOp(), WithOutMask(), stream); } } // namespace transform_points @@ -114,12 +114,12 @@ namespace cv { namespace gpu { namespace cuda const float* transl, const float* proj, PtrStepSz dst, cudaStream_t stream) { - cudaSafeCall(cudaMemcpyToSymbol(crot0, rot, sizeof(float) * 3)); - cudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3)); - cudaSafeCall(cudaMemcpyToSymbol(crot2, rot + 6, sizeof(float) * 3)); - cudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3)); - cudaSafeCall(cudaMemcpyToSymbol(cproj0, proj, sizeof(float) * 3)); - cudaSafeCall(cudaMemcpyToSymbol(cproj1, proj + 3, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(crot0, rot, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(crot2, rot + 6, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(cproj0, proj, sizeof(float) * 3)); + cvCudaSafeCall(cudaMemcpyToSymbol(cproj1, proj + 3, sizeof(float) * 3)); cv::gpu::cuda::transform(src, dst, ProjectOp(), WithOutMask(), stream); } } // namespace project_points @@ -174,17 +174,17 @@ namespace cv { namespace gpu { namespace cuda const float3* transl_vectors, const float3* object, const float2* image, const float dist_threshold, int* hypothesis_scores) { - cudaSafeCall(cudaMemcpyToSymbol(crot_matrices, rot_matrices, num_hypotheses * 3 * sizeof(float3))); - cudaSafeCall(cudaMemcpyToSymbol(ctransl_vectors, transl_vectors, num_hypotheses * sizeof(float3))); + cvCudaSafeCall(cudaMemcpyToSymbol(crot_matrices, rot_matrices, num_hypotheses * 3 * sizeof(float3))); + cvCudaSafeCall(cudaMemcpyToSymbol(ctransl_vectors, transl_vectors, num_hypotheses * sizeof(float3))); dim3 threads(256); dim3 grid(num_hypotheses); computeHypothesisScoresKernel<256><<>>( num_points, object, image, dist_threshold, hypothesis_scores); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } } // namespace solvepnp_ransac }}} // namespace cv { namespace gpu { namespace cuda diff --git a/modules/gpu/src/cuda/canny.cu b/modules/gpu/src/cuda/canny.cu index 151f234f0..62c7242ed 100644 --- a/modules/gpu/src/cuda/canny.cu +++ b/modules/gpu/src/cuda/canny.cu @@ -141,9 +141,9 @@ namespace canny calcMagnitudeKernel<<>>(src, dx, dy, mag, norm); } - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); - cudaSafeCall(cudaThreadSynchronize()); + cvCudaSafeCall(cudaThreadSynchronize()); } void calcMagnitude(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, bool L2Grad) @@ -227,9 +227,9 @@ namespace canny bindTexture(&tex_mag, mag); calcMapKernel<<>>(dx, dy, map, low_thresh, high_thresh); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } } @@ -324,17 +324,17 @@ namespace canny void edgesHysteresisLocal(PtrStepSzi map, ushort2* st1) { void* counter_ptr; - cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, counter) ); + cvCudaSafeCall( cudaGetSymbolAddress(&counter_ptr, counter) ); - cudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) ); + cvCudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) ); const dim3 block(16, 16); const dim3 grid(divUp(map.cols, block.x), divUp(map.rows, block.y)); edgesHysteresisLocalKernel<<>>(map, st1); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } } @@ -435,24 +435,24 @@ namespace canny void edgesHysteresisGlobal(PtrStepSzi map, ushort2* st1, ushort2* st2) { void* counter_ptr; - cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, canny::counter) ); + cvCudaSafeCall( cudaGetSymbolAddress(&counter_ptr, canny::counter) ); int count; - cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); + cvCudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); while (count > 0) { - cudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) ); + cvCudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) ); const dim3 block(128); const dim3 grid(::min(count, 65535u), divUp(count, 65535), 1); edgesHysteresisGlobalKernel<<>>(map, st1, st2, count); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); - cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); + cvCudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); std::swap(st1, st2); } diff --git a/modules/gpu/src/cuda/ccomponetns.cu b/modules/gpu/src/cuda/ccomponetns.cu index 2e52ff2bf..45928a084 100644 --- a/modules/gpu/src/cuda/ccomponetns.cu +++ b/modules/gpu/src/cuda/ccomponetns.cu @@ -215,9 +215,9 @@ namespace cv { namespace gpu { namespace cuda Int_t inInt(lo, hi); computeConnectivity<<>>(static_cast >(image), edges, inInt); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template void computeEdges (const PtrStepSzb& image, PtrStepSzb edges, const float4& lo, const float4& hi, cudaStream_t stream); @@ -503,7 +503,7 @@ namespace cv { namespace gpu { namespace cuda dim3 grid(divUp(edges.cols, TILE_COLS), divUp(edges.rows, TILE_ROWS)); lableTiles<<>>(edges, comps); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); int tileSizeX = TILE_COLS, tileSizeY = TILE_ROWS; while (grid.x > 1 || grid.y > 1) @@ -517,16 +517,16 @@ namespace cv { namespace gpu { namespace cuda tileSizeY <<= 1; grid = mergeGrid; - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); } grid.x = divUp(edges.cols, block.x); grid.y = divUp(edges.rows, block.y); flatten<<>>(edges, comps); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } } } } } diff --git a/modules/gpu/src/cuda/clahe.cu b/modules/gpu/src/cuda/clahe.cu index c5ecbf4e8..340e15b45 100644 --- a/modules/gpu/src/cuda/clahe.cu +++ b/modules/gpu/src/cuda/clahe.cu @@ -128,10 +128,10 @@ namespace clahe calcLutKernel<<>>(src, lut, tileSize, tilesX, clipLimit, lutScale); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } __global__ void tranformKernel(const PtrStepSzb src, PtrStepb dst, const PtrStepb lut, const int2 tileSize, const int tilesX, const int tilesY) @@ -173,13 +173,13 @@ namespace clahe const dim3 block(32, 8); const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y)); - cudaSafeCall( cudaFuncSetCacheConfig(tranformKernel, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(tranformKernel, cudaFuncCachePreferL1) ); tranformKernel<<>>(src, dst, lut, tileSize, tilesX, tilesY); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } } diff --git a/modules/gpu/src/cuda/column_filter.h b/modules/gpu/src/cuda/column_filter.h index 00278103e..1f4276740 100644 --- a/modules/gpu/src/cuda/column_filter.h +++ b/modules/gpu/src/cuda/column_filter.h @@ -169,10 +169,10 @@ namespace column_filter linearColumnFilter<<>>(src, dst, anchor, brd); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } } @@ -363,9 +363,9 @@ namespace filter }; if (stream == 0) - cudaSafeCall( cudaMemcpyToSymbol(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); + cvCudaSafeCall( cudaMemcpyToSymbol(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); else - cudaSafeCall( cudaMemcpyToSymbolAsync(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); + cvCudaSafeCall( cudaMemcpyToSymbolAsync(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); callers[brd_type][ksize]((PtrStepSz)src, (PtrStepSz)dst, anchor, cc, stream); } diff --git a/modules/gpu/src/cuda/copy_make_border.cu b/modules/gpu/src/cuda/copy_make_border.cu index 43544658f..30bbe0f40 100644 --- a/modules/gpu/src/cuda/copy_make_border.cu +++ b/modules/gpu/src/cuda/copy_make_border.cu @@ -70,10 +70,10 @@ namespace cv { namespace gpu { namespace cuda BorderReader< PtrStep, B > brdSrc(src, brd); copyMakeBorder<<>>(brdSrc, dst, top, left); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } }; diff --git a/modules/gpu/src/cuda/debayer.cu b/modules/gpu/src/cuda/debayer.cu index a079bd583..6ade873f3 100644 --- a/modules/gpu/src/cuda/debayer.cu +++ b/modules/gpu/src/cuda/debayer.cu @@ -347,13 +347,13 @@ namespace cv { namespace gpu { namespace cuda const dim3 block(32, 8); const dim3 grid(divUp(src.cols, 4 * block.x), divUp(src.rows, block.y)); - cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u, cudaFuncCachePreferL1) ); Bayer2BGR_8u<<>>(src, (PtrStepSz)dst, blue_last, start_with_green); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template @@ -364,13 +364,13 @@ namespace cv { namespace gpu { namespace cuda const dim3 block(32, 8); const dim3 grid(divUp(src.cols, 2 * block.x), divUp(src.rows, block.y)); - cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u, cudaFuncCachePreferL1) ); + cvCudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u, cudaFuncCachePreferL1) ); Bayer2BGR_16u<<>>(src, (PtrStepSz)dst, blue_last, start_with_green); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template void Bayer2BGR_8u_gpu<1>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); @@ -530,10 +530,10 @@ namespace cv { namespace gpu { namespace cuda bindTexture(&sourceTex, src); MHCdemosaic<<>>((PtrStepSz)dst, sourceOffset, firstRed); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template void MHCdemosaic<1>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream); diff --git a/modules/gpu/src/cuda/disp_bilateral_filter.cu b/modules/gpu/src/cuda/disp_bilateral_filter.cu index c3edff320..7758e4ce0 100644 --- a/modules/gpu/src/cuda/disp_bilateral_filter.cu +++ b/modules/gpu/src/cuda/disp_bilateral_filter.cu @@ -61,16 +61,16 @@ namespace cv { namespace gpu { namespace cuda void disp_load_constants(float* table_color, PtrStepSzf table_space, int ndisp, int radius, short edge_disc, short max_disc) { - cudaSafeCall( cudaMemcpyToSymbol(ctable_color, &table_color, sizeof(table_color)) ); - cudaSafeCall( cudaMemcpyToSymbol(ctable_space, &table_space.data, sizeof(table_space.data)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(ctable_color, &table_color, sizeof(table_color)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(ctable_space, &table_space.data, sizeof(table_space.data)) ); size_t table_space_step = table_space.step / sizeof(float); - cudaSafeCall( cudaMemcpyToSymbol(ctable_space_step, &table_space_step, sizeof(size_t)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(ctable_space_step, &table_space_step, sizeof(size_t)) ); - cudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); - cudaSafeCall( cudaMemcpyToSymbol(cradius, &radius, sizeof(int)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(cradius, &radius, sizeof(int)) ); - cudaSafeCall( cudaMemcpyToSymbol(cedge_disc, &edge_disc, sizeof(short)) ); - cudaSafeCall( cudaMemcpyToSymbol(cmax_disc, &max_disc, sizeof(short)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(cedge_disc, &edge_disc, sizeof(short)) ); + cvCudaSafeCall( cudaMemcpyToSymbol(cmax_disc, &max_disc, sizeof(short)) ); } template @@ -191,28 +191,28 @@ namespace cv { namespace gpu { namespace cuda for (int i = 0; i < iters; ++i) { disp_bilateral_filter<1><<>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); disp_bilateral_filter<1><<>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); } break; case 3: for (int i = 0; i < iters; ++i) { disp_bilateral_filter<3><<>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); disp_bilateral_filter<3><<>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); - cudaSafeCall( cudaGetLastError() ); + cvCudaSafeCall( cudaGetLastError() ); } break; default: - cv::gpu::error("Unsupported channels count", __FILE__, __LINE__, "disp_bilateral_filter"); + CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); + cvCudaSafeCall( cudaDeviceSynchronize() ); } template void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); diff --git a/modules/gpu/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu index 1d3b7ca4c..edd17aea7 100644 --- a/modules/gpu/src/cuda/element_operations.cu +++ b/modules/gpu/src/cuda/element_operations.cu @@ -216,21 +216,21 @@ namespace arithm { void addMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VAdd4(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VAdd4(), WithOutMask(), stream); } void addMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VAdd2(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VAdd2(), WithOutMask(), stream); } template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) { if (mask.data) - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), mask, stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), mask, stream); else - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), WithOutMask(), stream); } template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); @@ -323,9 +323,9 @@ namespace arithm AddScalar op(static_cast(val)); if (mask.data) - transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); else - transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); @@ -451,21 +451,21 @@ namespace arithm { void subMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VSub4(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VSub4(), WithOutMask(), stream); } void subMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VSub2(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VSub2(), WithOutMask(), stream); } template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) { if (mask.data) - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), mask, stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), mask, stream); else - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), WithOutMask(), stream); } template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); @@ -536,9 +536,9 @@ namespace arithm AddScalar op(-static_cast(val)); if (mask.data) - transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); else - transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void subScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); @@ -676,12 +676,12 @@ namespace arithm { void mulMat_8uc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, Mul_8uc4_32f(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, Mul_8uc4_32f(), WithOutMask(), stream); } void mulMat_16sc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, Mul_16sc4_32f(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, Mul_16sc4_32f(), WithOutMask(), stream); } template @@ -690,12 +690,12 @@ namespace arithm if (scale == 1) { Mul op; - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); } else { MulScale op(static_cast(scale)); - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); } } @@ -787,7 +787,7 @@ namespace arithm void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream) { MulScalar op(static_cast(val)); - transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); @@ -944,12 +944,12 @@ namespace arithm { void divMat_8uc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, Div_8uc4_32f(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, Div_8uc4_32f(), WithOutMask(), stream); } void divMat_16sc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, Div_16sc4_32f(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, Div_16sc4_32f(), WithOutMask(), stream); } template @@ -958,12 +958,12 @@ namespace arithm if (scale == 1) { Div op; - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); } else { DivScale op(static_cast(scale)); - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); } } @@ -1033,7 +1033,7 @@ namespace arithm void divScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream) { MulScalar op(static_cast(1.0 / val)); - transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void divScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); @@ -1124,7 +1124,7 @@ namespace arithm void divInv(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream) { DivInv op(static_cast(val)); - transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void divInv(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); @@ -1263,18 +1263,18 @@ namespace arithm { void absDiffMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VAbsDiff4(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VAbsDiff4(), WithOutMask(), stream); } void absDiffMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VAbsDiff2(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VAbsDiff2(), WithOutMask(), stream); } template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AbsDiffMat(), WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AbsDiffMat(), WithOutMask(), stream); } template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); @@ -1319,7 +1319,7 @@ namespace arithm { AbsDiffScalar op(static_cast(val)); - transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); @@ -1346,7 +1346,7 @@ namespace arithm template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src, (PtrStepSz) dst, abs_func(), WithOutMask(), stream); + cuda::transform((PtrStepSz) src, (PtrStepSz) dst, abs_func(), WithOutMask(), stream); } template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1387,7 +1387,7 @@ namespace arithm template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src, (PtrStepSz) dst, Sqr(), WithOutMask(), stream); + cuda::transform((PtrStepSz) src, (PtrStepSz) dst, Sqr(), WithOutMask(), stream); } template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1414,7 +1414,7 @@ namespace arithm template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src, (PtrStepSz) dst, sqrt_func(), WithOutMask(), stream); + cuda::transform((PtrStepSz) src, (PtrStepSz) dst, sqrt_func(), WithOutMask(), stream); } template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1441,7 +1441,7 @@ namespace arithm template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src, (PtrStepSz) dst, log_func(), WithOutMask(), stream); + cuda::transform((PtrStepSz) src, (PtrStepSz) dst, log_func(), WithOutMask(), stream); } template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1483,7 +1483,7 @@ namespace arithm template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src, (PtrStepSz) dst, Exp(), WithOutMask(), stream); + cuda::transform((PtrStepSz) src, (PtrStepSz) dst, Exp(), WithOutMask(), stream); } template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1580,26 +1580,26 @@ namespace arithm { void cmpMatEq_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VCmpEq4(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VCmpEq4(), WithOutMask(), stream); } void cmpMatNe_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VCmpNe4(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VCmpNe4(), WithOutMask(), stream); } void cmpMatLt_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VCmpLt4(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VCmpLt4(), WithOutMask(), stream); } void cmpMatLe_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - transform(src1, src2, dst, VCmpLe4(), WithOutMask(), stream); + cuda::transform(src1, src2, dst, VCmpLe4(), WithOutMask(), stream); } template