diff --git a/modules/gpu/src/cuda/copy_make_border.cu b/modules/gpu/src/cuda/copy_make_border.cu index 5603742302..aafcdf5c43 100644 --- a/modules/gpu/src/cuda/copy_make_border.cu +++ b/modules/gpu/src/cuda/copy_make_border.cu @@ -45,7 +45,7 @@ BEGIN_OPENCV_DEVICE_NAMESPACE -namespace copy_make_border { +namespace imgproc { template __global__ void copyMakeBorder(const Ptr2D src, DevMem2D_ dst, int top, int left) { @@ -124,6 +124,6 @@ template void copyMakeBorder_gpu(const DevMem2Db& src, const DevMem2Db template void copyMakeBorder_gpu(const DevMem2Db& src, const DevMem2Db& dst, int top, int left, int borderMode, const float* borderValue, cudaStream_t stream); template void copyMakeBorder_gpu(const DevMem2Db& src, const DevMem2Db& dst, int top, int left, int borderMode, const float* borderValue, cudaStream_t stream); -} // namespace copy_make_border +} // namespace imgproc END_OPENCV_DEVICE_NAMESPACE diff --git a/modules/gpu/src/cuda/pyr_down.cu b/modules/gpu/src/cuda/pyr_down.cu index 7c5077e293..75cb70d9c4 100644 --- a/modules/gpu/src/cuda/pyr_down.cu +++ b/modules/gpu/src/cuda/pyr_down.cu @@ -48,7 +48,7 @@ BEGIN_OPENCV_DEVICE_NAMESPACE -namespace pyr_down { +namespace imgproc { template __global__ void pyrDown(const PtrStep src, PtrStep dst, const B b, int dst_cols) { @@ -182,6 +182,6 @@ template void pyrDown_gpu(const DevMem2Db& src, const DevMem2Db& dst, template void pyrDown_gpu(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream); template void pyrDown_gpu(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream); -} // namespace pyr_down +} // namespace imgproc END_OPENCV_DEVICE_NAMESPACE diff --git a/modules/gpu/src/cuda/pyr_up.cu b/modules/gpu/src/cuda/pyr_up.cu index 35ddb3a865..0d24581a36 100644 --- a/modules/gpu/src/cuda/pyr_up.cu +++ b/modules/gpu/src/cuda/pyr_up.cu @@ -48,7 +48,7 @@ BEGIN_OPENCV_DEVICE_NAMESPACE -namespace pyr_up { +namespace imgproc { template __global__ void pyrUp(const PtrStep src, DevMem2D_ dst, const B b) { @@ -177,6 +177,6 @@ template void pyrUp_gpu(const DevMem2Db& src, const DevMem2Db& dst, in template void pyrUp_gpu(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream); template void pyrUp_gpu(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream); -} // namespace pyr_up +} // namespace imgproc END_OPENCV_DEVICE_NAMESPACE diff --git a/modules/gpu/src/cuda/remap.cu b/modules/gpu/src/cuda/remap.cu index 0fda25e884..650357d6cc 100644 --- a/modules/gpu/src/cuda/remap.cu +++ b/modules/gpu/src/cuda/remap.cu @@ -49,7 +49,7 @@ BEGIN_OPENCV_DEVICE_NAMESPACE -namespace remap { +namespace imgproc { template __global__ void remap(const Ptr2D src, const PtrStepf mapx, const PtrStepf mapy, DevMem2D_ dst) { @@ -249,6 +249,6 @@ template void remap_gpu(const DevMem2Db& src, const DevMem2Df& xmap, con template void remap_gpu(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc); template void remap_gpu(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc); -} // namespace remap +} // namespace imgproc END_OPENCV_DEVICE_NAMESPACE diff --git a/modules/gpu/src/cuda/resize.cu b/modules/gpu/src/cuda/resize.cu index b79716282b..4294dc1095 100644 --- a/modules/gpu/src/cuda/resize.cu +++ b/modules/gpu/src/cuda/resize.cu @@ -49,7 +49,7 @@ BEGIN_OPENCV_DEVICE_NAMESPACE -namespace resize { +namespace imgproc { template __global__ void resize(const Ptr2D src, float fx, float fy, DevMem2D_ dst) { @@ -260,6 +260,6 @@ template void resize_gpu(const DevMem2Db& src, float fx, float fy, const template void resize_gpu(const DevMem2Db& src, float fx, float fy, const DevMem2Db& dst, int interpolation, cudaStream_t stream); template void resize_gpu(const DevMem2Db& src, float fx, float fy, const DevMem2Db& dst, int interpolation, cudaStream_t stream); -} // namespace resize +} // namespace imgproc END_OPENCV_DEVICE_NAMESPACE diff --git a/modules/gpu/src/imgproc.cpp b/modules/gpu/src/imgproc.cpp index 8973280072..7d171c6844 100644 --- a/modules/gpu/src/imgproc.cpp +++ b/modules/gpu/src/imgproc.cpp @@ -109,7 +109,7 @@ void cv::gpu::CannyBuf::release() { throw_nogpu(); } BEGIN_OPENCV_DEVICE_NAMESPACE -namespace remap +namespace imgproc { template void remap_gpu(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, @@ -120,7 +120,7 @@ END_OPENCV_DEVICE_NAMESPACE void cv::gpu::remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap, int interpolation, int borderMode, const Scalar& borderValue, Stream& stream) { - using namespace OPENCV_DEVICE_NAMESPACE_ remap; + using namespace OPENCV_DEVICE_NAMESPACE_ imgproc; typedef void (*caller_t)(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc); @@ -311,7 +311,7 @@ void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, BEGIN_OPENCV_DEVICE_NAMESPACE -namespace resize +namespace imgproc { template void resize_gpu(const DevMem2Db& src, float fx, float fy, const DevMem2Db& dst, int interpolation, cudaStream_t stream); } @@ -380,7 +380,7 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub } else { - using namespace OPENCV_DEVICE_NAMESPACE_ resize; + using namespace OPENCV_DEVICE_NAMESPACE_ imgproc; typedef void (*caller_t)(const DevMem2Db& src, float fx, float fy, const DevMem2Db& dst, int interpolation, cudaStream_t stream); static const caller_t callers[6][4] = @@ -402,7 +402,7 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub BEGIN_OPENCV_DEVICE_NAMESPACE -namespace copy_make_border +namespace imgproc { template void copyMakeBorder_gpu(const DevMem2Db& src, const DevMem2Db& dst, int top, int left, int borderMode, const T* borderValue, cudaStream_t stream); } @@ -413,7 +413,7 @@ namespace { template void copyMakeBorder_caller(const DevMem2Db& src, const DevMem2Db& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream) { - using namespace OPENCV_DEVICE_NAMESPACE_ copy_make_border; + using namespace OPENCV_DEVICE_NAMESPACE_ imgproc; Scalar_ val(saturate_cast(value[0]), saturate_cast(value[1]), saturate_cast(value[2]), saturate_cast(value[3])); @@ -1813,7 +1813,7 @@ void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, BEGIN_OPENCV_DEVICE_NAMESPACE -namespace pyr_down +namespace imgproc { template void pyrDown_gpu(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream); } @@ -1822,7 +1822,7 @@ END_OPENCV_DEVICE_NAMESPACE void cv::gpu::pyrDown(const GpuMat& src, GpuMat& dst, int borderType, Stream& stream) { - using namespace OPENCV_DEVICE_NAMESPACE_ pyr_down; + using namespace OPENCV_DEVICE_NAMESPACE_ imgproc; typedef void (*func_t)(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream); @@ -1853,7 +1853,7 @@ void cv::gpu::pyrDown(const GpuMat& src, GpuMat& dst, int borderType, Stream& st BEGIN_OPENCV_DEVICE_NAMESPACE -namespace pyr_up +namespace imgproc { template void pyrUp_gpu(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream); } @@ -1862,7 +1862,7 @@ END_OPENCV_DEVICE_NAMESPACE void cv::gpu::pyrUp(const GpuMat& src, GpuMat& dst, int borderType, Stream& stream) { - using namespace OPENCV_DEVICE_NAMESPACE_ pyr_up; + using namespace OPENCV_DEVICE_NAMESPACE_ imgproc; typedef void (*func_t)(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream); diff --git a/modules/gpu/src/opencv2/gpu/device/functional.hpp b/modules/gpu/src/opencv2/gpu/device/functional.hpp index 5f29c0de6f..d885c1aa69 100644 --- a/modules/gpu/src/opencv2/gpu/device/functional.hpp +++ b/modules/gpu/src/opencv2/gpu/device/functional.hpp @@ -272,61 +272,61 @@ OPENCV_GPU_IMPLEMENT_MINMAX(minimum, double, ::fmin) // Math functions -#define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(func) \ - template struct func ## _func : unary_function \ +#define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \ + template struct name ## _func : unary_function \ { \ __device__ __forceinline__ float operator ()(typename TypeTraits::ParameterType v) const \ { \ - return :: ## func ## f(v); \ + return func ## f(v); \ } \ }; \ - template <> struct func ## _func : unary_function \ + template <> struct name ## _func : unary_function \ { \ __device__ __forceinline__ double operator ()(double v) const \ { \ - return :: ## func(v); \ + return func(v); \ } \ }; -#define OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(func) \ - template struct func ## _func : binary_function \ +#define OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(name, func) \ + template struct name ## _func : binary_function \ { \ __device__ __forceinline__ float operator ()(typename TypeTraits::ParameterType v1, typename TypeTraits::ParameterType v2) const \ { \ - return :: ## func ## f(v1, v2); \ + return func ## f(v1, v2); \ } \ }; \ - template <> struct func ## _func : binary_function \ + template <> struct name ## _func : binary_function \ { \ __device__ __forceinline__ double operator ()(double v1, double v2) const \ { \ - return :: ## func(v1, v2); \ + return func(v1, v2); \ } \ }; -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(fabs) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sqrt) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp2) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp10) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log2) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log10) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sin) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cos) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tan) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asin) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acos) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atan) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sinh) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cosh) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tanh) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asinh) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acosh) -OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atanh) - -OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(hypot) -OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(atan2) -OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(pow) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(fabs, ::fabs) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp, ::exp) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp10, ::exp10) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log, ::log) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log2, ::log2) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log10, ::log10) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sin, ::sin) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cos, ::cos) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tan, ::tan) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asin, ::asin) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acos, ::acos) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atan, ::atan) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sinh, ::sinh) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cosh, ::cosh) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tanh, ::tanh) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asinh, ::asinh) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acosh, ::acosh) +OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atanh, ::atanh) + +OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(hypot, ::hypot) +OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(atan2, ::atan2) +OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(pow, ::pow) #undef OPENCV_GPU_IMPLEMENT_UN_FUNCTOR #undef OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR