From b50090f850acc27d30ff1398fd351f1ce432b57d Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 8 Apr 2013 12:37:36 +0400 Subject: [PATCH] restore cudaSafeCall --- .../core/include/opencv2/core/cuda/common.hpp | 12 +- .../core/cuda/detail/transform_detail.hpp | 16 +- modules/core/src/cuda/matrix_operations.cu | 26 +-- modules/core/src/cudastream.cpp | 24 +-- modules/core/src/gpumat.cpp | 50 +++--- modules/core/src/matrix_operations.cpp | 18 +- modules/core/src/opengl_interop.cpp | 20 +-- modules/gpu/src/arithm.cpp | 8 +- modules/gpu/src/cascadeclassifier.cpp | 10 +- modules/gpu/src/color.cpp | 4 +- modules/gpu/src/cuda/NV12ToARGB.cu | 6 +- modules/gpu/src/cuda/bf_knnmatch.cu | 36 ++-- modules/gpu/src/cuda/bf_match.cu | 24 +-- modules/gpu/src/cuda/bf_radius_match.cu | 16 +- modules/gpu/src/cuda/bgfg_gmg.cu | 24 +-- modules/gpu/src/cuda/bgfg_mog.cu | 50 +++--- modules/gpu/src/cuda/bilateral_filter.cu | 6 +- modules/gpu/src/cuda/blend.cu | 8 +- modules/gpu/src/cuda/calib3d.cu | 28 +-- modules/gpu/src/cuda/canny.cu | 28 +-- modules/gpu/src/cuda/ccomponetns.cu | 12 +- modules/gpu/src/cuda/clahe.cu | 10 +- modules/gpu/src/cuda/column_filter.h | 8 +- modules/gpu/src/cuda/copy_make_border.cu | 4 +- modules/gpu/src/cuda/debayer.cu | 16 +- modules/gpu/src/cuda/disp_bilateral_filter.cu | 24 +-- modules/gpu/src/cuda/fast.cu | 20 +-- modules/gpu/src/cuda/fgd_bgfg.cu | 24 +-- modules/gpu/src/cuda/gftt.cu | 10 +- modules/gpu/src/cuda/global_motion.cu | 8 +- modules/gpu/src/cuda/hist.cu | 8 +- modules/gpu/src/cuda/hog.cu | 66 +++---- modules/gpu/src/cuda/hough.cu | 164 +++++++++--------- modules/gpu/src/cuda/imgproc.cu | 94 +++++----- modules/gpu/src/cuda/integral_image.cu | 14 +- modules/gpu/src/cuda/lbp.cu | 2 +- modules/gpu/src/cuda/match_template.cu | 56 +++--- modules/gpu/src/cuda/mathfunc.cu | 8 +- modules/gpu/src/cuda/matrix_reductions.cu | 44 ++--- modules/gpu/src/cuda/nlm.cu | 18 +- modules/gpu/src/cuda/optflowbm.cu | 8 +- modules/gpu/src/cuda/optical_flow.cu | 8 +- .../gpu/src/cuda/optical_flow_farneback.cu | 46 ++--- modules/gpu/src/cuda/orb.cu | 18 +- modules/gpu/src/cuda/pyr_down.cu | 4 +- modules/gpu/src/cuda/pyr_up.cu | 4 +- modules/gpu/src/cuda/pyrlk.cu | 20 +-- modules/gpu/src/cuda/remap.cu | 14 +- modules/gpu/src/cuda/resize.cu | 18 +- modules/gpu/src/cuda/rgb_to_yv12.cu | 8 +- modules/gpu/src/cuda/row_filter.h | 8 +- modules/gpu/src/cuda/split_merge.cu | 24 +-- modules/gpu/src/cuda/stereobm.cu | 32 ++-- modules/gpu/src/cuda/stereobp.cu | 56 +++--- modules/gpu/src/cuda/stereocsbp.cu | 68 ++++---- modules/gpu/src/cuda/tvl1flow.cu | 16 +- modules/gpu/src/cuda/warp.cu | 26 +-- modules/gpu/src/element_operations.cpp | 24 +-- modules/gpu/src/filtering.cpp | 16 +- modules/gpu/src/graphcuts.cpp | 4 +- modules/gpu/src/hough.cpp | 18 +- modules/gpu/src/imgproc.cpp | 22 +-- modules/gpu/src/matrix_reductions.cpp | 12 +- modules/gpu/src/nvidia/core/NCVPyramid.cu | 8 +- modules/gpu/src/optflowbm.cpp | 4 +- modules/gpu/src/optical_flow.cpp | 4 +- modules/gpu/src/resize.cpp | 2 +- modules/gpu/src/warp.cpp | 2 +- modules/nonfree/src/cuda/surf.cu | 54 +++--- modules/nonfree/src/cuda/vibe.cu | 20 +-- modules/nonfree/src/surf_gpu.cpp | 4 +- modules/softcascade/src/detector_cuda.cpp | 2 +- modules/superres/src/cuda/btv_l1_gpu.cu | 14 +- 73 files changed, 807 insertions(+), 805 deletions(-) diff --git a/modules/core/include/opencv2/core/cuda/common.hpp b/modules/core/include/opencv2/core/cuda/common.hpp index acc8adbdfb..774500e649 100644 --- a/modules/core/include/opencv2/core/cuda/common.hpp +++ b/modules/core/include/opencv2/core/cuda/common.hpp @@ -64,10 +64,12 @@ namespace cv { namespace gpu { } }} -#if defined(__GNUC__) - #define cvCudaSafeCall(expr) cv::gpu::checkCudaError(expr, __FILE__, __LINE__, __func__) -#else /* defined(__CUDACC__) || defined(__MSVC__) */ - #define cvCudaSafeCall(expr) cv::gpu::checkCudaError(expr, __FILE__, __LINE__, "") +#ifndef cudaSafeCall + #if defined(__GNUC__) + #define cudaSafeCall(expr) cv::gpu::checkCudaError(expr, __FILE__, __LINE__, __func__) + #else /* defined(__CUDACC__) || defined(__MSVC__) */ + #define cudaSafeCall(expr) cv::gpu::checkCudaError(expr, __FILE__, __LINE__, "") + #endif #endif namespace cv { namespace gpu @@ -104,7 +106,7 @@ namespace cv { namespace gpu template inline void bindTexture(const textureReference* tex, const PtrStepSz& img) { cudaChannelFormatDesc desc = cudaCreateChannelDesc(); - cvCudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) ); + cudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) ); } } }} 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 906e0e0203..2ac309b0c6 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 cudev const dim3 grid(divUp(src.cols, threads.x), divUp(src.rows, threads.y), 1); transformSimple<<>>(src, dst, mask, op); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -332,10 +332,10 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(src1.cols, threads.x), divUp(src1.rows, threads.y), 1); transformSimple<<>>(src1, src2, dst, mask, op); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; template<> struct TransformDispatcher @@ -358,10 +358,10 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(src.cols, threads.x * ft::smart_shift), divUp(src.rows, threads.y), 1); transformSmart<<>>(src, dst, mask, op); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -383,10 +383,10 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(src1.cols, threads.x * ft::smart_shift), divUp(src1.rows, threads.y), 1); transformSmart<<>>(src1, src2, dst, mask, op); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; } // namespace transform_detail diff --git a/modules/core/src/cuda/matrix_operations.cu b/modules/core/src/cuda/matrix_operations.cu index e0b4ebc422..521ee1a2ca 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 cudev void writeScalar(const uchar* vals) { - cvCudaSafeCall( cudaMemcpyToSymbol(scalar_8u, vals, sizeof(uchar) * 4) ); + cudaSafeCall( cudaMemcpyToSymbol(scalar_8u, vals, sizeof(uchar) * 4) ); } void writeScalar(const schar* vals) { - cvCudaSafeCall( cudaMemcpyToSymbol(scalar_8s, vals, sizeof(schar) * 4) ); + cudaSafeCall( cudaMemcpyToSymbol(scalar_8s, vals, sizeof(schar) * 4) ); } void writeScalar(const ushort* vals) { - cvCudaSafeCall( cudaMemcpyToSymbol(scalar_16u, vals, sizeof(ushort) * 4) ); + cudaSafeCall( cudaMemcpyToSymbol(scalar_16u, vals, sizeof(ushort) * 4) ); } void writeScalar(const short* vals) { - cvCudaSafeCall( cudaMemcpyToSymbol(scalar_16s, vals, sizeof(short) * 4) ); + cudaSafeCall( cudaMemcpyToSymbol(scalar_16s, vals, sizeof(short) * 4) ); } void writeScalar(const int* vals) { - cvCudaSafeCall( cudaMemcpyToSymbol(scalar_32s, vals, sizeof(int) * 4) ); + cudaSafeCall( cudaMemcpyToSymbol(scalar_32s, vals, sizeof(int) * 4) ); } void writeScalar(const float* vals) { - cvCudaSafeCall( cudaMemcpyToSymbol(scalar_32f, vals, sizeof(float) * 4) ); + cudaSafeCall( cudaMemcpyToSymbol(scalar_32f, vals, sizeof(float) * 4) ); } void writeScalar(const double* vals) { - cvCudaSafeCall( cudaMemcpyToSymbol(scalar_64f, vals, sizeof(double) * 4) ); + cudaSafeCall( cudaMemcpyToSymbol(scalar_64f, vals, sizeof(double) * 4) ); } template @@ -186,10 +186,10 @@ namespace cv { namespace gpu { namespace cudev 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall ( cudaDeviceSynchronize() ); + cudaSafeCall ( 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 cudev 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall ( cudaDeviceSynchronize() ); + cudaSafeCall ( 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 cudev template void cvt_(PtrStepSzb src, PtrStepSzb dst, double alpha, double beta, cudaStream_t stream) { - cvCudaSafeCall( cudaSetDoubleForDevice(&alpha) ); - cvCudaSafeCall( cudaSetDoubleForDevice(&beta) ); + cudaSafeCall( cudaSetDoubleForDevice(&alpha) ); + cudaSafeCall( cudaSetDoubleForDevice(&beta) ); Convertor op(static_cast(alpha), static_cast(beta)); cv::gpu::cudev::transform((PtrStepSz)src, (PtrStepSz)dst, op, WithOutMask(), stream); } diff --git a/modules/core/src/cudastream.cpp b/modules/core/src/cudastream.cpp index a6d1a41e60..346204dd5e 100644 --- a/modules/core/src/cudastream.cpp +++ b/modules/core/src/cudastream.cpp @@ -131,14 +131,14 @@ bool cv::gpu::Stream::queryIfComplete() if (err == cudaErrorNotReady || err == cudaSuccess) return err == cudaSuccess; - cvCudaSafeCall(err); + cudaSafeCall(err); return false; } void cv::gpu::Stream::waitForCompletion() { cudaStream_t stream = Impl::getStream(impl); - cvCudaSafeCall( cudaStreamSynchronize(stream) ); + cudaSafeCall( cudaStreamSynchronize(stream) ); } void cv::gpu::Stream::enqueueDownload(const GpuMat& src, Mat& dst) @@ -148,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(); - cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) ); + cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) ); } void cv::gpu::Stream::enqueueDownload(const GpuMat& src, CudaMem& dst) @@ -157,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(); - cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) ); + cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) ); } void cv::gpu::Stream::enqueueUpload(const CudaMem& src, GpuMat& dst) @@ -166,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(); - cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) ); + cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) ); } void cv::gpu::Stream::enqueueUpload(const Mat& src, GpuMat& dst) @@ -175,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(); - cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) ); + cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) ); } void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst) @@ -184,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(); - cvCudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToDevice, stream) ); + cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToDevice, stream) ); } void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val) @@ -201,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) { - cvCudaSafeCall( cudaMemset2DAsync(src.data, src.step, 0, src.cols * src.elemSize(), src.rows, stream) ); + cudaSafeCall( cudaMemset2DAsync(src.data, src.step, 0, src.cols * src.elemSize(), src.rows, stream) ); return; } @@ -212,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]); - cvCudaSafeCall( cudaMemset2DAsync(src.data, src.step, ival, src.cols * src.elemSize(), src.rows, stream) ); + cudaSafeCall( cudaMemset2DAsync(src.data, src.step, ival, src.cols * src.elemSize(), src.rows, stream) ); return; } } @@ -299,7 +299,7 @@ void cv::gpu::Stream::enqueueHostCallback(StreamCallback callback, void* userDat cudaStream_t stream = Impl::getStream(impl); - cvCudaSafeCall( cudaStreamAddCallback(stream, cudaStreamCallback, data, 0) ); + cudaSafeCall( cudaStreamAddCallback(stream, cudaStreamCallback, data, 0) ); #else (void) callback; (void) userData; @@ -328,7 +328,7 @@ void cv::gpu::Stream::create() release(); cudaStream_t stream; - cvCudaSafeCall( cudaStreamCreate( &stream ) ); + cudaSafeCall( cudaStreamCreate( &stream ) ); impl = (Stream::Impl*) fastMalloc(sizeof(Stream::Impl)); @@ -340,7 +340,7 @@ void cv::gpu::Stream::release() { if (impl && CV_XADD(&impl->ref_counter, -1) == 1) { - cvCudaSafeCall( cudaStreamDestroy(impl->stream) ); + cudaSafeCall( cudaStreamDestroy(impl->stream) ); cv::fastFree(impl); } } diff --git a/modules/core/src/gpumat.cpp b/modules/core/src/gpumat.cpp index 384adef732..0db8536879 100644 --- a/modules/core/src/gpumat.cpp +++ b/modules/core/src/gpumat.cpp @@ -91,25 +91,25 @@ int cv::gpu::getCudaEnabledDeviceCount() if (error == cudaErrorNoDevice) return 0; - cvCudaSafeCall( error ); + cudaSafeCall( error ); return count; } void cv::gpu::setDevice(int device) { - cvCudaSafeCall( cudaSetDevice( device ) ); + cudaSafeCall( cudaSetDevice( device ) ); } int cv::gpu::getDevice() { int device; - cvCudaSafeCall( cudaGetDevice( &device ) ); + cudaSafeCall( cudaGetDevice( &device ) ); return device; } void cv::gpu::resetDevice() { - cvCudaSafeCall( cudaDeviceReset() ); + cudaSafeCall( cudaDeviceReset() ); } namespace @@ -302,7 +302,7 @@ namespace if (!props_[devID]) { props_[devID] = new cudaDeviceProp; - cvCudaSafeCall( cudaGetDeviceProperties(props_[devID], devID) ); + cudaSafeCall( cudaGetDeviceProperties(props_[devID], devID) ); } return props_[devID]; @@ -322,7 +322,7 @@ void cv::gpu::DeviceInfo::queryMemory(size_t& _totalMemory, size_t& _freeMemory) if (prevDeviceID != device_id_) setDevice(device_id_); - cvCudaSafeCall( cudaMemGetInfo(&_freeMemory, &_totalMemory) ); + cudaSafeCall( cudaMemGetInfo(&_freeMemory, &_totalMemory) ); if (prevDeviceID != device_id_) setDevice(prevDeviceID); @@ -408,8 +408,8 @@ void cv::gpu::printCudaDeviceInfo(int device) printf("Device count: %d\n", count); int driverVersion = 0, runtimeVersion = 0; - cvCudaSafeCall( cudaDriverGetVersion(&driverVersion) ); - cvCudaSafeCall( cudaRuntimeGetVersion(&runtimeVersion) ); + cudaSafeCall( cudaDriverGetVersion(&driverVersion) ); + cudaSafeCall( cudaRuntimeGetVersion(&runtimeVersion) ); const char *computeMode[] = { "Default (multiple host threads can use ::cudaSetDevice() with device simultaneously)", @@ -423,7 +423,7 @@ void cv::gpu::printCudaDeviceInfo(int device) for(int dev = beg; dev < end; ++dev) { cudaDeviceProp prop; - cvCudaSafeCall( cudaGetDeviceProperties(&prop, dev) ); + cudaSafeCall( 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); @@ -485,13 +485,13 @@ void cv::gpu::printShortCudaDeviceInfo(int device) int end = valid ? device+1 : count; int driverVersion = 0, runtimeVersion = 0; - cvCudaSafeCall( cudaDriverGetVersion(&driverVersion) ); - cvCudaSafeCall( cudaRuntimeGetVersion(&runtimeVersion) ); + cudaSafeCall( cudaDriverGetVersion(&driverVersion) ); + cudaSafeCall( cudaRuntimeGetVersion(&runtimeVersion) ); for(int dev = beg; dev < end; ++dev) { cudaDeviceProp prop; - cvCudaSafeCall( cudaGetDeviceProperties(&prop, dev) ); + cudaSafeCall( 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); @@ -983,7 +983,7 @@ namespace nppSafeCall( func(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz) ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; template::func_ptr func> struct NppCvt @@ -998,7 +998,7 @@ namespace nppSafeCall( func(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz, NPP_RND_NEAR) ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -1040,7 +1040,7 @@ namespace nppSafeCall( func(nppS.val, src.ptr(), static_cast(src.step), sz) ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; template::func_ptr func> struct NppSet @@ -1057,7 +1057,7 @@ namespace nppSafeCall( func(nppS[0], src.ptr(), static_cast(src.step), sz) ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -1088,7 +1088,7 @@ namespace nppSafeCall( func(nppS.val, src.ptr(), static_cast(src.step), sz, mask.ptr(), static_cast(mask.step)) ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; template::func_ptr func> struct NppSetMask @@ -1105,7 +1105,7 @@ namespace nppSafeCall( func(nppS[0], src.ptr(), static_cast(src.step), sz, mask.ptr(), static_cast(mask.step)) ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -1131,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)) ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -1148,15 +1148,15 @@ namespace public: void copy(const Mat& src, GpuMat& dst) const { - cvCudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyHostToDevice) ); + cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyHostToDevice) ); } void copy(const GpuMat& src, Mat& dst) const { - cvCudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToHost) ); } void copy(const GpuMat& src, GpuMat& dst) const { - cvCudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToDevice) ); + cudaSafeCall( 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 @@ -1301,7 +1301,7 @@ namespace { if (s[0] == 0.0 && s[1] == 0.0 && s[2] == 0.0 && s[3] == 0.0) { - cvCudaSafeCall( cudaMemset2D(m.data, m.step, 0, m.cols * m.elemSize(), m.rows) ); + cudaSafeCall( cudaMemset2D(m.data, m.step, 0, m.cols * m.elemSize(), m.rows) ); return; } @@ -1312,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]); - cvCudaSafeCall( cudaMemset2D(m.data, m.step, val, m.cols * m.elemSize(), m.rows) ); + cudaSafeCall( cudaMemset2D(m.data, m.step, val, m.cols * m.elemSize(), m.rows) ); return; } } @@ -1367,7 +1367,7 @@ namespace void mallocPitch(void** devPtr, size_t* step, size_t width, size_t height) const { - cvCudaSafeCall( cudaMallocPitch(devPtr, step, width, height) ); + cudaSafeCall( cudaMallocPitch(devPtr, step, width, height) ); } void free(void* devPtr) const diff --git a/modules/core/src/matrix_operations.cpp b/modules/core/src/matrix_operations.cpp index ef09ef62ac..723c38aa04 100644 --- a/modules/core/src/matrix_operations.cpp +++ b/modules/core/src/matrix_operations.cpp @@ -191,18 +191,18 @@ GpuMat cv::gpu::CudaMem::createGpuMatHeader () const { throw_no_cuda(); return G void cv::gpu::registerPageLocked(Mat& m) { - cvCudaSafeCall( cudaHostRegister(m.ptr(), m.step * m.rows, cudaHostRegisterPortable) ); + cudaSafeCall( cudaHostRegister(m.ptr(), m.step * m.rows, cudaHostRegisterPortable) ); } void cv::gpu::unregisterPageLocked(Mat& m) { - cvCudaSafeCall( cudaHostUnregister(m.ptr()) ); + cudaSafeCall( cudaHostUnregister(m.ptr()) ); } bool cv::gpu::CudaMem::canMapHostMemory() { cudaDeviceProp prop; - cvCudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) ); + cudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) ); return (prop.canMapHostMemory != 0) ? true : false; } @@ -237,7 +237,7 @@ void cv::gpu::CudaMem::create(int _rows, int _cols, int _type, int _alloc_type) if (_alloc_type == ALLOC_ZEROCOPY) { cudaDeviceProp prop; - cvCudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) ); + cudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) ); step = alignUpStep(step, prop.textureAlignment); } int64 _nettosize = (int64)step*rows; @@ -252,9 +252,9 @@ void cv::gpu::CudaMem::create(int _rows, int _cols, int _type, int _alloc_type) switch (alloc_type) { - 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; + 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_Error(cv::Error::StsBadFlag, "Invalid alloc type"); } @@ -273,7 +273,7 @@ GpuMat cv::gpu::CudaMem::createGpuMatHeader () const GpuMat res; void *pdev; - cvCudaSafeCall( cudaHostGetDevicePointer( &pdev, data, 0 ) ); + cudaSafeCall( cudaHostGetDevicePointer( &pdev, data, 0 ) ); res = GpuMat(rows, cols, type(), pdev, step); return res; @@ -283,7 +283,7 @@ void cv::gpu::CudaMem::release() { if( refcount && CV_XADD(refcount, -1) == 1 ) { - cvCudaSafeCall( cudaFreeHost(datastart ) ); + cudaSafeCall( 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 19eabfa7d3..7c28d73ba6 100644 --- a/modules/core/src/opengl_interop.cpp +++ b/modules/core/src/opengl_interop.cpp @@ -133,7 +133,7 @@ void cv::gpu::setGlDevice(int device) (void) device; throw_no_cuda(); #else - cvCudaSafeCall( cudaGLSetGLDevice(device) ); + cudaSafeCall( cudaGLSetGLDevice(device) ); #endif #endif } @@ -184,7 +184,7 @@ namespace return; cudaGraphicsResource_t resource; - cvCudaSafeCall( cudaGraphicsGLRegisterBuffer(&resource, buffer, cudaGraphicsMapFlagsNone) ); + cudaSafeCall( cudaGraphicsGLRegisterBuffer(&resource, buffer, cudaGraphicsMapFlagsNone) ); release(); @@ -217,7 +217,7 @@ namespace CudaResource::GraphicsMapHolder::GraphicsMapHolder(cudaGraphicsResource_t* resource, cudaStream_t stream) : resource_(resource), stream_(stream) { if (resource_) - cvCudaSafeCall( cudaGraphicsMapResources(1, resource_, stream_) ); + cudaSafeCall( cudaGraphicsMapResources(1, resource_, stream_) ); } CudaResource::GraphicsMapHolder::~GraphicsMapHolder() @@ -240,14 +240,14 @@ namespace void* dst; size_t size; - cvCudaSafeCall( cudaGraphicsResourceGetMappedPointer(&dst, &size, resource_) ); + cudaSafeCall( cudaGraphicsResourceGetMappedPointer(&dst, &size, resource_) ); CV_DbgAssert( width * height == size ); if (stream == 0) - cvCudaSafeCall( cudaMemcpy2D(dst, width, src, spitch, width, height, cudaMemcpyDeviceToDevice) ); + cudaSafeCall( cudaMemcpy2D(dst, width, src, spitch, width, height, cudaMemcpyDeviceToDevice) ); else - cvCudaSafeCall( cudaMemcpy2DAsync(dst, width, src, spitch, width, height, cudaMemcpyDeviceToDevice, stream) ); + cudaSafeCall( 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) @@ -259,14 +259,14 @@ namespace void* src; size_t size; - cvCudaSafeCall( cudaGraphicsResourceGetMappedPointer(&src, &size, resource_) ); + cudaSafeCall( cudaGraphicsResourceGetMappedPointer(&src, &size, resource_) ); CV_DbgAssert( width * height == size ); if (stream == 0) - cvCudaSafeCall( cudaMemcpy2D(dst, dpitch, src, width, width, height, cudaMemcpyDeviceToDevice) ); + cudaSafeCall( cudaMemcpy2D(dst, dpitch, src, width, width, height, cudaMemcpyDeviceToDevice) ); else - cvCudaSafeCall( cudaMemcpy2DAsync(dst, dpitch, src, width, width, height, cudaMemcpyDeviceToDevice, stream) ); + cudaSafeCall( cudaMemcpy2DAsync(dst, dpitch, src, width, width, height, cudaMemcpyDeviceToDevice, stream) ); } void* CudaResource::map(cudaStream_t stream) @@ -277,7 +277,7 @@ namespace void* ptr; size_t size; - cvCudaSafeCall( cudaGraphicsResourceGetMappedPointer(&ptr, &size, resource_) ); + cudaSafeCall( cudaGraphicsResourceGetMappedPointer(&ptr, &size, resource_) ); h.reset(); diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index 5f32cdf1a8..f60c244c05 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -246,7 +246,7 @@ void cv::gpu::transpose(const GpuMat& src, GpuMat& dst, Stream& s) } if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////// @@ -287,7 +287,7 @@ namespace (flipCode == 0 ? NPP_HORIZONTAL_AXIS : (flipCode > 0 ? NPP_VERTICAL_AXIS : NPP_BOTH_AXIS))) ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; } @@ -402,7 +402,7 @@ void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst, Stream& s) } if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////// @@ -427,7 +427,7 @@ namespace nppSafeCall( func(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz) ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp index 2a65ab520f..770627017d 100644 --- a/modules/gpu/src/cascadeclassifier.cpp +++ b/modules/gpu/src/cascadeclassifier.cpp @@ -403,7 +403,7 @@ public: unsigned int classified = 0; GpuMat dclassified(1, 1, CV_32S); - cvCudaSafeCall( cudaMemcpy(dclassified.ptr(), &classified, sizeof(int), cudaMemcpyHostToDevice) ); + cudaSafeCall( 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; - cvCudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); cudev::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified.ptr()); - cvCudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaDeviceSynchronize() ); return classified; } @@ -481,7 +481,7 @@ private: roiSize.height = frame.height; cudaDeviceProp prop; - cvCudaSafeCall( cudaGetDeviceProperties(&prop, cv::gpu::getDevice()) ); + cudaSafeCall( 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 570efbcb87..5b503814a3 100644 --- a/modules/gpu/src/color.cpp +++ b/modules/gpu/src/color.cpp @@ -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) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 6c307c7fdf..09906613ff 100644 --- a/modules/gpu/src/cuda/NV12ToARGB.cu +++ b/modules/gpu/src/cuda/NV12ToARGB.cu @@ -60,7 +60,7 @@ namespace cv { namespace gpu { namespace cudev { void loadHueCSC(float hueCSC[9]) { - cvCudaSafeCall( cudaMemcpyToSymbol(constHueColorSpaceMat, hueCSC, 9 * sizeof(float)) ); + cudaSafeCall( 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 cudev { NV12ToARGB<<>>(decodedFrame.data, decodedFrame.step, interopFrame.data, interopFrame.step, interopFrame.cols, interopFrame.rows); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } }}} diff --git a/modules/gpu/src/cuda/bf_knnmatch.cu b/modules/gpu/src/cuda/bf_knnmatch.cu index 5fe40e5bd8..d5d17bb8a0 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 cudev 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -478,10 +478,10 @@ namespace cv { namespace gpu { namespace cudev 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -594,10 +594,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolled<<>>(query, train, mask, trainIdx.data, distance.data); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -653,10 +653,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolled<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -768,10 +768,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); match<<>>(query, train, mask, trainIdx.data, distance.data); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -827,10 +827,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); match<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -959,10 +959,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); calcDistanceUnrolled<<>>(query, train, mask, allDist); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -1022,10 +1022,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); calcDistance<<>>(query, train, mask, allDist); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -1115,11 +1115,11 @@ namespace cv { namespace gpu { namespace cudev for (int i = 0; i < k; ++i) { findBestMatch<<>>(allDist, i, trainIdx, distance); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 b6820aa531..338fefcb69 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 cudev 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -236,10 +236,10 @@ namespace cv { namespace gpu { namespace cudev 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -335,10 +335,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolled<<>>(query, train, mask, trainIdx.data, distance.data); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -392,10 +392,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); matchUnrolled<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -490,10 +490,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); match<<>>(query, train, mask, trainIdx.data, distance.data); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -546,10 +546,10 @@ namespace cv { namespace gpu { namespace cudev const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); match<<>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// diff --git a/modules/gpu/src/cuda/bf_radius_match.cu b/modules/gpu/src/cuda/bf_radius_match.cu index cb44768efb..3c714d63f9 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 cudev matchUnrolled<<>>(query, 0, train, maxDistance, mask, trainIdx, PtrStepi(), distance, nMatches.data, trainIdx.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -153,11 +153,11 @@ namespace cv { namespace gpu { namespace cudev matchUnrolled<<>>(query, i, train, maxDistance, WithOutMask(), trainIdx, imgIdx, distance, nMatches.data, trainIdx.cols); } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// @@ -230,10 +230,10 @@ namespace cv { namespace gpu { namespace cudev match<<>>(query, 0, train, maxDistance, mask, trainIdx, PtrStepi(), distance, nMatches.data, trainIdx.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -261,11 +261,11 @@ namespace cv { namespace gpu { namespace cudev match<<>>(query, i, train, maxDistance, WithOutMask(), trainIdx, imgIdx, distance, nMatches.data, trainIdx.cols); } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////////////////////// diff --git a/modules/gpu/src/cuda/bgfg_gmg.cu b/modules/gpu/src/cuda/bgfg_gmg.cu index 0e4421374f..8ae9b037b2 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 cudev { void loadConstants(int width, int height, float minVal, float maxVal, int quantizationLevels, float backgroundPrior, float decisionThreshold, int maxFeatures, int 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)) ); + 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)) ); } __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 cudev { const dim3 block(32, 8); const dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(update, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(update, cudaFuncCachePreferL1) ); update<<>>((PtrStepSz) frame, fgmask, colors, weights, nfeatures, frameNum, learningRate, updateBackgroundModel); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 1078343469..6508262d28 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 cudev dim3 block(32, 8); dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(mog_withoutLearning, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(mog_withoutLearning, cudaFuncCachePreferL1) ); mog_withoutLearning<<>>((PtrStepSz) frame, fgmask, weight, (PtrStepSz) mean, (PtrStepSz) var, nmixtures, varThreshold, backgroundRatio); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////// @@ -333,16 +333,16 @@ namespace cv { namespace gpu { namespace cudev dim3 block(32, 8); dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(mog_withLearning, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(mog_withLearning, cudaFuncCachePreferL1) ); mog_withLearning<<>>((PtrStepSz) frame, fgmask, weight, sortKey, (PtrStepSz) mean, (PtrStepSz) var, nmixtures, varThreshold, backgroundRatio, learningRate, minVar); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////////////////////////// @@ -406,13 +406,13 @@ namespace cv { namespace gpu { namespace cudev dim3 block(32, 8); dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage, cudaFuncCachePreferL1) ); getBackgroundImage<<>>(weight, (PtrStepSz) mean, (PtrStepSz) dst, nmixtures, backgroundRatio); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 cudev varMin = ::fminf(varMin, varMax); varMax = ::fmaxf(varMin, varMax); - 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)) ); + 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)) ); } template @@ -665,7 +665,7 @@ namespace cv { namespace gpu { namespace cudev if (detectShadows) { - cvCudaSafeCall( cudaFuncSetCacheConfig(mog2, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(mog2, cudaFuncCachePreferL1) ); mog2<<>>((PtrStepSz) frame, fgmask, modesUsed, weight, variance, (PtrStepSz) mean, @@ -673,17 +673,17 @@ namespace cv { namespace gpu { namespace cudev } else { - cvCudaSafeCall( cudaFuncSetCacheConfig(mog2, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(mog2, cudaFuncCachePreferL1) ); mog2<<>>((PtrStepSz) frame, fgmask, modesUsed, weight, variance, (PtrStepSz) mean, alphaT, alpha1, prune); } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 cudev dim3 block(32, 8); dim3 grid(divUp(modesUsed.cols, block.x), divUp(modesUsed.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage2, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage2, cudaFuncCachePreferL1) ); getBackgroundImage2<<>>(modesUsed, weight, (PtrStepSz) mean, (PtrStepSz) dst); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 8d87116180..4449274548 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 cudev float sigma_spatial2_inv_half = -0.5f/(sigma_spatial * sigma_spatial); float sigma_color2_inv_half = -0.5f/(sigma_color * sigma_color); - cvCudaSafeCall( cudaFuncSetCacheConfig (bilateral_kernel >, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig (bilateral_kernel >, cudaFuncCachePreferL1) ); bilateral_kernel<<>>((PtrStepSz)src, (PtrStepSz)dst, b, kernel_size, sigma_spatial2_inv_half, sigma_color2_inv_half); - cvCudaSafeCall ( cudaGetLastError () ); + cudaSafeCall ( cudaGetLastError () ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template diff --git a/modules/gpu/src/cuda/blend.cu b/modules/gpu/src/cuda/blend.cu index d8054d6b5f..be8c0b2f35 100644 --- a/modules/gpu/src/cuda/blend.cu +++ b/modules/gpu/src/cuda/blend.cu @@ -73,10 +73,10 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(cols * cn, threads.x), divUp(rows, threads.y)); blendLinearKernel<<>>(rows, cols * cn, cn, img1, img2, weights1, weights2, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } template void blendLinearCaller(int, int, int, PtrStep, PtrStep, PtrStepf, PtrStepf, PtrStep, cudaStream_t stream); @@ -109,10 +109,10 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); blendLinearKernel8UC4<<>>(rows, cols, img1, img2, weights1, weights2, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } } // namespace blend }}} // namespace cv { namespace gpu { namespace cudev diff --git a/modules/gpu/src/cuda/calib3d.cu b/modules/gpu/src/cuda/calib3d.cu index c85a2c7fe1..6085e716de 100644 --- a/modules/gpu/src/cuda/calib3d.cu +++ b/modules/gpu/src/cuda/calib3d.cu @@ -75,10 +75,10 @@ namespace cv { namespace gpu { namespace cudev const float* transl, PtrStepSz dst, cudaStream_t stream) { - 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)); + 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)); cv::gpu::cudev::transform(src, dst, TransformOp(), WithOutMask(), stream); } } // namespace transform_points @@ -114,12 +114,12 @@ namespace cv { namespace gpu { namespace cudev const float* transl, const float* proj, PtrStepSz dst, cudaStream_t stream) { - 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)); + 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)); cv::gpu::cudev::transform(src, dst, ProjectOp(), WithOutMask(), stream); } } // namespace project_points @@ -174,17 +174,17 @@ namespace cv { namespace gpu { namespace cudev const float3* transl_vectors, const float3* object, const float2* image, const float dist_threshold, int* hypothesis_scores) { - cvCudaSafeCall(cudaMemcpyToSymbol(crot_matrices, rot_matrices, num_hypotheses * 3 * sizeof(float3))); - cvCudaSafeCall(cudaMemcpyToSymbol(ctransl_vectors, transl_vectors, num_hypotheses * sizeof(float3))); + cudaSafeCall(cudaMemcpyToSymbol(crot_matrices, rot_matrices, num_hypotheses * 3 * sizeof(float3))); + cudaSafeCall(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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } // namespace solvepnp_ransac }}} // namespace cv { namespace gpu { namespace cudev diff --git a/modules/gpu/src/cuda/canny.cu b/modules/gpu/src/cuda/canny.cu index cbd792fd0b..042e9afcc6 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); } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall(cudaThreadSynchronize()); + cudaSafeCall(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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } @@ -324,17 +324,17 @@ namespace canny void edgesHysteresisLocal(PtrStepSzi map, ushort2* st1) { void* counter_ptr; - cvCudaSafeCall( cudaGetSymbolAddress(&counter_ptr, counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, counter) ); - cvCudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) ); + cudaSafeCall( 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } @@ -435,24 +435,24 @@ namespace canny void edgesHysteresisGlobal(PtrStepSzi map, ushort2* st1, ushort2* st2) { void* counter_ptr; - cvCudaSafeCall( cudaGetSymbolAddress(&counter_ptr, canny::counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, canny::counter) ); int count; - cvCudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); while (count > 0) { - cvCudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) ); + cudaSafeCall( 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); - cvCudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( 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 23dcfe9ba2..9552f1b06f 100644 --- a/modules/gpu/src/cuda/ccomponetns.cu +++ b/modules/gpu/src/cuda/ccomponetns.cu @@ -215,9 +215,9 @@ namespace cv { namespace gpu { namespace cudev Int_t inInt(lo, hi); computeConnectivity<<>>(static_cast >(image), edges, inInt); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 cudev dim3 grid(divUp(edges.cols, TILE_COLS), divUp(edges.rows, TILE_ROWS)); lableTiles<<>>(edges, comps); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); int tileSizeX = TILE_COLS, tileSizeY = TILE_ROWS; while (grid.x > 1 || grid.y > 1) @@ -517,16 +517,16 @@ namespace cv { namespace gpu { namespace cudev tileSizeY <<= 1; grid = mergeGrid; - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } grid.x = divUp(edges.cols, block.x); grid.y = divUp(edges.rows, block.y); flatten<<>>(edges, comps); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } } } } diff --git a/modules/gpu/src/cuda/clahe.cu b/modules/gpu/src/cuda/clahe.cu index 16afd1d5a7..7c6645749b 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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)); - cvCudaSafeCall( cudaFuncSetCacheConfig(tranformKernel, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(tranformKernel, cudaFuncCachePreferL1) ); tranformKernel<<>>(src, dst, lut, tileSize, tilesX, tilesY); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } diff --git a/modules/gpu/src/cuda/column_filter.h b/modules/gpu/src/cuda/column_filter.h index ecb6da16aa..39b6d47622 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } @@ -363,9 +363,9 @@ namespace filter }; if (stream == 0) - cvCudaSafeCall( cudaMemcpyToSymbol(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); + cudaSafeCall( cudaMemcpyToSymbol(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); else - cvCudaSafeCall( cudaMemcpyToSymbolAsync(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); + cudaSafeCall( 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 a9ad4f9888..ed90e9e80d 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 cudev BorderReader< PtrStep, B > brdSrc(src, brd); copyMakeBorder<<>>(brdSrc, dst, top, left); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; diff --git a/modules/gpu/src/cuda/debayer.cu b/modules/gpu/src/cuda/debayer.cu index 6c258a83aa..46a1c14ef4 100644 --- a/modules/gpu/src/cuda/debayer.cu +++ b/modules/gpu/src/cuda/debayer.cu @@ -347,13 +347,13 @@ namespace cv { namespace gpu { namespace cudev const dim3 block(32, 8); const dim3 grid(divUp(src.cols, 4 * block.x), divUp(src.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u, cudaFuncCachePreferL1) ); Bayer2BGR_8u<<>>(src, (PtrStepSz)dst, blue_last, start_with_green); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -364,13 +364,13 @@ namespace cv { namespace gpu { namespace cudev const dim3 block(32, 8); const dim3 grid(divUp(src.cols, 2 * block.x), divUp(src.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u, cudaFuncCachePreferL1) ); Bayer2BGR_16u<<>>(src, (PtrStepSz)dst, blue_last, start_with_green); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 cudev bindTexture(&sourceTex, src); MHCdemosaic<<>>((PtrStepSz)dst, sourceOffset, firstRed); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( 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 58dc4e80b8..cfea880ecd 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 cudev void disp_load_constants(float* table_color, PtrStepSzf table_space, int ndisp, int radius, short edge_disc, short max_disc) { - cvCudaSafeCall( cudaMemcpyToSymbol(ctable_color, &table_color, sizeof(table_color)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(ctable_space, &table_space.data, sizeof(table_space.data)) ); + cudaSafeCall( cudaMemcpyToSymbol(ctable_color, &table_color, sizeof(table_color)) ); + cudaSafeCall( cudaMemcpyToSymbol(ctable_space, &table_space.data, sizeof(table_space.data)) ); size_t table_space_step = table_space.step / sizeof(float); - cvCudaSafeCall( cudaMemcpyToSymbol(ctable_space_step, &table_space_step, sizeof(size_t)) ); + cudaSafeCall( cudaMemcpyToSymbol(ctable_space_step, &table_space_step, sizeof(size_t)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(cradius, &radius, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(cradius, &radius, sizeof(int)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(cedge_disc, &edge_disc, sizeof(short)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(cmax_disc, &max_disc, sizeof(short)) ); + cudaSafeCall( cudaMemcpyToSymbol(cedge_disc, &edge_disc, sizeof(short)) ); + cudaSafeCall( cudaMemcpyToSymbol(cmax_disc, &max_disc, sizeof(short)) ); } template @@ -191,20 +191,20 @@ namespace cv { namespace gpu { namespace cudev 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); disp_bilateral_filter<1><<>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( 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); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); disp_bilateral_filter<3><<>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } break; default: @@ -212,7 +212,7 @@ namespace cv { namespace gpu { namespace cudev } if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); diff --git a/modules/gpu/src/cuda/fast.cu b/modules/gpu/src/cuda/fast.cu index d0da17fbaf..39d66d0eda 100644 --- a/modules/gpu/src/cuda/fast.cu +++ b/modules/gpu/src/cuda/fast.cu @@ -282,7 +282,7 @@ namespace cv { namespace gpu { namespace cudev int calcKeypoints_gpu(PtrStepSzb img, PtrStepSzb mask, short2* kpLoc, int maxKeypoints, PtrStepSzi score, int threshold) { void* counter_ptr; - cvCudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) ); dim3 block(32, 8); @@ -290,7 +290,7 @@ namespace cv { namespace gpu { namespace cudev grid.x = divUp(img.cols - 6, block.x); grid.y = divUp(img.rows - 6, block.y); - cvCudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(unsigned int)) ); + cudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(unsigned int)) ); if (score.data) { @@ -307,12 +307,12 @@ namespace cv { namespace gpu { namespace cudev calcKeypoints<<>>(img, WithOutMask(), kpLoc, maxKeypoints, score, threshold); } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); unsigned int count; - cvCudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); return count; } @@ -359,22 +359,22 @@ namespace cv { namespace gpu { namespace cudev int nonmaxSupression_gpu(const short2* kpLoc, int count, PtrStepSzi score, short2* loc, float* response) { void* counter_ptr; - cvCudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) ); dim3 block(256); dim3 grid; grid.x = divUp(count, block.x); - cvCudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(unsigned int)) ); + cudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(unsigned int)) ); nonmaxSupression<<>>(kpLoc, count, score, loc, response); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); unsigned int new_count; - cvCudaSafeCall( cudaMemcpy(&new_count, counter_ptr, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&new_count, counter_ptr, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); return new_count; } diff --git a/modules/gpu/src/cuda/fgd_bgfg.cu b/modules/gpu/src/cuda/fgd_bgfg.cu index 30612b476f..d8c1df1900 100644 --- a/modules/gpu/src/cuda/fgd_bgfg.cu +++ b/modules/gpu/src/cuda/fgd_bgfg.cu @@ -205,13 +205,13 @@ namespace bgfg calcPartialHistogram<<>>( (PtrStepSz)prevFrame, (PtrStepSz)curFrame, partialBuf0, partialBuf1, partialBuf2); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); mergeHistogram<<>>(partialBuf0, partialBuf1, partialBuf2, hist0, hist1, hist2); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void calcDiffHistogram_gpu(PtrStepSzb prevFrame, PtrStepSzb curFrame, unsigned int* hist0, unsigned int* hist1, unsigned int* hist2, unsigned int* partialBuf0, unsigned int* partialBuf1, unsigned int* partialBuf2, bool cc20, cudaStream_t stream); @@ -251,10 +251,10 @@ namespace bgfg dim3 grid(divUp(prevFrame.cols, block.x), divUp(prevFrame.rows, block.y)); calcDiffThreshMask<<>>((PtrStepSz)prevFrame, (PtrStepSz)curFrame, bestThres, changeMask); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void calcDiffThreshMask_gpu(PtrStepSzb prevFrame, PtrStepSzb curFrame, uchar3 bestThres, PtrStepSzb changeMask, cudaStream_t stream); @@ -269,7 +269,7 @@ namespace bgfg void setBGPixelStat(const BGPixelStat& stat) { - cvCudaSafeCall( cudaMemcpyToSymbol(c_stat, &stat, sizeof(BGPixelStat)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_stat, &stat, sizeof(BGPixelStat)) ); } template struct Output; @@ -374,15 +374,15 @@ namespace bgfg dim3 block(32, 8); dim3 grid(divUp(prevFrame.cols, block.x), divUp(prevFrame.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(bgfgClassification, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(bgfgClassification, cudaFuncCachePreferL1) ); bgfgClassification<<>>((PtrStepSz)prevFrame, (PtrStepSz)curFrame, Ftd, Fbd, foreground, deltaC, deltaCC, alpha2, N1c, N1cc); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void bgfgClassification_gpu(PtrStepSzb prevFrame, PtrStepSzb curFrame, PtrStepSzb Ftd, PtrStepSzb Fbd, PtrStepSzb foreground, int deltaC, int deltaCC, float alpha2, int N1c, int N1cc, cudaStream_t stream); @@ -765,17 +765,17 @@ namespace bgfg dim3 block(32, 8); dim3 grid(divUp(prevFrame.cols, block.x), divUp(prevFrame.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(updateBackgroundModel, PtrStep, PtrStepb, PtrStepb>, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(updateBackgroundModel, PtrStep, PtrStepb, PtrStepb>, cudaFuncCachePreferL1) ); updateBackgroundModel, PtrStep, PtrStepb, PtrStepb><<>>( prevFrame.cols, prevFrame.rows, prevFrame, curFrame, Ftd, Fbd, foreground, background, deltaC, deltaCC, alpha1, alpha2, alpha3, N1c, N1cc, N2c, N2cc, T); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; diff --git a/modules/gpu/src/cuda/gftt.cu b/modules/gpu/src/cuda/gftt.cu index 2ab579cc90..b4af9e5dbc 100644 --- a/modules/gpu/src/cuda/gftt.cu +++ b/modules/gpu/src/cuda/gftt.cu @@ -94,9 +94,9 @@ namespace cv { namespace gpu { namespace cudev int findCorners_gpu(PtrStepSzf eig, float threshold, PtrStepSzb mask, float2* corners, int max_count) { void* counter_ptr; - cvCudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) ); - cvCudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) ); bindTexture(&eigTex, eig); @@ -108,12 +108,12 @@ namespace cv { namespace gpu { namespace cudev else findCorners<<>>(threshold, WithOutMask(), corners, max_count, eig.rows, eig.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int count; - cvCudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) ); return std::min(count, max_count); } diff --git a/modules/gpu/src/cuda/global_motion.cu b/modules/gpu/src/cuda/global_motion.cu index 07952ca3c2..5685c6750c 100644 --- a/modules/gpu/src/cuda/global_motion.cu +++ b/modules/gpu/src/cuda/global_motion.cu @@ -98,8 +98,8 @@ void calcWobbleSuppressionMaps( int left, int idx, int right, int width, int height, const float *ml, const float *mr, PtrStepSzf mapx, PtrStepSzf mapy) { - cvCudaSafeCall(cudaMemcpyToSymbol(cml, ml, 9*sizeof(float))); - cvCudaSafeCall(cudaMemcpyToSymbol(cmr, mr, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(cml, ml, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(cmr, mr, 9*sizeof(float))); dim3 threads(32, 8); dim3 grid(divUp(width, threads.x), divUp(height, threads.y)); @@ -107,8 +107,8 @@ void calcWobbleSuppressionMaps( calcWobbleSuppressionMapsKernel<<>>( left, idx, right, width, height, mapx, mapy); - cvCudaSafeCall(cudaGetLastError()); - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaDeviceSynchronize()); } }}}} diff --git a/modules/gpu/src/cuda/hist.cu b/modules/gpu/src/cuda/hist.cu index 09f3e51736..474c27cf76 100644 --- a/modules/gpu/src/cuda/hist.cu +++ b/modules/gpu/src/cuda/hist.cu @@ -100,10 +100,10 @@ namespace hist const dim3 grid(divUp(src.rows, block.y)); histogram256Kernel<<>>(src.data, src.cols, src.rows, src.step, hist); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } @@ -140,9 +140,9 @@ namespace hist void equalizeHist(PtrStepSzb src, PtrStepSzb dst, const int* lut, cudaStream_t stream) { if (stream == 0) - cvCudaSafeCall( cudaMemcpyToSymbol(c_lut, lut, 256 * sizeof(int), 0, cudaMemcpyDeviceToDevice) ); + cudaSafeCall( cudaMemcpyToSymbol(c_lut, lut, 256 * sizeof(int), 0, cudaMemcpyDeviceToDevice) ); else - cvCudaSafeCall( cudaMemcpyToSymbolAsync(c_lut, lut, 256 * sizeof(int), 0, cudaMemcpyDeviceToDevice, stream) ); + cudaSafeCall( cudaMemcpyToSymbolAsync(c_lut, lut, 256 * sizeof(int), 0, cudaMemcpyDeviceToDevice, stream) ); const float scale = 255.0f / (src.cols * src.rows); diff --git a/modules/gpu/src/cuda/hog.cu b/modules/gpu/src/cuda/hog.cu index 82682bb10c..48d656a744 100644 --- a/modules/gpu/src/cuda/hog.cu +++ b/modules/gpu/src/cuda/hog.cu @@ -90,23 +90,23 @@ namespace cv { namespace gpu { namespace cudev void set_up_constants(int nbins, int block_stride_x, int block_stride_y, int nblocks_win_x, int nblocks_win_y) { - cvCudaSafeCall( cudaMemcpyToSymbol(cnbins, &nbins, sizeof(nbins)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(cblock_stride_x, &block_stride_x, sizeof(block_stride_x)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(cblock_stride_y, &block_stride_y, sizeof(block_stride_y)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(cnblocks_win_x, &nblocks_win_x, sizeof(nblocks_win_x)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(cnblocks_win_y, &nblocks_win_y, sizeof(nblocks_win_y)) ); + cudaSafeCall( cudaMemcpyToSymbol(cnbins, &nbins, sizeof(nbins)) ); + cudaSafeCall( cudaMemcpyToSymbol(cblock_stride_x, &block_stride_x, sizeof(block_stride_x)) ); + cudaSafeCall( cudaMemcpyToSymbol(cblock_stride_y, &block_stride_y, sizeof(block_stride_y)) ); + cudaSafeCall( cudaMemcpyToSymbol(cnblocks_win_x, &nblocks_win_x, sizeof(nblocks_win_x)) ); + cudaSafeCall( cudaMemcpyToSymbol(cnblocks_win_y, &nblocks_win_y, sizeof(nblocks_win_y)) ); int block_hist_size = nbins * CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y; - cvCudaSafeCall( cudaMemcpyToSymbol(cblock_hist_size, &block_hist_size, sizeof(block_hist_size)) ); + cudaSafeCall( cudaMemcpyToSymbol(cblock_hist_size, &block_hist_size, sizeof(block_hist_size)) ); int block_hist_size_2up = power_2up(block_hist_size); - cvCudaSafeCall( cudaMemcpyToSymbol(cblock_hist_size_2up, &block_hist_size_2up, sizeof(block_hist_size_2up)) ); + cudaSafeCall( cudaMemcpyToSymbol(cblock_hist_size_2up, &block_hist_size_2up, sizeof(block_hist_size_2up)) ); int descr_width = nblocks_win_x * block_hist_size; - cvCudaSafeCall( cudaMemcpyToSymbol(cdescr_width, &descr_width, sizeof(descr_width)) ); + cudaSafeCall( cudaMemcpyToSymbol(cdescr_width, &descr_width, sizeof(descr_width)) ); int descr_size = descr_width * nblocks_win_y; - cvCudaSafeCall( cudaMemcpyToSymbol(cdescr_size, &descr_size, sizeof(descr_size)) ); + cudaSafeCall( cudaMemcpyToSymbol(cdescr_size, &descr_size, sizeof(descr_size)) ); } @@ -206,7 +206,7 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(img_block_width, nblocks), img_block_height); dim3 threads(32, 2, nblocks); - cvCudaSafeCall(cudaFuncSetCacheConfig(compute_hists_kernel_many_blocks, + cudaSafeCall(cudaFuncSetCacheConfig(compute_hists_kernel_many_blocks, cudaFuncCachePreferL1)); // Precompute gaussian spatial window parameter @@ -217,9 +217,9 @@ namespace cv { namespace gpu { namespace cudev int smem = hists_size + final_hists_size; compute_hists_kernel_many_blocks<<>>( img_block_width, grad, qangle, scale, block_hists); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -318,9 +318,9 @@ namespace cv { namespace gpu { namespace cudev else CV_Error(cv::Error::StsBadArg, "normalize_hists: histogram's size is too big, try to decrease number of bins"); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -378,7 +378,7 @@ namespace cv { namespace gpu { namespace cudev dim3 threads(nthreads, 1, nblocks); dim3 grid(divUp(img_win_width, nblocks), img_win_height); - cvCudaSafeCall(cudaFuncSetCacheConfig(compute_confidence_hists_kernel_many_blocks, + cudaSafeCall(cudaFuncSetCacheConfig(compute_confidence_hists_kernel_many_blocks, cudaFuncCachePreferL1)); int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) / @@ -386,7 +386,7 @@ namespace cv { namespace gpu { namespace cudev compute_confidence_hists_kernel_many_blocks<<>>( img_win_width, img_block_width, win_block_stride_x, win_block_stride_y, block_hists, coefs, free_coef, threshold, confidences); - cvCudaSafeCall(cudaThreadSynchronize()); + cudaSafeCall(cudaThreadSynchronize()); } @@ -440,15 +440,15 @@ namespace cv { namespace gpu { namespace cudev dim3 threads(nthreads, 1, nblocks); dim3 grid(divUp(img_win_width, nblocks), img_win_height); - cvCudaSafeCall(cudaFuncSetCacheConfig(classify_hists_kernel_many_blocks, cudaFuncCachePreferL1)); + cudaSafeCall(cudaFuncSetCacheConfig(classify_hists_kernel_many_blocks, cudaFuncCachePreferL1)); int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) / block_stride_x; classify_hists_kernel_many_blocks<<>>( img_win_width, img_block_width, win_block_stride_x, win_block_stride_y, block_hists, coefs, free_coef, threshold, labels); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //---------------------------------------------------------------------------- @@ -491,9 +491,9 @@ namespace cv { namespace gpu { namespace cudev int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) / block_stride_x; extract_descrs_by_rows_kernel<<>>( img_block_width, win_block_stride_x, win_block_stride_y, block_hists, descriptors); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -540,9 +540,9 @@ namespace cv { namespace gpu { namespace cudev int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) / block_stride_x; extract_descrs_by_cols_kernel<<>>( img_block_width, win_block_stride_x, win_block_stride_y, block_hists, descriptors); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //---------------------------------------------------------------------------- @@ -666,9 +666,9 @@ namespace cv { namespace gpu { namespace cudev else compute_gradients_8UC4_kernel<<>>(height, width, img, angle_scale, grad, qangle); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -739,9 +739,9 @@ namespace cv { namespace gpu { namespace cudev else compute_gradients_8UC1_kernel<<>>(height, width, img, angle_scale, grad, qangle); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -782,13 +782,13 @@ namespace cv { namespace gpu { namespace cudev int colOfs = 0; cudaChannelFormatDesc desc = cudaCreateChannelDesc(); - cvCudaSafeCall( cudaBindTexture2D(&texOfs, tex, src.data, desc, src.cols, src.rows, src.step) ); + cudaSafeCall( cudaBindTexture2D(&texOfs, tex, src.data, desc, src.cols, src.rows, src.step) ); if (texOfs != 0) { colOfs = static_cast( texOfs/sizeof(T) ); - cvCudaSafeCall( cudaUnbindTexture(tex) ); - cvCudaSafeCall( cudaBindTexture2D(&texOfs, tex, src.data, desc, src.cols, src.rows, src.step) ); + cudaSafeCall( cudaUnbindTexture(tex) ); + cudaSafeCall( cudaBindTexture2D(&texOfs, tex, src.data, desc, src.cols, src.rows, src.step) ); } dim3 threads(32, 8); @@ -798,11 +798,11 @@ namespace cv { namespace gpu { namespace cudev float sy = static_cast(src.rows) / dst.rows; resize_for_hog_kernel<<>>(sx, sy, (PtrStepSz)dst, colOfs); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); - cvCudaSafeCall( cudaUnbindTexture(tex) ); + cudaSafeCall( cudaUnbindTexture(tex) ); } void resize_8UC1(const PtrStepSzb& src, PtrStepSzb dst) { resize_for_hog (src, dst, resize8UC1_tex); } diff --git a/modules/gpu/src/cuda/hough.cu b/modules/gpu/src/cuda/hough.cu index d1cdf82e9e..5a4481b6e5 100644 --- a/modules/gpu/src/cuda/hough.cu +++ b/modules/gpu/src/cuda/hough.cu @@ -122,22 +122,22 @@ namespace cv { namespace gpu { namespace cudev const int PIXELS_PER_THREAD = 16; void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(32, 4); const dim3 grid(divUp(src.cols, block.x * PIXELS_PER_THREAD), divUp(src.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(buildPointList, cudaFuncCachePreferShared) ); + cudaSafeCall( cudaFuncSetCacheConfig(buildPointList, cudaFuncCachePreferShared) ); buildPointList<<>>(src, list); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); return totalCount; } @@ -225,9 +225,9 @@ namespace cv { namespace gpu { namespace cudev else linesAccumGlobal<<>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////// @@ -264,22 +264,22 @@ namespace cv { namespace gpu { namespace cudev int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort) { void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(32, 8); const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(linesGetResult, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(linesGetResult, cudaFuncCachePreferL1) ); linesGetResult<<>>(accum, out, votes, maxSize, rho, theta, threshold, accum.cols - 2); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); totalCount = ::min(totalCount, maxSize); @@ -462,9 +462,9 @@ namespace cv { namespace gpu { namespace cudev int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength) { void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(32, 8); const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); @@ -476,12 +476,12 @@ namespace cv { namespace gpu { namespace cudev rho, theta, lineGap, lineLength, mask.rows, mask.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); totalCount = ::min(totalCount, maxSize); @@ -548,12 +548,12 @@ namespace cv { namespace gpu { namespace cudev const dim3 block(256); const dim3 grid(divUp(count, block.x)); - cvCudaSafeCall( cudaFuncSetCacheConfig(circlesAccumCenters, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(circlesAccumCenters, cudaFuncCachePreferL1) ); circlesAccumCenters<<>>(list, count, dx, dy, accum, accum.cols - 2, accum.rows - 2, minRadius, maxRadius, idp); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////// @@ -586,22 +586,22 @@ namespace cv { namespace gpu { namespace cudev int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold) { void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(32, 8); const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(buildCentersList, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(buildCentersList, cudaFuncCachePreferL1) ); buildCentersList<<>>(accum, centers, threshold); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); return totalCount; } @@ -662,9 +662,9 @@ namespace cv { namespace gpu { namespace cudev float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20) { void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(has20 ? 1024 : 512); const dim3 grid(centersCount); @@ -673,12 +673,12 @@ namespace cv { namespace gpu { namespace cudev size_t smemSize = (histSize + 2) * sizeof(int); circlesAccumRadius<<>>(centers, list, count, circles, maxCircles, dp, minRadius, maxRadius, histSize, threshold); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); totalCount = ::min(totalCount, maxCircles); @@ -768,22 +768,22 @@ namespace cv { namespace gpu { namespace cudev const int PIXELS_PER_THREAD = 8; void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(32, 4); const dim3 grid(divUp(edges.cols, block.x * PIXELS_PER_THREAD), divUp(edges.rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(buildEdgePointList, cudaFuncCachePreferShared) ); + cudaSafeCall( cudaFuncSetCacheConfig(buildEdgePointList, cudaFuncCachePreferShared) ); buildEdgePointList<<>>(edges, (PtrStepSz) dx, (PtrStepSz) dy, coordList, thetaList); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); return totalCount; } @@ -824,9 +824,9 @@ namespace cv { namespace gpu { namespace cudev const float thetaScale = levels / (2.0f * CV_PI_F); buildRTable<<>>(coordList, thetaList, pointsCount, r_table, r_sizes, r_table.cols, templCenter, thetaScale); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////// @@ -877,9 +877,9 @@ namespace cv { namespace gpu { namespace cudev const float thetaScale = levels / (2.0f * CV_PI_F); GHT_Ballard_Pos_calcHist<<>>(coordList, thetaList, pointsCount, r_table, r_sizes, hist, idp, thetaScale); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void GHT_Ballard_Pos_findPosInHist(const PtrStepSzi hist, float4* out, int3* votes, const int maxSize, const float dp, const int threshold) @@ -911,22 +911,22 @@ namespace cv { namespace gpu { namespace cudev int GHT_Ballard_Pos_findPosInHist_gpu(PtrStepSzi hist, float4* out, int3* votes, int maxSize, float dp, int threshold) { void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(32, 8); const dim3 grid(divUp(hist.cols - 2, block.x), divUp(hist.rows - 2, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(GHT_Ballard_Pos_findPosInHist, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(GHT_Ballard_Pos_findPosInHist, cudaFuncCachePreferL1) ); GHT_Ballard_Pos_findPosInHist<<>>(hist, out, votes, maxSize, dp, threshold); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); totalCount = ::min(totalCount, maxSize); @@ -989,9 +989,9 @@ namespace cv { namespace gpu { namespace cudev hist, rows, cols, minScale, scaleStep, scaleRange, idp, thetaScale); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void GHT_Ballard_PosScale_findPosInHist(const PtrStepi hist, const int rows, const int cols, const int scaleRange, @@ -1037,22 +1037,22 @@ namespace cv { namespace gpu { namespace cudev float minScale, float scaleStep, float dp, int threshold) { void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(32, 8); const dim3 grid(divUp(cols, block.x), divUp(rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(GHT_Ballard_PosScale_findPosInHist, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(GHT_Ballard_PosScale_findPosInHist, cudaFuncCachePreferL1) ); GHT_Ballard_PosScale_findPosInHist<<>>(hist, rows, cols, scaleRange, out, votes, maxSize, minScale, scaleStep, dp, threshold); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); totalCount = ::min(totalCount, maxSize); @@ -1123,9 +1123,9 @@ namespace cv { namespace gpu { namespace cudev hist, rows, cols, minAngle, angleStep, angleRange, idp, thetaScale); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void GHT_Ballard_PosRotation_findPosInHist(const PtrStepi hist, const int rows, const int cols, const int angleRange, @@ -1171,22 +1171,22 @@ namespace cv { namespace gpu { namespace cudev float minAngle, float angleStep, float dp, int threshold) { void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); const dim3 block(32, 8); const dim3 grid(divUp(cols, block.x), divUp(rows, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(GHT_Ballard_PosRotation_findPosInHist, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(GHT_Ballard_PosRotation_findPosInHist, cudaFuncCachePreferL1) ); GHT_Ballard_PosRotation_findPosInHist<<>>(hist, rows, cols, angleRange, out, votes, maxSize, minAngle, angleStep, dp, threshold); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); totalCount = ::min(totalCount, maxSize); @@ -1242,7 +1242,7 @@ namespace cv { namespace gpu { namespace cudev tbl.r2_data = r2.data; tbl.r2_step = r2.step; - cvCudaSafeCall( cudaMemcpyToSymbol(c_templFeatures, &tbl, sizeof(FeatureTable)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_templFeatures, &tbl, sizeof(FeatureTable)) ); } void GHT_Guil_Full_setImageFeatures(PtrStepb p1_pos, PtrStepb p1_theta, PtrStepb p2_pos, PtrStepb d12, PtrStepb r1, PtrStepb r2) { @@ -1266,7 +1266,7 @@ namespace cv { namespace gpu { namespace cudev tbl.r2_data = r2.data; tbl.r2_step = r2.step; - cvCudaSafeCall( cudaMemcpyToSymbol(c_imageFeatures, &tbl, sizeof(FeatureTable)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_imageFeatures, &tbl, sizeof(FeatureTable)) ); } struct TemplFeatureTable @@ -1419,9 +1419,9 @@ namespace cv { namespace gpu { namespace cudev sizes, maxSize, xi * (CV_PI_F / 180.0f), angleEpsilon * (CV_PI_F / 180.0f), alphaScale, center, maxDist); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); thrust::device_ptr sizesPtr(sizes); thrust::transform(sizesPtr, sizesPtr + levels + 1, sizesPtr, cudev::bind2nd(cudev::minimum(), maxSize)); @@ -1501,9 +1501,9 @@ namespace cv { namespace gpu { namespace cudev GHT_Guil_Full_calcOHist<<>>(templSizes, imageSizes, OHist, minAngle, maxAngle, 1.0f / angleStep, angleRange); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void GHT_Guil_Full_calcSHist(const int* templSizes, const int* imageSizes, int* SHist, @@ -1566,9 +1566,9 @@ namespace cv { namespace gpu { namespace cudev GHT_Guil_Full_calcSHist<<>>(templSizes, imageSizes, SHist, angle, angleEpsilon, minScale, maxScale, iScaleStep, scaleRange); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void GHT_Guil_Full_calcPHist(const int* templSizes, const int* imageSizes, PtrStepSzi PHist, @@ -1636,14 +1636,14 @@ namespace cv { namespace gpu { namespace cudev const float sinVal = ::sinf(angle); const float cosVal = ::cosf(angle); - cvCudaSafeCall( cudaFuncSetCacheConfig(GHT_Guil_Full_calcPHist, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(GHT_Guil_Full_calcPHist, cudaFuncCachePreferL1) ); GHT_Guil_Full_calcPHist<<>>(templSizes, imageSizes, PHist, angle, sinVal, cosVal, angleEpsilon, scale, 1.0f / dp); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void GHT_Guil_Full_findPosInHist(const PtrStepSzi hist, float4* out, int3* votes, const int maxSize, @@ -1679,24 +1679,24 @@ namespace cv { namespace gpu { namespace cudev float dp, int threshold) { void* counterPtr; - cvCudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - cvCudaSafeCall( cudaMemcpy(counterPtr, &curSize, sizeof(int), cudaMemcpyHostToDevice) ); + cudaSafeCall( cudaMemcpy(counterPtr, &curSize, sizeof(int), cudaMemcpyHostToDevice) ); const dim3 block(32, 8); const dim3 grid(divUp(hist.cols - 2, block.x), divUp(hist.rows - 2, block.y)); - cvCudaSafeCall( cudaFuncSetCacheConfig(GHT_Guil_Full_findPosInHist, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(GHT_Guil_Full_findPosInHist, cudaFuncCachePreferL1) ); GHT_Guil_Full_findPosInHist<<>>(hist, out, votes, maxSize, angle, angleVotes, scale, scaleVotes, dp, threshold); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); int totalCount; - cvCudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); totalCount = ::min(totalCount, maxSize); diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index b145235f7e..fc27ec19a9 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -154,13 +154,13 @@ namespace cv { namespace gpu { namespace cudev grid.y = divUp(src.rows, threads.y); cudaChannelFormatDesc desc = cudaCreateChannelDesc(); - cvCudaSafeCall( cudaBindTexture2D( 0, tex_meanshift, src.data, desc, src.cols, src.rows, src.step ) ); + cudaSafeCall( cudaBindTexture2D( 0, tex_meanshift, src.data, desc, src.cols, src.rows, src.step ) ); meanshift_kernel<<< grid, threads, 0, stream >>>( dst.data, dst.step, dst.cols, dst.rows, sp, sr, maxIter, eps ); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); //cudaSafeCall( cudaUnbindTexture( tex_meanshift ) ); } @@ -173,13 +173,13 @@ namespace cv { namespace gpu { namespace cudev grid.y = divUp(src.rows, threads.y); cudaChannelFormatDesc desc = cudaCreateChannelDesc(); - cvCudaSafeCall( cudaBindTexture2D( 0, tex_meanshift, src.data, desc, src.cols, src.rows, src.step ) ); + cudaSafeCall( cudaBindTexture2D( 0, tex_meanshift, src.data, desc, src.cols, src.rows, src.step ) ); meanshiftproc_kernel<<< grid, threads, 0, stream >>>( dstr.data, dstr.step, dstsp.data, dstsp.step, dstr.cols, dstr.rows, sp, sr, maxIter, eps ); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); //cudaSafeCall( cudaUnbindTexture( tex_meanshift ) ); } @@ -295,10 +295,10 @@ namespace cv { namespace gpu { namespace cudev grid.y = divUp(src.rows, threads.y); drawColorDisp<<>>(src.data, src.step, dst.data, dst.step, src.cols, src.rows, ndisp); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } void drawColorDisp_gpu(const PtrStepSz& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream) @@ -309,10 +309,10 @@ namespace cv { namespace gpu { namespace cudev grid.y = divUp(src.rows, threads.y); drawColorDisp<<>>(src.data, src.step / sizeof(short), dst.data, dst.step, src.cols, src.rows, ndisp); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////// reprojectImageTo3D /////////////////////////////////////////////// @@ -351,13 +351,13 @@ namespace cv { namespace gpu { namespace cudev dim3 block(32, 8); dim3 grid(divUp(disp.cols, block.x), divUp(disp.rows, block.y)); - cvCudaSafeCall( cudaMemcpyToSymbol(cq, q, 16 * sizeof(float)) ); + cudaSafeCall( cudaMemcpyToSymbol(cq, q, 16 * sizeof(float)) ); reprojectImageTo3D<<>>((PtrStepSz)disp, (PtrStepSz)xyz); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); @@ -464,10 +464,10 @@ namespace cv { namespace gpu { namespace cudev break; } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } /////////////////////////////////////////// Corner Min Eigen Val ///////////////////////////////////////////////// @@ -576,10 +576,10 @@ namespace cv { namespace gpu { namespace cudev break; } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } ////////////////////////////// Column Sum ////////////////////////////////////// @@ -611,9 +611,9 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(src.cols, threads.x)); column_sumKernel_32F<<>>(src.cols, src.rows, src, dst); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -638,10 +638,10 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(c.cols, threads.x), divUp(c.rows, threads.y)); mulSpectrumsKernel<<>>(a, b, c); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -666,10 +666,10 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(c.cols, threads.x), divUp(c.rows, threads.y)); mulSpectrumsKernel_CONJ<<>>(a, b, c); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -695,10 +695,10 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(c.cols, threads.x), divUp(c.rows, threads.y)); mulAndScaleSpectrumsKernel<<>>(a, b, scale, c); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -724,10 +724,10 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(c.cols, threads.x), divUp(c.rows, threads.y)); mulAndScaleSpectrumsKernel_CONJ<<>>(a, b, scale, c); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } ////////////////////////////////////////////////////////////////////////// @@ -837,10 +837,10 @@ namespace cv { namespace gpu { namespace cudev const float k_rinv[9], const float r_kinv[9], const float t[3], float scale, cudaStream_t stream) { - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ct, t, 3*sizeof(float))); - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ct, t, 3*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); int cols = map_x.cols; int rows = map_x.rows; @@ -849,9 +849,9 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } @@ -859,9 +859,9 @@ namespace cv { namespace gpu { namespace cudev const float k_rinv[9], const float r_kinv[9], float scale, cudaStream_t stream) { - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); int cols = map_x.cols; int rows = map_x.rows; @@ -870,9 +870,9 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } @@ -880,9 +880,9 @@ namespace cv { namespace gpu { namespace cudev const float k_rinv[9], const float r_kinv[9], float scale, cudaStream_t stream) { - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); - cvCudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); int cols = map_x.cols; int rows = map_x.rows; @@ -891,9 +891,9 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } ////////////////////////////////////////////////////////////////////////// @@ -955,9 +955,9 @@ namespace cv { namespace gpu { namespace cudev Brd brd(dst.rows, dst.cols, VecTraits::make(borderValue)); \ BorderReader< tex_filter2D_ ## type ##_reader, Brd > brdSrc(texSrc, brd); \ filter2D<<>>(brdSrc, dst, kWidth, kHeight, anchorX, anchorY); \ - cvCudaSafeCall( cudaGetLastError() ); \ + cudaSafeCall( cudaGetLastError() ); \ if (stream == 0) \ - cvCudaSafeCall( cudaDeviceSynchronize() ); \ + cudaSafeCall( cudaDeviceSynchronize() ); \ } \ }; @@ -988,9 +988,9 @@ namespace cv { namespace gpu { namespace cudev }; if (stream == 0) - cvCudaSafeCall( cudaMemcpyToSymbol(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); + cudaSafeCall( cudaMemcpyToSymbol(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); else - cvCudaSafeCall( cudaMemcpyToSymbolAsync(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); + cudaSafeCall( cudaMemcpyToSymbolAsync(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); funcs[borderMode](static_cast< PtrStepSz >(srcWhole), ofsX, ofsY, static_cast< PtrStepSz >(dst), kWidth, kHeight, anchorX, anchorY, borderValue, stream); } diff --git a/modules/gpu/src/cuda/integral_image.cu b/modules/gpu/src/cuda/integral_image.cu index 231ee659fa..d8276b2c9a 100644 --- a/modules/gpu/src/cuda/integral_image.cu +++ b/modules/gpu/src/cuda/integral_image.cu @@ -367,10 +367,10 @@ namespace cv { namespace gpu { namespace cudev // launch 1 block / row const int grid = img.rows; - cvCudaSafeCall( cudaFuncSetCacheConfig(shfl_integral_horizontal, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(shfl_integral_horizontal, cudaFuncCachePreferL1) ); shfl_integral_horizontal<<>>((const PtrStepSz) img, (PtrStepSz) integral); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } { @@ -378,11 +378,11 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(integral.cols, block.x), 1); shfl_integral_vertical<<>>(integral); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void shfl_integral_vertical(PtrStepSz buffer, PtrStepSz integral) @@ -452,10 +452,10 @@ namespace cv { namespace gpu { namespace cudev const int block = blockStep; const int grid = img.rows; - cvCudaSafeCall( cudaFuncSetCacheConfig(shfl_integral_horizontal, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(shfl_integral_horizontal, cudaFuncCachePreferL1) ); shfl_integral_horizontal<<>>((PtrStepSz) img, buffer); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } { @@ -463,7 +463,7 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(integral.cols, block.x), 1); shfl_integral_vertical<<>>((PtrStepSz)buffer, integral); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } } } diff --git a/modules/gpu/src/cuda/lbp.cu b/modules/gpu/src/cuda/lbp.cu index fbff332f1e..fb6267f2fe 100644 --- a/modules/gpu/src/cuda/lbp.cu +++ b/modules/gpu/src/cuda/lbp.cu @@ -189,7 +189,7 @@ namespace cv { namespace gpu { namespace cudev int block = ncandidates; int smem = block * ( sizeof(int) + sizeof(int4) ); disjoin<<<1, block, smem>>>(candidates, objects, ncandidates, groupThreshold, grouping_eps, nclasses); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } struct Cascade diff --git a/modules/gpu/src/cuda/match_template.cu b/modules/gpu/src/cuda/match_template.cu index 81387bd6a0..6670639290 100644 --- a/modules/gpu/src/cuda/match_template.cu +++ b/modules/gpu/src/cuda/match_template.cu @@ -114,10 +114,10 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(result.cols, threads.x), divUp(result.rows, threads.y)); matchTemplateNaiveKernel_CCORR<<>>(templ.cols, templ.rows, image, templ, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } void matchTemplateNaive_CCORR_32F(const PtrStepSzb image, const PtrStepSzb templ, PtrStepSzf result, int cn, cudaStream_t stream) @@ -184,10 +184,10 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(result.cols, threads.x), divUp(result.rows, threads.y)); matchTemplateNaiveKernel_SQDIFF<<>>(templ.cols, templ.rows, image, templ, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } void matchTemplateNaive_SQDIFF_32F(const PtrStepSzb image, const PtrStepSzb templ, PtrStepSzf result, int cn, cudaStream_t stream) @@ -240,10 +240,10 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(result.cols, threads.x), divUp(result.rows, threads.y)); matchTemplatePreparedKernel_SQDIFF_8U<<>>(w, h, image_sqsum, templ_sqsum, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } void matchTemplatePrepared_SQDIFF_8U(int w, int h, const PtrStepSz image_sqsum, unsigned long long templ_sqsum, PtrStepSzf result, int cn, @@ -312,10 +312,10 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(result.cols, threads.x), divUp(result.rows, threads.y)); matchTemplatePreparedKernel_SQDIFF_NORMED_8U<<>>(w, h, image_sqsum, templ_sqsum, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -355,10 +355,10 @@ namespace cv { namespace gpu { namespace cudev dim3 grid(divUp(result.cols, threads.x), divUp(result.rows, threads.y)); matchTemplatePreparedKernel_CCOFF_8U<<>>(w, h, (float)templ_sum / (w * h), image_sum, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -399,10 +399,10 @@ namespace cv { namespace gpu { namespace cudev matchTemplatePreparedKernel_CCOFF_8UC2<<>>( w, h, (float)templ_sum_r / (w * h), (float)templ_sum_g / (w * h), image_sum_r, image_sum_g, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -457,10 +457,10 @@ namespace cv { namespace gpu { namespace cudev (float)templ_sum_g / (w * h), (float)templ_sum_b / (w * h), image_sum_r, image_sum_g, image_sum_b, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -525,10 +525,10 @@ namespace cv { namespace gpu { namespace cudev (float)templ_sum_a / (w * h), image_sum_r, image_sum_g, image_sum_b, image_sum_a, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } ////////////////////////////////////////////////////////////////////// @@ -574,10 +574,10 @@ namespace cv { namespace gpu { namespace cudev matchTemplatePreparedKernel_CCOFF_NORMED_8U<<>>( w, h, weight, templ_sum_scale, templ_sqsum_scale, image_sum, image_sqsum, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -640,10 +640,10 @@ namespace cv { namespace gpu { namespace cudev image_sum_r, image_sqsum_r, image_sum_g, image_sqsum_g, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -720,10 +720,10 @@ namespace cv { namespace gpu { namespace cudev image_sum_g, image_sqsum_g, image_sum_b, image_sqsum_b, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } @@ -812,10 +812,10 @@ namespace cv { namespace gpu { namespace cudev image_sum_b, image_sqsum_b, image_sum_a, image_sqsum_a, result); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } ////////////////////////////////////////////////////////////////////// @@ -860,10 +860,10 @@ namespace cv { namespace gpu { namespace cudev break; } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } ////////////////////////////////////////////////////////////////////// @@ -904,10 +904,10 @@ namespace cv { namespace gpu { namespace cudev extractFirstChannel_32F<4><<>>(image, result); break; } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } //namespace match_template }}} // namespace cv { namespace gpu { namespace cudev diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu index e8caab66de..88626c5d2e 100644 --- a/modules/gpu/src/cuda/mathfunc.cu +++ b/modules/gpu/src/cuda/mathfunc.cu @@ -145,10 +145,10 @@ namespace cv { namespace gpu { namespace cudev cartToPolar<<>>( x.data, x.step/x.elemSize(), y.data, y.step/y.elemSize(), mag.data, mag.step/mag.elemSize(), angle.data, angle.step/angle.elemSize(), scale, x.cols, x.rows); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } void cartToPolar_gpu(PtrStepSzf x, PtrStepSzf y, PtrStepSzf mag, bool magSqr, PtrStepSzf angle, bool angleInDegrees, cudaStream_t stream) @@ -194,10 +194,10 @@ namespace cv { namespace gpu { namespace cudev polarToCart<<>>(mag.data, mag.step/mag.elemSize(), angle.data, angle.step/angle.elemSize(), scale, x.data, x.step/x.elemSize(), y.data, y.step/y.elemSize(), mag.cols, mag.rows); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } void polarToCart_gpu(PtrStepSzf mag, PtrStepSzf angle, PtrStepSzf x, PtrStepSzf y, bool angleInDegrees, cudaStream_t stream) diff --git a/modules/gpu/src/cuda/matrix_reductions.cu b/modules/gpu/src/cuda/matrix_reductions.cu index 56191da28e..8eac195131 100644 --- a/modules/gpu/src/cuda/matrix_reductions.cu +++ b/modules/gpu/src/cuda/matrix_reductions.cu @@ -432,12 +432,12 @@ namespace sum kernel<<>>(src, buf, SingleMask(mask), op, twidth, theight); else kernel<<>>(src, buf, WithOutMask(), op, twidth, theight); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); R result[4] = {0, 0, 0, 0}; - cvCudaSafeCall( cudaMemcpy(&result, buf, sizeof(result_type), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&result, buf, sizeof(result_type), cudaMemcpyDeviceToHost) ); out[0] = result[0]; out[1] = result[1]; @@ -761,13 +761,13 @@ namespace minMax else kernel<<>>((PtrStepSz) src, WithOutMask(), minval_buf, maxval_buf, twidth, theight); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); R minval_, maxval_; - cvCudaSafeCall( cudaMemcpy(&minval_, minval_buf, sizeof(R), cudaMemcpyDeviceToHost) ); - cvCudaSafeCall( cudaMemcpy(&maxval_, maxval_buf, sizeof(R), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&minval_, minval_buf, sizeof(R), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&maxval_, maxval_buf, sizeof(R), cudaMemcpyDeviceToHost) ); *minval = minval_; *maxval = maxval_; } @@ -934,22 +934,22 @@ namespace minMaxLoc else kernel_pass_1<<>>((PtrStepSz) src, WithOutMask(), minval_buf, maxval_buf, minloc_buf, maxloc_buf, twidth, theight); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); kernel_pass_2<<<1, threads_x * threads_y>>>(minval_buf, maxval_buf, minloc_buf, maxloc_buf, grid.x * grid.y); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); T minval_, maxval_; - cvCudaSafeCall( cudaMemcpy(&minval_, minval_buf, sizeof(T), cudaMemcpyDeviceToHost) ); - cvCudaSafeCall( cudaMemcpy(&maxval_, maxval_buf, sizeof(T), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&minval_, minval_buf, sizeof(T), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&maxval_, maxval_buf, sizeof(T), cudaMemcpyDeviceToHost) ); *minval = minval_; *maxval = maxval_; unsigned int minloc_, maxloc_; - cvCudaSafeCall( cudaMemcpy(&minloc_, minloc_buf, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); - cvCudaSafeCall( cudaMemcpy(&maxloc_, maxloc_buf, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&minloc_, minloc_buf, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); + cudaSafeCall( cudaMemcpy(&maxloc_, maxloc_buf, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); minloc[1] = minloc_ / src.cols; minloc[0] = minloc_ - minloc[1] * src.cols; maxloc[1] = maxloc_ / src.cols; maxloc[0] = maxloc_ - maxloc[1] * src.cols; } @@ -1065,15 +1065,15 @@ namespace countNonZero unsigned int* count_buf = buf.ptr(0); - cvCudaSafeCall( cudaMemset(count_buf, 0, sizeof(unsigned int)) ); + cudaSafeCall( cudaMemset(count_buf, 0, sizeof(unsigned int)) ); kernel<<>>((PtrStepSz) src, count_buf, twidth, theight); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); unsigned int count; - cvCudaSafeCall(cudaMemcpy(&count, count_buf, sizeof(unsigned int), cudaMemcpyDeviceToHost)); + cudaSafeCall(cudaMemcpy(&count, count_buf, sizeof(unsigned int), cudaMemcpyDeviceToHost)); return count; } @@ -1236,10 +1236,10 @@ namespace reduce Op op; rowsKernel<<>>(src, dst, op); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -1316,10 +1316,10 @@ namespace reduce Op op; colsKernel<<>>((PtrStepSz::vec_type>) src, (typename TypeVec::vec_type*) dst, op); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } diff --git a/modules/gpu/src/cuda/nlm.cu b/modules/gpu/src/cuda/nlm.cu index b62513318b..92bfccf37c 100644 --- a/modules/gpu/src/cuda/nlm.cu +++ b/modules/gpu/src/cuda/nlm.cu @@ -146,12 +146,12 @@ namespace cv { namespace gpu { namespace cudev float minus_h2_inv = -1.f/(h * h * VecTraits::cn); float noise_mult = minus_h2_inv/(block_window * block_window); - cvCudaSafeCall( cudaFuncSetCacheConfig (nlm_kernel >, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig (nlm_kernel >, cudaFuncCachePreferL1) ); nlm_kernel<<>>((PtrStepSz)src, (PtrStepSz)dst, b, search_radius, block_radius, noise_mult); - cvCudaSafeCall ( cudaGetLastError () ); + cudaSafeCall ( cudaGetLastError () ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -505,9 +505,9 @@ namespace cv { namespace gpu { namespace cudev fast_nlm_kernel<<>>(fnlm, (PtrStepSz)dst); - cvCudaSafeCall ( cudaGetLastError () ); + cudaSafeCall ( cudaGetLastError () ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void nlm_fast_gpu(const PtrStepSzb&, PtrStepSzb, PtrStepi, int, int, float, cudaStream_t); @@ -535,9 +535,9 @@ namespace cv { namespace gpu { namespace cudev dim3 g(divUp(lab.cols, b.x), divUp(lab.rows, b.y)); fnlm_split_kernel<<>>(lab, l, ab); - cvCudaSafeCall ( cudaGetLastError () ); + cudaSafeCall ( cudaGetLastError () ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void fnlm_merge_kernel(const PtrStepb l, const PtrStep ab, PtrStepSz lab) @@ -558,9 +558,9 @@ namespace cv { namespace gpu { namespace cudev dim3 g(divUp(lab.cols, b.x), divUp(lab.rows, b.y)); fnlm_merge_kernel<<>>(l, ab, lab); - cvCudaSafeCall ( cudaGetLastError () ); + cudaSafeCall ( cudaGetLastError () ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } }}} diff --git a/modules/gpu/src/cuda/optflowbm.cu b/modules/gpu/src/cuda/optflowbm.cu index 21b9b0e7f1..8f5b72efad 100644 --- a/modules/gpu/src/cuda/optflowbm.cu +++ b/modules/gpu/src/cuda/optflowbm.cu @@ -159,10 +159,10 @@ namespace optflowbm calcOptFlowBM<<>>(velx, vely, blockSize, shiftSize, usePrevious, maxX, maxY, acceptLevel, escapeLevel, ss, ssCount); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } @@ -402,10 +402,10 @@ namespace optflowbm_fast size_t smem = search_window * search_window * sizeof(int); optflowbm_fast_kernel<<>>(fbm, velx, vely); - cvCudaSafeCall ( cudaGetLastError () ); + cudaSafeCall ( cudaGetLastError () ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void calc(PtrStepSzb I0, PtrStepSzb I1, PtrStepSzf velx, PtrStepSzf vely, PtrStepi buffer, int search_window, int block_window, cudaStream_t stream); diff --git a/modules/gpu/src/cuda/optical_flow.cu b/modules/gpu/src/cuda/optical_flow.cu index 532cade920..d361bcfc63 100644 --- a/modules/gpu/src/cuda/optical_flow.cu +++ b/modules/gpu/src/cuda/optical_flow.cu @@ -119,9 +119,9 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(u_avg.cols, u_avg.rows); NeedleMapAverageKernel<<>>(u, v, u_avg, v_avg); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } __global__ void NeedleMapVertexKernel(const PtrStepSzf u_avg, const PtrStepf v_avg, float* vertex_data, float* color_data, float max_flow, float xscale, float yscale) @@ -210,9 +210,9 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(u_avg.cols, block.x), divUp(u_avg.rows, block.y)); NeedleMapVertexKernel<<>>(u_avg, v_avg, vertex_buffer, color_data, max_flow, xscale, yscale); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } }}} diff --git a/modules/gpu/src/cuda/optical_flow_farneback.cu b/modules/gpu/src/cuda/optical_flow_farneback.cu index 3d4b0fdec4..e7ff3a02f6 100644 --- a/modules/gpu/src/cuda/optical_flow_farneback.cu +++ b/modules/gpu/src/cuda/optical_flow_farneback.cu @@ -123,13 +123,13 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback int polyN, const float *g, const float *xg, const float *xxg, float ig11, float ig03, float ig33, float ig55) { - cvCudaSafeCall(cudaMemcpyToSymbol(c_g, g, (polyN + 1) * sizeof(*g))); - cvCudaSafeCall(cudaMemcpyToSymbol(c_xg, xg, (polyN + 1) * sizeof(*xg))); - cvCudaSafeCall(cudaMemcpyToSymbol(c_xxg, xxg, (polyN + 1) * sizeof(*xxg))); - cvCudaSafeCall(cudaMemcpyToSymbol(c_ig11, &ig11, sizeof(ig11))); - cvCudaSafeCall(cudaMemcpyToSymbol(c_ig03, &ig03, sizeof(ig03))); - cvCudaSafeCall(cudaMemcpyToSymbol(c_ig33, &ig33, sizeof(ig33))); - cvCudaSafeCall(cudaMemcpyToSymbol(c_ig55, &ig55, sizeof(ig55))); + cudaSafeCall(cudaMemcpyToSymbol(c_g, g, (polyN + 1) * sizeof(*g))); + cudaSafeCall(cudaMemcpyToSymbol(c_xg, xg, (polyN + 1) * sizeof(*xg))); + cudaSafeCall(cudaMemcpyToSymbol(c_xxg, xxg, (polyN + 1) * sizeof(*xxg))); + cudaSafeCall(cudaMemcpyToSymbol(c_ig11, &ig11, sizeof(ig11))); + cudaSafeCall(cudaMemcpyToSymbol(c_ig03, &ig03, sizeof(ig03))); + cudaSafeCall(cudaMemcpyToSymbol(c_ig33, &ig33, sizeof(ig33))); + cudaSafeCall(cudaMemcpyToSymbol(c_ig55, &ig55, sizeof(ig55))); } @@ -144,10 +144,10 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback else if (polyN == 7) polynomialExpansion<7><<>>(src.rows, src.cols, src, dst); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } @@ -244,7 +244,7 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback void setUpdateMatricesConsts() { static const float border[BORDER_SIZE + 1] = {0.14f, 0.14f, 0.4472f, 0.4472f, 0.4472f, 1.f}; - cvCudaSafeCall(cudaMemcpyToSymbol(c_border, border, (BORDER_SIZE + 1) * sizeof(*border))); + cudaSafeCall(cudaMemcpyToSymbol(c_border, border, (BORDER_SIZE + 1) * sizeof(*border))); } @@ -257,10 +257,10 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback updateMatrices<<>>(flowx.rows, flowx.cols, flowx, flowy, R0, R1, M); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } @@ -293,10 +293,10 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback updateFlow<<>>(flowx.rows, flowx.cols, M, flowx, flowy); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } @@ -424,10 +424,10 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback float boxAreaInv = 1.f / ((1 + 2*ksizeHalf) * (1 + 2*ksizeHalf)); boxFilter5<<>>(height, width, src, ksizeHalf, boxAreaInv, dst); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } @@ -443,10 +443,10 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback float boxAreaInv = 1.f / ((1 + 2*ksizeHalf) * (1 + 2*ksizeHalf)); boxFilter5<<>>(height, width, src, ksizeHalf, boxAreaInv, dst); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } @@ -494,7 +494,7 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback void setGaussianBlurKernel(const float *gKer, int ksizeHalf) { - cvCudaSafeCall(cudaMemcpyToSymbol(c_gKer, gKer, (ksizeHalf + 1) * sizeof(*gKer))); + cudaSafeCall(cudaMemcpyToSymbol(c_gKer, gKer, (ksizeHalf + 1) * sizeof(*gKer))); } @@ -511,10 +511,10 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback gaussianBlur<<>>(height, width, src, ksizeHalf, b, dst); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } @@ -606,10 +606,10 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback gaussianBlur5<<>>(height, width, src, ksizeHalf, b, dst); - cvCudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaGetLastError()); if (stream == 0) - cvCudaSafeCall(cudaDeviceSynchronize()); + cudaSafeCall(cudaDeviceSynchronize()); } diff --git a/modules/gpu/src/cuda/orb.cu b/modules/gpu/src/cuda/orb.cu index 2065c60e97..1e88648014 100644 --- a/modules/gpu/src/cuda/orb.cu +++ b/modules/gpu/src/cuda/orb.cu @@ -132,10 +132,10 @@ namespace cv { namespace gpu { namespace cudev HarrisResponses<<>>(img, loc, response, npoints, blockSize, harris_k); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////////////////////////////////////// @@ -145,7 +145,7 @@ namespace cv { namespace gpu { namespace cudev void loadUMax(const int* u_max, int count) { - cvCudaSafeCall( cudaMemcpyToSymbol(c_u_max, u_max, count * sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_u_max, u_max, count * sizeof(int)) ); } __global__ void IC_Angle(const PtrStepb image, const short2* loc_, float* angle, const int npoints, const int half_k) @@ -214,10 +214,10 @@ namespace cv { namespace gpu { namespace cudev IC_Angle<<>>(image, loc, angle, npoints, half_k); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////////////////////////////////////// @@ -382,10 +382,10 @@ namespace cv { namespace gpu { namespace cudev break; } - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////////////////////////////////////// @@ -413,10 +413,10 @@ namespace cv { namespace gpu { namespace cudev mergeLocation<<>>(loc, x, y, npoints, scale); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } }}} diff --git a/modules/gpu/src/cuda/pyr_down.cu b/modules/gpu/src/cuda/pyr_down.cu index c8bcb7c999..904f549bad 100644 --- a/modules/gpu/src/cuda/pyr_down.cu +++ b/modules/gpu/src/cuda/pyr_down.cu @@ -181,10 +181,10 @@ namespace cv { namespace gpu { namespace cudev B b(src.rows, src.cols); pyrDown<<>>(src, dst, b, dst.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) diff --git a/modules/gpu/src/cuda/pyr_up.cu b/modules/gpu/src/cuda/pyr_up.cu index 279d02421a..36a72274cf 100644 --- a/modules/gpu/src/cuda/pyr_up.cu +++ b/modules/gpu/src/cuda/pyr_up.cu @@ -150,10 +150,10 @@ namespace cv { namespace gpu { namespace cudev const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); pyrUp<<>>(src, dst); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template void pyrUp_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) diff --git a/modules/gpu/src/cuda/pyrlk.cu b/modules/gpu/src/cuda/pyrlk.cu index 5b99becce6..410666fcfb 100644 --- a/modules/gpu/src/cuda/pyrlk.cu +++ b/modules/gpu/src/cuda/pyrlk.cu @@ -320,10 +320,10 @@ namespace pyrlk else sparseKernel<<>>(prevPts, nextPts, status, err, level, rows, cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } template @@ -474,14 +474,14 @@ namespace pyrlk void loadConstants(int2 winSize, int iters) { - cvCudaSafeCall( cudaMemcpyToSymbol(c_winSize_x, &winSize.x, sizeof(int)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(c_winSize_y, &winSize.y, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_winSize_x, &winSize.x, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_winSize_y, &winSize.y, sizeof(int)) ); int2 halfWin = make_int2((winSize.x - 1) / 2, (winSize.y - 1) / 2); - cvCudaSafeCall( cudaMemcpyToSymbol(c_halfWin_x, &halfWin.x, sizeof(int)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(c_halfWin_y, &halfWin.y, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_halfWin_x, &halfWin.x, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_halfWin_y, &halfWin.y, sizeof(int)) ); - cvCudaSafeCall( cudaMemcpyToSymbol(c_iters, &iters, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_iters, &iters, sizeof(int)) ); } void sparse1(PtrStepSzf I, PtrStepSzf J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, @@ -544,16 +544,16 @@ namespace pyrlk if (err.data) { denseKernel<<>>(u, v, prevU, prevV, err, I.rows, I.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } else { denseKernel<<>>(u, v, prevU, prevV, PtrStepf(), I.rows, I.cols); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } if (stream == 0) - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } } diff --git a/modules/gpu/src/cuda/remap.cu b/modules/gpu/src/cuda/remap.cu index dcc6da22c1..dd2c669159 100644 --- a/modules/gpu/src/cuda/remap.cu +++ b/modules/gpu/src/cuda/remap.cu @@ -81,7 +81,7 @@ namespace cv { namespace gpu { namespace cudev Filter< BorderReader< PtrStep, B > > filter_src(brdSrc); remap<<>>(filter_src, mapx, mapy, dst); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); } }; @@ -102,9 +102,9 @@ namespace cv { namespace gpu { namespace cudev Filter< BorderReader< PtrStep, B > > filter_src(brdSrc); remap<<>>(filter_src, mapx, mapy, dst); - cvCudaSafeCall( cudaGetLastError() ); + cudaSafeCall( cudaGetLastError() ); - cvCudaSafeCall( cudaDeviceSynchronize() ); + cudaSafeCall( cudaDeviceSynchronize() ); } }; @@ -135,8 +135,8 @@ namespace cv { namespace gpu { namespace cudev BorderReader< tex_remap_ ## type ##_reader, B > brdSrc(texSrc, brd); \ Filter< BorderReader< tex_remap_ ## type ##_reader, B > > filter_src(brdSrc); \ remap<<>>(filter_src, mapx, mapy, dst); \ - cvCudaSafeCall( cudaGetLastError() ); \ - cvCudaSafeCall( cudaDeviceSynchronize() ); \ + cudaSafeCall( cudaGetLastError() ); \ + cudaSafeCall( cudaDeviceSynchronize() ); \ } \ }; \ template