diff --git a/modules/core/include/opencv2/core/cuda/block.hpp b/modules/core/include/opencv2/core/cuda/block.hpp index b2bcb9f503..a6464079ad 100644 --- a/modules/core/include/opencv2/core/cuda/block.hpp +++ b/modules/core/include/opencv2/core/cuda/block.hpp @@ -43,7 +43,7 @@ #ifndef __OPENCV_GPU_DEVICE_BLOCK_HPP__ #define __OPENCV_GPU_DEVICE_BLOCK_HPP__ -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { struct Block { diff --git a/modules/core/include/opencv2/core/cuda/border_interpolate.hpp b/modules/core/include/opencv2/core/cuda/border_interpolate.hpp index 3854e8429a..1347a2f849 100644 --- a/modules/core/include/opencv2/core/cuda/border_interpolate.hpp +++ b/modules/core/include/opencv2/core/cuda/border_interpolate.hpp @@ -47,7 +47,7 @@ #include "vec_traits.hpp" #include "vec_math.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { ////////////////////////////////////////////////////////////// // BrdConstant @@ -709,6 +709,6 @@ namespace cv { namespace gpu { namespace cuda const int width; const D val; }; -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_BORDER_INTERPOLATE_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/color.hpp b/modules/core/include/opencv2/core/cuda/color.hpp index c625308f6b..a2b772d8b3 100644 --- a/modules/core/include/opencv2/core/cuda/color.hpp +++ b/modules/core/include/opencv2/core/cuda/color.hpp @@ -45,7 +45,7 @@ #include "detail/color_detail.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { // All OPENCV_GPU_IMPLEMENT_*_TRAITS(ColorSpace1_to_ColorSpace2, ...) macros implements // template class ColorSpace1_to_ColorSpace2_traits @@ -296,6 +296,6 @@ namespace cv { namespace gpu { namespace cuda OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lbgra, 4, 4, false, 0) #undef OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_BORDER_INTERPOLATE_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/common.hpp b/modules/core/include/opencv2/core/cuda/common.hpp index 446cfef85c..acc8adbdfb 100644 --- a/modules/core/include/opencv2/core/cuda/common.hpp +++ b/modules/core/include/opencv2/core/cuda/common.hpp @@ -56,18 +56,18 @@ #endif #endif -namespace cv { namespace gpu { namespace cuda { - static inline void checkError(cudaError_t err, const char* file, const int line, const char* func) +namespace cv { namespace gpu { + static inline void checkCudaError(cudaError_t err, const char* file, const int line, const char* func) { if (cudaSuccess != err) cv::error(cv::Error::GpuApiCallError, cudaGetErrorString(err), func, file, line); } -}}} +}} #if defined(__GNUC__) - #define cvCudaSafeCall(expr) cv::gpu::cuda::checkError(expr, __FILE__, __LINE__, __func__) + #define cvCudaSafeCall(expr) cv::gpu::checkCudaError(expr, __FILE__, __LINE__, __func__) #else /* defined(__CUDACC__) || defined(__MSVC__) */ - #define cvCudaSafeCall(expr) cv::gpu::cuda::checkError(expr, __FILE__, __LINE__, "") + #define cvCudaSafeCall(expr) cv::gpu::checkCudaError(expr, __FILE__, __LINE__, "") #endif namespace cv { namespace gpu @@ -94,8 +94,7 @@ namespace cv { namespace gpu BORDER_WRAP_GPU }; -#ifdef __CUDACC__ - namespace cuda + namespace cudev { __host__ __device__ __forceinline__ int divUp(int total, int grain) { @@ -108,7 +107,6 @@ namespace cv { namespace gpu cvCudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) ); } } -#endif // __CUDACC__ }} diff --git a/modules/core/include/opencv2/core/cuda/datamov_utils.hpp b/modules/core/include/opencv2/core/cuda/datamov_utils.hpp index d2aadf9f12..10df540936 100644 --- a/modules/core/include/opencv2/core/cuda/datamov_utils.hpp +++ b/modules/core/include/opencv2/core/cuda/datamov_utils.hpp @@ -45,7 +45,7 @@ #include "common.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 @@ -100,6 +100,6 @@ namespace cv { namespace gpu { namespace cuda #undef OPENCV_GPU_ASM_PTR #endif // __CUDA_ARCH__ >= 200 -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_DATAMOV_UTILS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/detail/color_detail.hpp b/modules/core/include/opencv2/core/cuda/detail/color_detail.hpp index 5853475d35..9246b0dafe 100644 --- a/modules/core/include/opencv2/core/cuda/detail/color_detail.hpp +++ b/modules/core/include/opencv2/core/cuda/detail/color_detail.hpp @@ -49,7 +49,7 @@ #include "../limits.hpp" #include "../functional.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { #ifndef CV_DESCALE #define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n)) @@ -149,7 +149,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -222,7 +222,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(name, scn, bidx, green_bits) \ struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2RGB5x5 functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2RGB5x5 functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -303,7 +303,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(name, dcn, bidx, green_bits) \ struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB5x52RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB5x52RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -350,7 +350,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_GRAY2RGB_TRAITS(name, dcn) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::Gray2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::Gray2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -392,7 +392,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_GRAY2RGB5x5_TRAITS(name, green_bits) \ struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::Gray2RGB5x5 functor_type; \ + typedef ::cv::gpu::cudev::color_detail::Gray2RGB5x5 functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -434,7 +434,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB5x52GRAY_TRAITS(name, green_bits) \ struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB5x52Gray functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB5x52Gray functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -486,7 +486,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS(name, scn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2Gray functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2Gray functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -539,7 +539,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2YUV functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2YUV functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -629,7 +629,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::YUV2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::YUV2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -710,7 +710,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2YCrCb functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2YCrCb functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -791,7 +791,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::YCrCb2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::YCrCb2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -869,7 +869,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2XYZ functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2XYZ functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -946,7 +946,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::XYZ2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::XYZ2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1086,7 +1086,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2HSV functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2HSV functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1094,7 +1094,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template struct name ## _full_traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2HSV functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2HSV functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1102,7 +1102,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template <> struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2HSV functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2HSV functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1110,7 +1110,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template <> struct name ## _full_traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2HSV functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2HSV functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1228,7 +1228,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::HSV2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::HSV2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1236,7 +1236,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template struct name ## _full_traits \ { \ - typedef ::cv::gpu::cuda::color_detail::HSV2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::HSV2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1244,7 +1244,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template <> struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::HSV2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::HSV2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1252,7 +1252,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template <> struct name ## _full_traits \ { \ - typedef ::cv::gpu::cuda::color_detail::HSV2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::HSV2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1363,7 +1363,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2HLS functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2HLS functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1371,7 +1371,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template struct name ## _full_traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2HLS functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2HLS functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1379,7 +1379,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template <> struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2HLS functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2HLS functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1387,7 +1387,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template <> struct name ## _full_traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2HLS functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2HLS functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1505,7 +1505,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::HLS2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::HLS2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1513,7 +1513,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template struct name ## _full_traits \ { \ - typedef ::cv::gpu::cuda::color_detail::HLS2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::HLS2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1521,7 +1521,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template <> struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::HLS2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::HLS2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1529,7 +1529,7 @@ namespace cv { namespace gpu { namespace cuda }; \ template <> struct name ## _full_traits \ { \ - typedef ::cv::gpu::cuda::color_detail::HLS2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::HLS2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1674,7 +1674,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(name, scn, dcn, srgb, blueIdx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2Lab functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2Lab functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1787,7 +1787,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::Lab2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::Lab2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1886,7 +1886,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(name, scn, dcn, srgb, blueIdx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::RGB2Luv functor_type; \ + typedef ::cv::gpu::cudev::color_detail::RGB2Luv functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1987,7 +1987,7 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::cuda::color_detail::Luv2RGB functor_type; \ + typedef ::cv::gpu::cudev::color_detail::Luv2RGB functor_type; \ static __host__ __device__ __forceinline__ functor_type create_functor() \ { \ return functor_type(); \ @@ -1996,6 +1996,6 @@ namespace cv { namespace gpu { namespace cuda #undef CV_DESCALE -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_COLOR_DETAIL_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/detail/reduce.hpp b/modules/core/include/opencv2/core/cuda/detail/reduce.hpp index 6866016c91..eba9b41a78 100644 --- a/modules/core/include/opencv2/core/cuda/detail/reduce.hpp +++ b/modules/core/include/opencv2/core/cuda/detail/reduce.hpp @@ -47,7 +47,7 @@ #include "../warp.hpp" #include "../warp_shuffle.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace reduce_detail { diff --git a/modules/core/include/opencv2/core/cuda/detail/reduce_key_val.hpp b/modules/core/include/opencv2/core/cuda/detail/reduce_key_val.hpp index dde37dc759..1049e6714c 100644 --- a/modules/core/include/opencv2/core/cuda/detail/reduce_key_val.hpp +++ b/modules/core/include/opencv2/core/cuda/detail/reduce_key_val.hpp @@ -47,7 +47,7 @@ #include "../warp.hpp" #include "../warp_shuffle.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace reduce_key_val_detail { 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 731aa15ce1..906e0e0203 100644 --- a/modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp +++ b/modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp @@ -47,7 +47,7 @@ #include "../vec_traits.hpp" #include "../functional.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace transform_detail { @@ -390,6 +390,6 @@ namespace cv { namespace gpu { namespace cuda } }; } // namespace transform_detail -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_TRANSFORM_DETAIL_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/detail/type_traits_detail.hpp b/modules/core/include/opencv2/core/cuda/detail/type_traits_detail.hpp index ee92184c3e..4292d88003 100644 --- a/modules/core/include/opencv2/core/cuda/detail/type_traits_detail.hpp +++ b/modules/core/include/opencv2/core/cuda/detail/type_traits_detail.hpp @@ -46,7 +46,7 @@ #include "../common.hpp" #include "../vec_traits.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace type_traits_detail { @@ -182,6 +182,6 @@ namespace cv { namespace gpu { namespace cuda enum { value = 1 }; }; } // namespace type_traits_detail -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_TYPE_TRAITS_DETAIL_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/detail/vec_distance_detail.hpp b/modules/core/include/opencv2/core/cuda/detail/vec_distance_detail.hpp index 435e69c8ef..a2a31a703c 100644 --- a/modules/core/include/opencv2/core/cuda/detail/vec_distance_detail.hpp +++ b/modules/core/include/opencv2/core/cuda/detail/vec_distance_detail.hpp @@ -45,7 +45,7 @@ #include "../datamov_utils.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace vec_distance_detail { @@ -112,6 +112,6 @@ namespace cv { namespace gpu { namespace cuda } }; } // namespace vec_distance_detail -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_VEC_DISTANCE_DETAIL_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp b/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp index 41f9ecce6b..aa20e53b82 100644 --- a/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp +++ b/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp @@ -43,7 +43,7 @@ #ifndef __OPENCV_GPU_DYNAMIC_SMEM_HPP__ #define __OPENCV_GPU_DYNAMIC_SMEM_HPP__ -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct DynamicSharedMem { diff --git a/modules/core/include/opencv2/core/cuda/emulation.hpp b/modules/core/include/opencv2/core/cuda/emulation.hpp index e2c7437886..3df26468b2 100644 --- a/modules/core/include/opencv2/core/cuda/emulation.hpp +++ b/modules/core/include/opencv2/core/cuda/emulation.hpp @@ -45,7 +45,7 @@ #include "warp_reduce.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { struct Emulation { @@ -133,6 +133,6 @@ namespace cv { namespace gpu { namespace cuda } }; }; -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* OPENCV_GPU_EMULATION_HPP_ */ diff --git a/modules/core/include/opencv2/core/cuda/filters.hpp b/modules/core/include/opencv2/core/cuda/filters.hpp index 607aecf2a0..19a8c5883f 100644 --- a/modules/core/include/opencv2/core/cuda/filters.hpp +++ b/modules/core/include/opencv2/core/cuda/filters.hpp @@ -48,7 +48,7 @@ #include "vec_math.hpp" #include "type_traits.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct PointFilter { @@ -273,6 +273,6 @@ namespace cv { namespace gpu { namespace cuda float scale_x, scale_y; int width, haight; }; -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_FILTERS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/funcattrib.hpp b/modules/core/include/opencv2/core/cuda/funcattrib.hpp index e97b5bec61..46ef81926c 100644 --- a/modules/core/include/opencv2/core/cuda/funcattrib.hpp +++ b/modules/core/include/opencv2/core/cuda/funcattrib.hpp @@ -45,7 +45,7 @@ #include -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template void printFuncAttrib(Func& func) @@ -66,6 +66,6 @@ namespace cv { namespace gpu { namespace cuda printf("\n"); fflush(stdout); } -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_ */ diff --git a/modules/core/include/opencv2/core/cuda/functional.hpp b/modules/core/include/opencv2/core/cuda/functional.hpp index 5ac90eb471..506ccd8768 100644 --- a/modules/core/include/opencv2/core/cuda/functional.hpp +++ b/modules/core/include/opencv2/core/cuda/functional.hpp @@ -49,7 +49,7 @@ #include "type_traits.hpp" #include "device_functions.h" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { // Function Objects template struct unary_function : public std::unary_function {}; @@ -786,6 +786,6 @@ namespace cv { namespace gpu { namespace cuda #define OPENCV_GPU_TRANSFORM_FUNCTOR_TRAITS(type) \ template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type > -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_FUNCTIONAL_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/limits.hpp b/modules/core/include/opencv2/core/cuda/limits.hpp index fc5e7b2340..4b265da0e5 100644 --- a/modules/core/include/opencv2/core/cuda/limits.hpp +++ b/modules/core/include/opencv2/core/cuda/limits.hpp @@ -46,7 +46,7 @@ #include #include "common.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct numeric_limits { @@ -230,6 +230,6 @@ namespace cv { namespace gpu { namespace cuda __device__ __forceinline__ static type signaling_NaN(); static const bool is_signed = true; }; -}}} // namespace cv { namespace gpu { namespace cuda { +}}} // namespace cv { namespace gpu { namespace cudev { #endif // __OPENCV_GPU_LIMITS_GPU_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/reduce.hpp b/modules/core/include/opencv2/core/cuda/reduce.hpp index ca16a20e0a..722e2bbeb6 100644 --- a/modules/core/include/opencv2/core/cuda/reduce.hpp +++ b/modules/core/include/opencv2/core/cuda/reduce.hpp @@ -47,7 +47,7 @@ #include "detail/reduce.hpp" #include "detail/reduce_key_val.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template __device__ __forceinline__ void reduce(volatile T* smem, T& val, unsigned int tid, const Op& op) diff --git a/modules/core/include/opencv2/core/cuda/saturate_cast.hpp b/modules/core/include/opencv2/core/cuda/saturate_cast.hpp index f80a43abe5..b30f5e7ce0 100644 --- a/modules/core/include/opencv2/core/cuda/saturate_cast.hpp +++ b/modules/core/include/opencv2/core/cuda/saturate_cast.hpp @@ -45,7 +45,7 @@ #include "common.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template __device__ __forceinline__ _Tp saturate_cast(uchar v) { return _Tp(v); } template __device__ __forceinline__ _Tp saturate_cast(schar v) { return _Tp(v); } diff --git a/modules/core/include/opencv2/core/cuda/scan.hpp b/modules/core/include/opencv2/core/cuda/scan.hpp index 0493ee3608..ecde123bb3 100644 --- a/modules/core/include/opencv2/core/cuda/scan.hpp +++ b/modules/core/include/opencv2/core/cuda/scan.hpp @@ -48,7 +48,7 @@ #include "opencv2/core/cuda/warp.hpp" #include "opencv2/core/cuda/warp_shuffle.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { enum ScanKind { EXCLUSIVE = 0, INCLUSIVE = 1 }; @@ -174,13 +174,13 @@ namespace cv { namespace gpu { namespace cuda __device__ T warpScanInclusive(T idata, volatile T* s_Data, unsigned int tid) { #if __CUDA_ARCH__ >= 300 - const unsigned int laneId = cv::gpu::cuda::Warp::laneId(); + const unsigned int laneId = cv::gpu::cudev::Warp::laneId(); // scan on shuffl functions #pragma unroll for (int i = 1; i <= (OPENCV_GPU_WARP_SIZE / 2); i *= 2) { - const T n = cv::gpu::cuda::shfl_up(idata, i); + const T n = cv::gpu::cudev::shfl_up(idata, i); if (laneId >= i) idata += n; } diff --git a/modules/core/include/opencv2/core/cuda/simd_functions.hpp b/modules/core/include/opencv2/core/cuda/simd_functions.hpp index b6d2fe13c9..aedd5632f2 100644 --- a/modules/core/include/opencv2/core/cuda/simd_functions.hpp +++ b/modules/core/include/opencv2/core/cuda/simd_functions.hpp @@ -123,7 +123,7 @@ vmin4(a,b) per-byte unsigned minimum: min(a, b) */ -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { // 2 diff --git a/modules/core/include/opencv2/core/cuda/static_check.hpp b/modules/core/include/opencv2/core/cuda/static_check.hpp index db1bb8e057..4f0ee8e17f 100644 --- a/modules/core/include/opencv2/core/cuda/static_check.hpp +++ b/modules/core/include/opencv2/core/cuda/static_check.hpp @@ -51,7 +51,7 @@ namespace cv { namespace gpu { - namespace cuda + namespace cudev { template struct Static {}; diff --git a/modules/core/include/opencv2/core/cuda/transform.hpp b/modules/core/include/opencv2/core/cuda/transform.hpp index bf9b68bedd..7c82e36469 100644 --- a/modules/core/include/opencv2/core/cuda/transform.hpp +++ b/modules/core/include/opencv2/core/cuda/transform.hpp @@ -47,7 +47,7 @@ #include "utility.hpp" #include "detail/transform_detail.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template static inline void transform(PtrStepSz src, PtrStepSz dst, UnOp op, const Mask& mask, cudaStream_t stream) diff --git a/modules/core/include/opencv2/core/cuda/type_traits.hpp b/modules/core/include/opencv2/core/cuda/type_traits.hpp index 0fb5fdac5b..8a58264bfc 100644 --- a/modules/core/include/opencv2/core/cuda/type_traits.hpp +++ b/modules/core/include/opencv2/core/cuda/type_traits.hpp @@ -45,7 +45,7 @@ #include "detail/type_traits_detail.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct IsSimpleParameter { diff --git a/modules/core/include/opencv2/core/cuda/utility.hpp b/modules/core/include/opencv2/core/cuda/utility.hpp index 79973850de..83fe388950 100644 --- a/modules/core/include/opencv2/core/cuda/utility.hpp +++ b/modules/core/include/opencv2/core/cuda/utility.hpp @@ -46,7 +46,7 @@ #include "saturate_cast.hpp" #include "datamov_utils.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { #define OPENCV_GPU_LOG_WARP_SIZE (5) #define OPENCV_GPU_WARP_SIZE (1 << OPENCV_GPU_LOG_WARP_SIZE) @@ -208,6 +208,6 @@ namespace cv { namespace gpu { namespace cuda return false; } -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_UTILITY_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/vec_distance.hpp b/modules/core/include/opencv2/core/cuda/vec_distance.hpp index 552100161f..4b88410207 100644 --- a/modules/core/include/opencv2/core/cuda/vec_distance.hpp +++ b/modules/core/include/opencv2/core/cuda/vec_distance.hpp @@ -47,7 +47,7 @@ #include "functional.hpp" #include "detail/vec_distance_detail.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct L1Dist { @@ -219,6 +219,6 @@ namespace cv { namespace gpu { namespace cuda U vec1Vals[MAX_LEN / THREAD_DIM]; }; -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_VEC_DISTANCE_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/vec_math.hpp b/modules/core/include/opencv2/core/cuda/vec_math.hpp index 238c71e28e..e4f981476e 100644 --- a/modules/core/include/opencv2/core/cuda/vec_math.hpp +++ b/modules/core/include/opencv2/core/cuda/vec_math.hpp @@ -47,7 +47,7 @@ #include "vec_traits.hpp" #include "functional.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace vec_math_detail { @@ -325,6 +325,6 @@ namespace cv { namespace gpu { namespace cuda #undef OPENCV_GPU_IMPLEMENT_VEC_BINOP #undef OPENCV_GPU_IMPLEMENT_VEC_OP #undef OPENCV_GPU_IMPLEMENT_VEC_INT_OP -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_VECMATH_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/vec_traits.hpp b/modules/core/include/opencv2/core/cuda/vec_traits.hpp index 9cdf64b7b3..304b05c919 100644 --- a/modules/core/include/opencv2/core/cuda/vec_traits.hpp +++ b/modules/core/include/opencv2/core/cuda/vec_traits.hpp @@ -45,7 +45,7 @@ #include "common.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TypeVec; @@ -275,6 +275,6 @@ namespace cv { namespace gpu { namespace cuda static __device__ __host__ __forceinline__ char8 make(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7) {return make_char8(a0, a1, a2, a3, a4, a5, a6, a7);} static __device__ __host__ __forceinline__ char8 make(const schar* v) {return make_char8(v[0], v[1], v[2], v[3], v[4], v[5], v[6], v[7]);} }; -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif // __OPENCV_GPU_VEC_TRAITS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/warp.hpp b/modules/core/include/opencv2/core/cuda/warp.hpp index 14cf6cba71..6d2b7745f1 100644 --- a/modules/core/include/opencv2/core/cuda/warp.hpp +++ b/modules/core/include/opencv2/core/cuda/warp.hpp @@ -43,7 +43,7 @@ #ifndef __OPENCV_GPU_DEVICE_WARP_HPP__ #define __OPENCV_GPU_DEVICE_WARP_HPP__ -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { struct Warp { @@ -126,6 +126,6 @@ namespace cv { namespace gpu { namespace cuda *t = value; } }; -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* __OPENCV_GPU_DEVICE_WARP_HPP__ */ diff --git a/modules/core/include/opencv2/core/cuda/warp_reduce.hpp b/modules/core/include/opencv2/core/cuda/warp_reduce.hpp index 7d4d838e8e..82185e8c0c 100644 --- a/modules/core/include/opencv2/core/cuda/warp_reduce.hpp +++ b/modules/core/include/opencv2/core/cuda/warp_reduce.hpp @@ -43,7 +43,7 @@ #ifndef OPENCV_GPU_WARP_REDUCE_HPP__ #define OPENCV_GPU_WARP_REDUCE_HPP__ -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template __device__ __forceinline__ T warp_reduce(volatile T *ptr , const unsigned int tid = threadIdx.x) @@ -63,6 +63,6 @@ namespace cv { namespace gpu { namespace cuda return ptr[tid - lane]; } -}}} // namespace cv { namespace gpu { namespace cuda { +}}} // namespace cv { namespace gpu { namespace cudev { #endif /* OPENCV_GPU_WARP_REDUCE_HPP__ */ diff --git a/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp b/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp index 954b618d3d..aabcacfa4b 100644 --- a/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp +++ b/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp @@ -43,7 +43,7 @@ #ifndef __OPENCV_GPU_WARP_SHUFFLE_HPP__ #define __OPENCV_GPU_WARP_SHUFFLE_HPP__ -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template __device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize) diff --git a/modules/core/src/cuda/matrix_operations.cu b/modules/core/src/cuda/matrix_operations.cu index 725321286b..e0b4ebc422 100644 --- a/modules/core/src/cuda/matrix_operations.cu +++ b/modules/core/src/cuda/matrix_operations.cu @@ -45,7 +45,7 @@ #include "opencv2/core/cuda/functional.hpp" #include "opencv2/core/cuda/type_traits.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { void writeScalar(const uchar*); void writeScalar(const schar*); @@ -58,7 +58,7 @@ namespace cv { namespace gpu { namespace cuda void convert_gpu(PtrStepSzb, int, PtrStepSzb, int, double, double, cudaStream_t); }}} -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct shift_and_sizeof; template <> struct shift_and_sizeof { enum { shift = 0 }; }; @@ -76,9 +76,9 @@ namespace cv { namespace gpu { namespace cuda template void copyToWithMask(PtrStepSzb src, PtrStepSzb dst, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream) { if (colorMask) - cv::gpu::cuda::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMask(mask), stream); + cv::gpu::cudev::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMask(mask), stream); else - cv::gpu::cuda::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMaskChannels(mask, cn), stream); + cv::gpu::cudev::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMaskChannels(mask, cn), stream); } void copyToWithMask_gpu(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream) @@ -293,7 +293,7 @@ namespace cv { namespace gpu { namespace cuda cvCudaSafeCall( cudaSetDoubleForDevice(&alpha) ); cvCudaSafeCall( cudaSetDoubleForDevice(&beta) ); Convertor op(static_cast(alpha), static_cast(beta)); - cv::gpu::cuda::transform((PtrStepSz)src, (PtrStepSz)dst, op, WithOutMask(), stream); + cv::gpu::cudev::transform((PtrStepSz)src, (PtrStepSz)dst, op, WithOutMask(), stream); } #if defined __clang__ @@ -379,4 +379,4 @@ namespace cv { namespace gpu { namespace cuda #if defined __clang__ # pragma clang diagnostic pop #endif -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev diff --git a/modules/core/src/gpumat.cpp b/modules/core/src/gpumat.cpp index bd4e38b077..384adef732 100644 --- a/modules/core/src/gpumat.cpp +++ b/modules/core/src/gpumat.cpp @@ -844,7 +844,7 @@ namespace #else // HAVE_CUDA -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { void copyToWithMask_gpu(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream); @@ -862,13 +862,13 @@ namespace template void kernelSetCaller(GpuMat& src, Scalar s, cudaStream_t stream) { Scalar_ sf = s; - cv::gpu::cuda::set_to_gpu(src, sf.val, src.channels(), stream); + cv::gpu::cudev::set_to_gpu(src, sf.val, src.channels(), stream); } template void kernelSetCaller(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream) { Scalar_ sf = s; - cv::gpu::cuda::set_to_gpu(src, sf.val, mask, src.channels(), stream); + cv::gpu::cudev::set_to_gpu(src, sf.val, mask, src.channels(), stream); } } @@ -892,17 +892,17 @@ namespace cv { namespace gpu CV_Assert(src.size() == dst.size() && src.type() == dst.type()); CV_Assert(src.size() == mask.size() && mask.depth() == CV_8U && (mask.channels() == 1 || mask.channels() == src.channels())); - cv::gpu::cuda::copyToWithMask_gpu(src.reshape(1), dst.reshape(1), src.elemSize1(), src.channels(), mask.reshape(1), mask.channels() != 1, stream); + cv::gpu::cudev::copyToWithMask_gpu(src.reshape(1), dst.reshape(1), src.elemSize1(), src.channels(), mask.reshape(1), mask.channels() != 1, stream); } void convertTo(const GpuMat& src, GpuMat& dst) { - cv::gpu::cuda::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), 1.0, 0.0, 0); + cv::gpu::cudev::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), 1.0, 0.0, 0); } void convertTo(const GpuMat& src, GpuMat& dst, double alpha, double beta, cudaStream_t stream = 0) { - cv::gpu::cuda::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), alpha, beta, stream); + cv::gpu::cudev::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), alpha, beta, stream); } void setTo(GpuMat& src, Scalar s, cudaStream_t stream) diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index 8f9113991c..5f32cdf1a8 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -444,7 +444,7 @@ void cv::gpu::magnitudeSqr(const GpuMat& src, GpuMat& dst, Stream& stream) //////////////////////////////////////////////////////////////////////// // Polar <-> Cart -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace mathfunc { @@ -457,7 +457,7 @@ namespace { inline void cartToPolar_caller(const GpuMat& x, const GpuMat& y, GpuMat* mag, bool magSqr, GpuMat* angle, bool angleInDegrees, cudaStream_t stream) { - using namespace ::cv::gpu::cuda::mathfunc; + using namespace ::cv::gpu::cudev::mathfunc; CV_Assert(x.size() == y.size() && x.type() == y.type()); CV_Assert(x.depth() == CV_32F); @@ -477,7 +477,7 @@ namespace inline void polarToCart_caller(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t stream) { - using namespace ::cv::gpu::cuda::mathfunc; + using namespace ::cv::gpu::cudev::mathfunc; CV_Assert((mag.empty() || mag.size() == angle.size()) && mag.type() == angle.type()); CV_Assert(mag.depth() == CV_32F); diff --git a/modules/gpu/src/bgfg_gmg.cpp b/modules/gpu/src/bgfg_gmg.cpp index 79777fcb7e..f29bf45519 100644 --- a/modules/gpu/src/bgfg_gmg.cpp +++ b/modules/gpu/src/bgfg_gmg.cpp @@ -51,7 +51,7 @@ void cv::gpu::GMG_GPU::release() {} #else -namespace cv { namespace gpu { namespace cuda { +namespace cv { namespace gpu { namespace cudev { namespace bgfg_gmg { void loadConstants(int width, int height, float minVal, float maxVal, int quantizationLevels, float backgroundPrior, @@ -77,7 +77,7 @@ cv::gpu::GMG_GPU::GMG_GPU() void cv::gpu::GMG_GPU::initialize(cv::Size frameSize, float min, float max) { - using namespace cv::gpu::cuda::bgfg_gmg; + using namespace cv::gpu::cudev::bgfg_gmg; CV_Assert(min < max); CV_Assert(maxFeatures > 0); @@ -107,7 +107,7 @@ void cv::gpu::GMG_GPU::initialize(cv::Size frameSize, float min, float max) void cv::gpu::GMG_GPU::operator ()(const cv::gpu::GpuMat& frame, cv::gpu::GpuMat& fgmask, float newLearningRate, cv::gpu::Stream& stream) { - using namespace cv::gpu::cuda::bgfg_gmg; + using namespace cv::gpu::cudev::bgfg_gmg; typedef void (*func_t)(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/bgfg_mog.cpp b/modules/gpu/src/bgfg_mog.cpp index 78cd9a4a0b..22b31074fe 100644 --- a/modules/gpu/src/bgfg_mog.cpp +++ b/modules/gpu/src/bgfg_mog.cpp @@ -58,7 +58,7 @@ void cv::gpu::MOG2_GPU::release() {} #else -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace mog { @@ -123,7 +123,7 @@ void cv::gpu::MOG_GPU::initialize(cv::Size frameSize, int frameType) void cv::gpu::MOG_GPU::operator()(const cv::gpu::GpuMat& frame, cv::gpu::GpuMat& fgmask, float learningRate, Stream& stream) { - using namespace cv::gpu::cuda::mog; + using namespace cv::gpu::cudev::mog; CV_Assert(frame.depth() == CV_8U); @@ -146,7 +146,7 @@ void cv::gpu::MOG_GPU::operator()(const cv::gpu::GpuMat& frame, cv::gpu::GpuMat& void cv::gpu::MOG_GPU::getBackgroundImage(GpuMat& backgroundImage, Stream& stream) const { - using namespace cv::gpu::cuda::mog; + using namespace cv::gpu::cudev::mog; backgroundImage.create(frameSize_, frameType_); @@ -208,7 +208,7 @@ cv::gpu::MOG2_GPU::MOG2_GPU(int nmixtures) : void cv::gpu::MOG2_GPU::initialize(cv::Size frameSize, int frameType) { - using namespace cv::gpu::cuda::mog; + using namespace cv::gpu::cudev::mog; CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3 || frameType == CV_8UC4); @@ -236,7 +236,7 @@ void cv::gpu::MOG2_GPU::initialize(cv::Size frameSize, int frameType) void cv::gpu::MOG2_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, float learningRate, Stream& stream) { - using namespace cv::gpu::cuda::mog; + using namespace cv::gpu::cudev::mog; int ch = frame.channels(); int work_ch = ch; @@ -256,7 +256,7 @@ void cv::gpu::MOG2_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, float le void cv::gpu::MOG2_GPU::getBackgroundImage(GpuMat& backgroundImage, Stream& stream) const { - using namespace cv::gpu::cuda::mog; + using namespace cv::gpu::cudev::mog; backgroundImage.create(frameSize_, frameType_); diff --git a/modules/gpu/src/bilateral_filter.cpp b/modules/gpu/src/bilateral_filter.cpp index 41b6275f2d..ef5be018da 100644 --- a/modules/gpu/src/bilateral_filter.cpp +++ b/modules/gpu/src/bilateral_filter.cpp @@ -54,7 +54,7 @@ void cv::gpu::DisparityBilateralFilter::operator()(const GpuMat&, const GpuMat&, #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace disp_bilateral_filter { @@ -65,7 +65,7 @@ namespace cv { namespace gpu { namespace cuda } }}} -using namespace ::cv::gpu::cuda::disp_bilateral_filter; +using namespace ::cv::gpu::cudev::disp_bilateral_filter; namespace { diff --git a/modules/gpu/src/blend.cpp b/modules/gpu/src/blend.cpp index 97469f6b8d..c667214ae7 100644 --- a/modules/gpu/src/blend.cpp +++ b/modules/gpu/src/blend.cpp @@ -51,7 +51,7 @@ void cv::gpu::blendLinear(const GpuMat&, const GpuMat&, const GpuMat&, const Gpu #else -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace blend { @@ -62,7 +62,7 @@ namespace cv { namespace gpu { namespace cuda } }}} -using namespace ::cv::gpu::cuda::blend; +using namespace ::cv::gpu::cudev::blend; void cv::gpu::blendLinear(const GpuMat& img1, const GpuMat& img2, const GpuMat& weights1, const GpuMat& weights2, GpuMat& result, Stream& stream) diff --git a/modules/gpu/src/brute_force_matcher.cpp b/modules/gpu/src/brute_force_matcher.cpp index 8fd4e11fd1..e350d48cfa 100644 --- a/modules/gpu/src/brute_force_matcher.cpp +++ b/modules/gpu/src/brute_force_matcher.cpp @@ -81,7 +81,7 @@ void cv::gpu::BFMatcher_GPU::radiusMatch(const GpuMat&, std::vector< std::vector #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace bf_match { @@ -197,7 +197,7 @@ void cv::gpu::BFMatcher_GPU::matchSingle(const GpuMat& query, const GpuMat& trai if (query.empty() || train.empty()) return; - using namespace cv::gpu::cuda::bf_match; + using namespace cv::gpu::cudev::bf_match; typedef void (*caller_t)(const PtrStepSzb& query, const PtrStepSzb& train, const PtrStepSzb& mask, const PtrStepSzi& trainIdx, const PtrStepSzf& distance, @@ -340,7 +340,7 @@ void cv::gpu::BFMatcher_GPU::matchCollection(const GpuMat& query, const GpuMat& if (query.empty() || trainCollection.empty()) return; - using namespace cv::gpu::cuda::bf_match; + using namespace cv::gpu::cudev::bf_match; typedef void (*caller_t)(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz& masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, @@ -451,7 +451,7 @@ void cv::gpu::BFMatcher_GPU::knnMatchSingle(const GpuMat& query, const GpuMat& t if (query.empty() || train.empty()) return; - using namespace cv::gpu::cuda::bf_knnmatch; + using namespace cv::gpu::cudev::bf_knnmatch; typedef void (*caller_t)(const PtrStepSzb& query, const PtrStepSzb& train, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, @@ -580,7 +580,7 @@ void cv::gpu::BFMatcher_GPU::knnMatch2Collection(const GpuMat& query, const GpuM if (query.empty() || trainCollection.empty()) return; - using namespace cv::gpu::cuda::bf_knnmatch; + using namespace cv::gpu::cudev::bf_knnmatch; typedef void (*caller_t)(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, @@ -761,7 +761,7 @@ void cv::gpu::BFMatcher_GPU::radiusMatchSingle(const GpuMat& query, const GpuMat if (query.empty() || train.empty()) return; - using namespace cv::gpu::cuda::bf_radius_match; + using namespace cv::gpu::cudev::bf_radius_match; typedef void (*caller_t)(const PtrStepSzb& query, const PtrStepSzb& train, float maxDistance, const PtrStepSzb& mask, const PtrStepSzi& trainIdx, const PtrStepSzf& distance, const PtrStepSz& nMatches, @@ -890,7 +890,7 @@ void cv::gpu::BFMatcher_GPU::radiusMatchCollection(const GpuMat& query, GpuMat& if (query.empty() || empty()) return; - using namespace cv::gpu::cuda::bf_radius_match; + using namespace cv::gpu::cudev::bf_radius_match; typedef void (*caller_t)(const PtrStepSzb& query, const PtrStepSzb* trains, int n, float maxDistance, const PtrStepSzb* masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, const PtrStepSz& nMatches, diff --git a/modules/gpu/src/calib3d.cpp b/modules/gpu/src/calib3d.cpp index 8c8f753b6a..abcc3423d8 100644 --- a/modules/gpu/src/calib3d.cpp +++ b/modules/gpu/src/calib3d.cpp @@ -55,7 +55,7 @@ void cv::gpu::solvePnPRansac(const Mat&, const Mat&, const Mat&, const Mat&, Mat #else -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace transform_points { @@ -78,7 +78,7 @@ namespace cv { namespace gpu { namespace cuda } }}} -using namespace ::cv::gpu::cuda; +using namespace ::cv::gpu::cudev; namespace { diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp index 4ccc86a694..2a65ab520f 100644 --- a/modules/gpu/src/cascadeclassifier.cpp +++ b/modules/gpu/src/cascadeclassifier.cpp @@ -340,7 +340,7 @@ struct PyrLavel cv::Size sWindow; }; -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace lbp { @@ -441,7 +441,7 @@ public: acc += level.sFrame.width + 1; } - cuda::lbp::classifyPyramid(image.cols, image.rows, NxM.width - 1, NxM.height - 1, iniScale, scaleFactor, total, stage_mat, stage_mat.cols / sizeof(Stage), nodes_mat, + cudev::lbp::classifyPyramid(image.cols, image.rows, NxM.width - 1, NxM.height - 1, iniScale, scaleFactor, total, stage_mat, stage_mat.cols / sizeof(Stage), nodes_mat, leaves_mat, subsets_mat, features_mat, subsetSize, candidates, dclassified.ptr(), integral); } @@ -449,7 +449,7 @@ public: return 0; cvCudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); - cuda::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified.ptr()); + cudev::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified.ptr()); cvCudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); cvCudaSafeCall( cudaDeviceSynchronize() ); diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp index 33d57360ce..570efbcb87 100644 --- a/modules/gpu/src/color.cpp +++ b/modules/gpu/src/color.cpp @@ -57,7 +57,7 @@ void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_no_ #include "cvt_color_internal.h" namespace cv { namespace gpu { - namespace cuda + namespace cudev { template void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); @@ -69,7 +69,7 @@ namespace cv { namespace gpu { } }} -using namespace ::cv::gpu::cuda; +using namespace ::cv::gpu::cudev; namespace { @@ -77,7 +77,7 @@ namespace void bgr_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {bgr_to_rgb_8u, 0, bgr_to_rgb_16u, 0, 0, bgr_to_rgb_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -90,7 +90,7 @@ namespace void bgr_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {bgr_to_bgra_8u, 0, bgr_to_bgra_16u, 0, 0, bgr_to_bgra_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -103,7 +103,7 @@ namespace void bgr_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {bgr_to_rgba_8u, 0, bgr_to_rgba_16u, 0, 0, bgr_to_rgba_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -116,7 +116,7 @@ namespace void bgra_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {bgra_to_bgr_8u, 0, bgra_to_bgr_16u, 0, 0, bgra_to_bgr_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -129,7 +129,7 @@ namespace void bgra_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {bgra_to_rgb_8u, 0, bgra_to_rgb_16u, 0, 0, bgra_to_rgb_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -142,7 +142,7 @@ namespace void bgra_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {bgra_to_rgba_8u, 0, bgra_to_rgba_16u, 0, 0, bgra_to_rgba_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -160,7 +160,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void bgr_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -170,7 +170,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void rgb_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -180,7 +180,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cudev::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void rgb_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -190,7 +190,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cudev::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void bgra_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -200,7 +200,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cudev::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void bgra_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -210,7 +210,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cudev::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void rgba_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -220,7 +220,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cudev::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void rgba_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -230,7 +230,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cudev::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -240,7 +240,7 @@ namespace dst.create(src.size(), CV_8UC3); - cuda::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -250,7 +250,7 @@ namespace dst.create(src.size(), CV_8UC3); - cuda::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -260,7 +260,7 @@ namespace dst.create(src.size(), CV_8UC3); - cuda::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -270,7 +270,7 @@ namespace dst.create(src.size(), CV_8UC3); - cuda::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -280,7 +280,7 @@ namespace dst.create(src.size(), CV_8UC4); - cuda::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -290,7 +290,7 @@ namespace dst.create(src.size(), CV_8UC4); - cuda::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -300,7 +300,7 @@ namespace dst.create(src.size(), CV_8UC4); - cuda::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -310,12 +310,12 @@ namespace dst.create(src.size(), CV_8UC4); - cuda::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream)); } void gray_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {gray_to_bgr_8u, 0, gray_to_bgr_16u, 0, 0, gray_to_bgr_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -328,7 +328,7 @@ namespace void gray_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {gray_to_bgra_8u, 0, gray_to_bgra_16u, 0, 0, gray_to_bgra_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -346,7 +346,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cudev::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void gray_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -356,7 +356,7 @@ namespace dst.create(src.size(), CV_8UC2); - cuda::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cudev::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -366,7 +366,7 @@ namespace dst.create(src.size(), CV_8UC1); - cuda::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) @@ -376,12 +376,12 @@ namespace dst.create(src.size(), CV_8UC1); - cuda::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream)); + cudev::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream)); } void rgb_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {rgb_to_gray_8u, 0, rgb_to_gray_16u, 0, 0, rgb_to_gray_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -394,7 +394,7 @@ namespace void bgr_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {bgr_to_gray_8u, 0, bgr_to_gray_16u, 0, 0, bgr_to_gray_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -407,7 +407,7 @@ namespace void rgba_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {rgba_to_gray_8u, 0, rgba_to_gray_16u, 0, 0, rgba_to_gray_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -420,7 +420,7 @@ namespace void bgra_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[] = {bgra_to_gray_8u, 0, bgra_to_gray_16u, 0, 0, bgra_to_gray_32f}; CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); @@ -433,7 +433,7 @@ namespace void rgb_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -459,7 +459,7 @@ namespace void bgr_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -485,7 +485,7 @@ namespace void yuv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -511,7 +511,7 @@ namespace void yuv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -537,7 +537,7 @@ namespace void rgb_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -563,7 +563,7 @@ namespace void bgr_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -589,7 +589,7 @@ namespace void YCrCb_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -615,7 +615,7 @@ namespace void YCrCb_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -641,7 +641,7 @@ namespace void rgb_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -667,7 +667,7 @@ namespace void bgr_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -693,7 +693,7 @@ namespace void xyz_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -719,7 +719,7 @@ namespace void xyz_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -745,7 +745,7 @@ namespace void rgb_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -771,7 +771,7 @@ namespace void bgr_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -797,7 +797,7 @@ namespace void hsv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -823,7 +823,7 @@ namespace void hsv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -849,7 +849,7 @@ namespace void rgb_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -875,7 +875,7 @@ namespace void bgr_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -901,7 +901,7 @@ namespace void hls_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -927,7 +927,7 @@ namespace void hls_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -953,7 +953,7 @@ namespace void rgb_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -979,7 +979,7 @@ namespace void bgr_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -1005,7 +1005,7 @@ namespace void hsv_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -1031,7 +1031,7 @@ namespace void hsv_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -1057,7 +1057,7 @@ namespace void rgb_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -1083,7 +1083,7 @@ namespace void bgr_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -1109,7 +1109,7 @@ namespace void hls_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -1135,7 +1135,7 @@ namespace void hls_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][6] = { { @@ -1161,7 +1161,7 @@ namespace void bgr_to_lab(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1187,7 +1187,7 @@ namespace void rgb_to_lab(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1213,7 +1213,7 @@ namespace void lbgr_to_lab(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1239,7 +1239,7 @@ namespace void lrgb_to_lab(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1265,7 +1265,7 @@ namespace void lab_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1291,7 +1291,7 @@ namespace void lab_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1317,7 +1317,7 @@ namespace void lab_to_lbgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1343,7 +1343,7 @@ namespace void lab_to_lrgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1369,7 +1369,7 @@ namespace void bgr_to_luv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1395,7 +1395,7 @@ namespace void rgb_to_luv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1421,7 +1421,7 @@ namespace void lbgr_to_luv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1447,7 +1447,7 @@ namespace void lrgb_to_luv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1473,7 +1473,7 @@ namespace void luv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1499,7 +1499,7 @@ namespace void luv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1525,7 +1525,7 @@ namespace void luv_to_lbgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1551,7 +1551,7 @@ namespace void luv_to_lrgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - using namespace cv::gpu::cuda; + using namespace cv::gpu::cudev; static const gpu_func_t funcs[2][2][2] = { { @@ -1895,9 +1895,9 @@ void cv::gpu::demosaicing(const GpuMat& src, GpuMat& dst, int code, int dcn, Str code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1); if (dcn == 3) - cuda::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); + cudev::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); else - cuda::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); + cudev::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); break; } @@ -1917,7 +1917,7 @@ void cv::gpu::demosaicing(const GpuMat& src, GpuMat& dst, int code, int dcn, Str const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1, code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1); - cuda::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); + cudev::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); break; } diff --git a/modules/gpu/src/cuda/NV12ToARGB.cu b/modules/gpu/src/cuda/NV12ToARGB.cu index cccf6eb465..6c307c7fdf 100644 --- a/modules/gpu/src/cuda/NV12ToARGB.cu +++ b/modules/gpu/src/cuda/NV12ToARGB.cu @@ -51,7 +51,7 @@ #include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace cuda { +namespace cv { namespace gpu { namespace cudev { namespace video_decoding { __constant__ uint constAlpha = ((uint)0xff << 24); diff --git a/modules/gpu/src/cuda/bf_knnmatch.cu b/modules/gpu/src/cuda/bf_knnmatch.cu index f03b1c4992..5fe40e5bd8 100644 --- a/modules/gpu/src/cuda/bf_knnmatch.cu +++ b/modules/gpu/src/cuda/bf_knnmatch.cu @@ -50,7 +50,7 @@ #include "opencv2/core/cuda/datamov_utils.hpp" #include "opencv2/core/cuda/warp_shuffle.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace bf_knnmatch { @@ -1249,7 +1249,7 @@ namespace cv { namespace gpu { namespace cuda //template void match2Hamming_gpu(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, cudaStream_t stream); template void match2Hamming_gpu(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, cudaStream_t stream); } // namespace bf_knnmatch -}}} // namespace cv { namespace gpu { namespace cuda { +}}} // namespace cv { namespace gpu { namespace cudev { #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/bf_match.cu b/modules/gpu/src/cuda/bf_match.cu index 2ecb48ead0..b6820aa531 100644 --- a/modules/gpu/src/cuda/bf_match.cu +++ b/modules/gpu/src/cuda/bf_match.cu @@ -49,7 +49,7 @@ #include "opencv2/core/cuda/vec_distance.hpp" #include "opencv2/core/cuda/datamov_utils.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace bf_match { @@ -768,7 +768,7 @@ namespace cv { namespace gpu { namespace cuda //template void matchHamming_gpu(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz& masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, cudaStream_t stream); template void matchHamming_gpu(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz& masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, cudaStream_t stream); } // namespace bf_match -}}} // namespace cv { namespace gpu { namespace cuda { +}}} // namespace cv { namespace gpu { namespace cudev { #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/bf_radius_match.cu b/modules/gpu/src/cuda/bf_radius_match.cu index 56b99faf5b..cb44768efb 100644 --- a/modules/gpu/src/cuda/bf_radius_match.cu +++ b/modules/gpu/src/cuda/bf_radius_match.cu @@ -48,7 +48,7 @@ #include "opencv2/core/cuda/vec_distance.hpp" #include "opencv2/core/cuda/datamov_utils.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace bf_radius_match { @@ -457,7 +457,7 @@ namespace cv { namespace gpu { namespace cuda //template void matchHamming_gpu(const PtrStepSzb& query, const PtrStepSzb* trains, int n, float maxDistance, const PtrStepSzb* masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, const PtrStepSz& nMatches, cudaStream_t stream); template void matchHamming_gpu(const PtrStepSzb& query, const PtrStepSzb* trains, int n, float maxDistance, const PtrStepSzb* masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, const PtrStepSz& nMatches, cudaStream_t stream); } // namespace bf_radius_match -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/bgfg_gmg.cu b/modules/gpu/src/cuda/bgfg_gmg.cu index 14e23df947..0e4421374f 100644 --- a/modules/gpu/src/cuda/bgfg_gmg.cu +++ b/modules/gpu/src/cuda/bgfg_gmg.cu @@ -46,7 +46,7 @@ #include "opencv2/core/cuda/vec_traits.hpp" #include "opencv2/core/cuda/limits.hpp" -namespace cv { namespace gpu { namespace cuda { +namespace cv { namespace gpu { namespace cudev { namespace bgfg_gmg { __constant__ int c_width; diff --git a/modules/gpu/src/cuda/bgfg_mog.cu b/modules/gpu/src/cuda/bgfg_mog.cu index 771f367cec..1078343469 100644 --- a/modules/gpu/src/cuda/bgfg_mog.cu +++ b/modules/gpu/src/cuda/bgfg_mog.cu @@ -47,7 +47,7 @@ #include "opencv2/core/cuda/vec_math.hpp" #include "opencv2/core/cuda/limits.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace mog { diff --git a/modules/gpu/src/cuda/bilateral_filter.cu b/modules/gpu/src/cuda/bilateral_filter.cu index 8a6f781118..8d87116180 100644 --- a/modules/gpu/src/cuda/bilateral_filter.cu +++ b/modules/gpu/src/cuda/bilateral_filter.cu @@ -55,7 +55,7 @@ typedef unsigned short ushort; ////////////////////////////////////////////////////////////////////////////////// /// Bilateral filtering -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace imgproc { @@ -163,7 +163,7 @@ namespace cv { namespace gpu { namespace cuda #define OCV_INSTANTIATE_BILATERAL_FILTER(T) \ - template void cv::gpu::cuda::imgproc::bilateral_filter_gpu(const PtrStepSzb&, PtrStepSzb, int, float, float, int, cudaStream_t); + template void cv::gpu::cudev::imgproc::bilateral_filter_gpu(const PtrStepSzb&, PtrStepSzb, int, float, float, int, cudaStream_t); OCV_INSTANTIATE_BILATERAL_FILTER(uchar) //OCV_INSTANTIATE_BILATERAL_FILTER(uchar2) diff --git a/modules/gpu/src/cuda/blend.cu b/modules/gpu/src/cuda/blend.cu index f5f734cc5c..d8054d6b5f 100644 --- a/modules/gpu/src/cuda/blend.cu +++ b/modules/gpu/src/cuda/blend.cu @@ -44,7 +44,7 @@ #include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace blend { @@ -115,7 +115,7 @@ namespace cv { namespace gpu { namespace cuda cvCudaSafeCall(cudaDeviceSynchronize()); } } // namespace blend -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/calib3d.cu b/modules/gpu/src/cuda/calib3d.cu index f8882cfc21..c85a2c7fe1 100644 --- a/modules/gpu/src/cuda/calib3d.cu +++ b/modules/gpu/src/cuda/calib3d.cu @@ -47,7 +47,7 @@ #include "opencv2/core/cuda/functional.hpp" #include "opencv2/core/cuda/reduce.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { #define SOLVE_PNP_RANSAC_MAX_NUM_ITERS 200 @@ -79,7 +79,7 @@ namespace cv { namespace gpu { namespace cuda cvCudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3)); cvCudaSafeCall(cudaMemcpyToSymbol(crot2, rot + 6, sizeof(float) * 3)); cvCudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3)); - cv::gpu::cuda::transform(src, dst, TransformOp(), WithOutMask(), stream); + cv::gpu::cudev::transform(src, dst, TransformOp(), WithOutMask(), stream); } } // namespace transform_points @@ -120,7 +120,7 @@ namespace cv { namespace gpu { namespace cuda cvCudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3)); cvCudaSafeCall(cudaMemcpyToSymbol(cproj0, proj, sizeof(float) * 3)); cvCudaSafeCall(cudaMemcpyToSymbol(cproj1, proj + 3, sizeof(float) * 3)); - cv::gpu::cuda::transform(src, dst, ProjectOp(), WithOutMask(), stream); + cv::gpu::cudev::transform(src, dst, ProjectOp(), WithOutMask(), stream); } } // namespace project_points @@ -187,7 +187,7 @@ namespace cv { namespace gpu { namespace cuda cvCudaSafeCall( cudaDeviceSynchronize() ); } } // namespace solvepnp_ransac -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/canny.cu b/modules/gpu/src/cuda/canny.cu index 62c7242ed2..cbd792fd0b 100644 --- a/modules/gpu/src/cuda/canny.cu +++ b/modules/gpu/src/cuda/canny.cu @@ -51,7 +51,7 @@ #include "opencv2/core/cuda/utility.hpp" using namespace cv::gpu; -using namespace cv::gpu::cuda; +using namespace cv::gpu::cudev; namespace canny { @@ -77,7 +77,7 @@ namespace canny }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template <> struct TransformFunctorTraits : DefaultTransformFunctorTraits { @@ -475,7 +475,7 @@ namespace canny }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template <> struct TransformFunctorTraits : DefaultTransformFunctorTraits { diff --git a/modules/gpu/src/cuda/ccomponetns.cu b/modules/gpu/src/cuda/ccomponetns.cu index 45928a0841..23dcfe9ba2 100644 --- a/modules/gpu/src/cuda/ccomponetns.cu +++ b/modules/gpu/src/cuda/ccomponetns.cu @@ -50,7 +50,7 @@ #include #include -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace ccl { diff --git a/modules/gpu/src/cuda/clahe.cu b/modules/gpu/src/cuda/clahe.cu index 340e15b450..16afd1d5a7 100644 --- a/modules/gpu/src/cuda/clahe.cu +++ b/modules/gpu/src/cuda/clahe.cu @@ -50,7 +50,7 @@ #include "opencv2/core/cuda/saturate_cast.hpp" using namespace cv::gpu; -using namespace cv::gpu::cuda; +using namespace cv::gpu::cudev; namespace clahe { diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu index a388735187..1a5d4865ed 100644 --- a/modules/gpu/src/cuda/color.cu +++ b/modules/gpu/src/cuda/color.cu @@ -47,7 +47,7 @@ #include "opencv2/core/cuda/color.hpp" #include "cvt_color_internal.h" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { OPENCV_GPU_TRANSFORM_FUNCTOR_TRAITS(bgra_to_rgba_traits::functor_type) { @@ -229,7 +229,7 @@ namespace cv { namespace gpu { namespace cuda traits::functor_type functor = traits::create_functor(); \ typedef typename traits::functor_type::argument_type src_t; \ typedef typename traits::functor_type::result_type dst_t; \ - cv::gpu::cuda::transform((PtrStepSz)src, (PtrStepSz)dst, functor, WithOutMask(), stream); \ + cv::gpu::cudev::transform((PtrStepSz)src, (PtrStepSz)dst, functor, WithOutMask(), stream); \ } #define OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(name) \ @@ -456,6 +456,6 @@ namespace cv { namespace gpu { namespace cuda #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F_FULL -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/column_filter.h b/modules/gpu/src/cuda/column_filter.h index 1f42767407..ecb6da16aa 100644 --- a/modules/gpu/src/cuda/column_filter.h +++ b/modules/gpu/src/cuda/column_filter.h @@ -46,7 +46,7 @@ #include "opencv2/core/cuda/border_interpolate.hpp" using namespace cv::gpu; -using namespace cv::gpu::cuda; +using namespace cv::gpu::cudev; namespace column_filter { diff --git a/modules/gpu/src/cuda/copy_make_border.cu b/modules/gpu/src/cuda/copy_make_border.cu index 30bbe0f40a..a9ad4f9888 100644 --- a/modules/gpu/src/cuda/copy_make_border.cu +++ b/modules/gpu/src/cuda/copy_make_border.cu @@ -45,7 +45,7 @@ #include "opencv2/core/cuda/common.hpp" #include "opencv2/core/cuda/border_interpolate.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace imgproc { @@ -126,6 +126,6 @@ namespace cv { namespace gpu { namespace cuda template void copyMakeBorder_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderMode, const float* borderValue, cudaStream_t stream); template void copyMakeBorder_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderMode, const float* borderValue, cudaStream_t stream); } // namespace imgproc -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/debayer.cu b/modules/gpu/src/cuda/debayer.cu index 6ade873f3e..6c258a83aa 100644 --- a/modules/gpu/src/cuda/debayer.cu +++ b/modules/gpu/src/cuda/debayer.cu @@ -49,7 +49,7 @@ #include "opencv2/core/cuda/color.hpp" #include "opencv2/core/cuda/saturate_cast.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct Bayer2BGR; diff --git a/modules/gpu/src/cuda/disp_bilateral_filter.cu b/modules/gpu/src/cuda/disp_bilateral_filter.cu index 7758e4ce01..58dc4e80b8 100644 --- a/modules/gpu/src/cuda/disp_bilateral_filter.cu +++ b/modules/gpu/src/cuda/disp_bilateral_filter.cu @@ -45,7 +45,7 @@ #include "opencv2/core/cuda/common.hpp" #include "opencv2/core/cuda/limits.hpp" -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { namespace disp_bilateral_filter { @@ -218,6 +218,6 @@ namespace cv { namespace gpu { namespace cuda template void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); template void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); } // namespace bilateral_filter -}}} // namespace cv { namespace gpu { namespace cuda +}}} // namespace cv { namespace gpu { namespace cudev #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu index edd17aea79..095d8bac06 100644 --- a/modules/gpu/src/cuda/element_operations.cu +++ b/modules/gpu/src/cuda/element_operations.cu @@ -51,7 +51,7 @@ #include "opencv2/core/cuda/simd_functions.hpp" using namespace cv::gpu; -using namespace cv::gpu::cuda; +using namespace cv::gpu::cudev; namespace arithm { @@ -193,7 +193,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template <> struct TransformFunctorTraits< arithm::VAdd4 > : arithm::ArithmFuncTraits { @@ -216,21 +216,21 @@ namespace arithm { void addMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VAdd4(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VAdd4(), WithOutMask(), stream); } void addMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VAdd2(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VAdd2(), WithOutMask(), stream); } template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) { if (mask.data) - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), mask, stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), mask, stream); else - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), WithOutMask(), stream); } template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); @@ -308,7 +308,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< arithm::AddScalar > : arithm::ArithmFuncTraits { @@ -323,9 +323,9 @@ namespace arithm AddScalar op(static_cast(val)); if (mask.data) - cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); else - cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); @@ -428,7 +428,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template <> struct TransformFunctorTraits< arithm::VSub4 > : arithm::ArithmFuncTraits { @@ -451,21 +451,21 @@ namespace arithm { void subMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VSub4(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VSub4(), WithOutMask(), stream); } void subMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VSub2(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VSub2(), WithOutMask(), stream); } template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) { if (mask.data) - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), mask, stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), mask, stream); else - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), WithOutMask(), stream); } template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); @@ -536,9 +536,9 @@ namespace arithm AddScalar op(-static_cast(val)); if (mask.data) - cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); else - cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void subScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); @@ -657,7 +657,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template <> struct TransformFunctorTraits : arithm::ArithmFuncTraits { @@ -676,12 +676,12 @@ namespace arithm { void mulMat_8uc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, Mul_8uc4_32f(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, Mul_8uc4_32f(), WithOutMask(), stream); } void mulMat_16sc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, Mul_16sc4_32f(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, Mul_16sc4_32f(), WithOutMask(), stream); } template @@ -690,12 +690,12 @@ namespace arithm if (scale == 1) { Mul op; - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); } else { MulScale op(static_cast(scale)); - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); } } @@ -774,7 +774,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< arithm::MulScalar > : arithm::ArithmFuncTraits { @@ -787,7 +787,7 @@ namespace arithm void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream) { MulScalar op(static_cast(val)); - cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); @@ -925,7 +925,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template <> struct TransformFunctorTraits : arithm::ArithmFuncTraits { @@ -944,12 +944,12 @@ namespace arithm { void divMat_8uc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, Div_8uc4_32f(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, Div_8uc4_32f(), WithOutMask(), stream); } void divMat_16sc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, Div_16sc4_32f(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, Div_16sc4_32f(), WithOutMask(), stream); } template @@ -958,12 +958,12 @@ namespace arithm if (scale == 1) { Div op; - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); } else { DivScale op(static_cast(scale)); - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); } } @@ -1033,7 +1033,7 @@ namespace arithm void divScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream) { MulScalar op(static_cast(1.0 / val)); - cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void divScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); @@ -1111,7 +1111,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< arithm::DivInv > : arithm::ArithmFuncTraits { @@ -1124,7 +1124,7 @@ namespace arithm void divInv(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream) { DivInv op(static_cast(val)); - cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void divInv(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); @@ -1240,7 +1240,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template <> struct TransformFunctorTraits< arithm::VAbsDiff4 > : arithm::ArithmFuncTraits { @@ -1263,18 +1263,18 @@ namespace arithm { void absDiffMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VAbsDiff4(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VAbsDiff4(), WithOutMask(), stream); } void absDiffMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VAbsDiff2(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VAbsDiff2(), WithOutMask(), stream); } template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream) { - cuda::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AbsDiffMat(), WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AbsDiffMat(), WithOutMask(), stream); } template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); @@ -1305,7 +1305,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< arithm::AbsDiffScalar > : arithm::ArithmFuncTraits { @@ -1319,7 +1319,7 @@ namespace arithm { AbsDiffScalar op(static_cast(val)); - cuda::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + cudev::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); } template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); @@ -1334,7 +1334,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////// // absMat -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< abs_func > : arithm::ArithmFuncTraits { @@ -1346,7 +1346,7 @@ namespace arithm template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - cuda::transform((PtrStepSz) src, (PtrStepSz) dst, abs_func(), WithOutMask(), stream); + cudev::transform((PtrStepSz) src, (PtrStepSz) dst, abs_func(), WithOutMask(), stream); } template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1375,7 +1375,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< arithm::Sqr > : arithm::ArithmFuncTraits { @@ -1387,7 +1387,7 @@ namespace arithm template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - cuda::transform((PtrStepSz) src, (PtrStepSz) dst, Sqr(), WithOutMask(), stream); + cudev::transform((PtrStepSz) src, (PtrStepSz) dst, Sqr(), WithOutMask(), stream); } template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1402,7 +1402,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////// // sqrtMat -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< sqrt_func > : arithm::ArithmFuncTraits { @@ -1414,7 +1414,7 @@ namespace arithm template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - cuda::transform((PtrStepSz) src, (PtrStepSz) dst, sqrt_func(), WithOutMask(), stream); + cudev::transform((PtrStepSz) src, (PtrStepSz) dst, sqrt_func(), WithOutMask(), stream); } template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1429,7 +1429,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////// // logMat -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< log_func > : arithm::ArithmFuncTraits { @@ -1441,7 +1441,7 @@ namespace arithm template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - cuda::transform((PtrStepSz) src, (PtrStepSz) dst, log_func(), WithOutMask(), stream); + cudev::transform((PtrStepSz) src, (PtrStepSz) dst, log_func(), WithOutMask(), stream); } template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1471,7 +1471,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template struct TransformFunctorTraits< arithm::Exp > : arithm::ArithmFuncTraits { @@ -1483,7 +1483,7 @@ namespace arithm template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) { - cuda::transform((PtrStepSz) src, (PtrStepSz) dst, Exp(), WithOutMask(), stream); + cudev::transform((PtrStepSz) src, (PtrStepSz) dst, Exp(), WithOutMask(), stream); } template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -1554,7 +1554,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace cuda +namespace cv { namespace gpu { namespace cudev { template <> struct TransformFunctorTraits< arithm::VCmpEq4 > : arithm::ArithmFuncTraits { @@ -1580,26 +1580,26 @@ namespace arithm { void cmpMatEq_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VCmpEq4(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VCmpEq4(), WithOutMask(), stream); } void cmpMatNe_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VCmpNe4(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VCmpNe4(), WithOutMask(), stream); } void cmpMatLt_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VCmpLt4(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VCmpLt4(), WithOutMask(), stream); } void cmpMatLe_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) { - cuda::transform(src1, src2, dst, VCmpLe4(), WithOutMask(), stream); + cudev::transform(src1, src2, dst, VCmpLe4(), WithOutMask(), stream); } template