From 28b1e81883653f59acf42fb268d8610d26707274 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Wed, 3 Apr 2013 16:04:04 +0400 Subject: [PATCH] renamed device -> cuda --- cmake/OpenCVModule.cmake | 16 +- modules/core/CMakeLists.txt | 6 +- .../opencv2/core/{device => cuda}/block.hpp | 2 +- .../core/{device => cuda}/border_interpolate.hpp | 4 +- .../opencv2/core/{device => cuda}/color.hpp | 4 +- .../opencv2/core/{device => cuda}/common.hpp | 2 +- .../core/{device => cuda}/datamov_utils.hpp | 4 +- .../core/{device => cuda}/detail/color_detail.hpp | 70 ++++----- .../core/{device => cuda}/detail/reduce.hpp | 2 +- .../{device => cuda}/detail/reduce_key_val.hpp | 2 +- .../{device => cuda}/detail/transform_detail.hpp | 4 +- .../{device => cuda}/detail/type_traits_detail.hpp | 4 +- .../detail/vec_distance_detail.hpp | 4 +- .../opencv2/core/{device => cuda}/dynamic_smem.hpp | 2 +- .../opencv2/core/{device => cuda}/emulation.hpp | 4 +- .../opencv2/core/{device => cuda}/filters.hpp | 4 +- .../opencv2/core/{device => cuda}/funcattrib.hpp | 4 +- .../opencv2/core/{device => cuda}/functional.hpp | 4 +- .../opencv2/core/{device => cuda}/limits.hpp | 4 +- .../opencv2/core/{device => cuda}/reduce.hpp | 2 +- .../core/{device => cuda}/saturate_cast.hpp | 2 +- .../include/opencv2/core/{device => cuda}/scan.hpp | 14 +- .../core/{device => cuda}/simd_functions.hpp | 2 +- .../opencv2/core/{device => cuda}/static_check.hpp | 2 +- .../opencv2/core/{device => cuda}/transform.hpp | 2 +- .../opencv2/core/{device => cuda}/type_traits.hpp | 2 +- .../opencv2/core/{device => cuda}/utility.hpp | 4 +- .../opencv2/core/{device => cuda}/vec_distance.hpp | 4 +- .../opencv2/core/{device => cuda}/vec_math.hpp | 4 +- .../opencv2/core/{device => cuda}/vec_traits.hpp | 4 +- .../include/opencv2/core/{device => cuda}/warp.hpp | 4 +- .../opencv2/core/{device => cuda}/warp_reduce.hpp | 4 +- .../opencv2/core/{device => cuda}/warp_shuffle.hpp | 2 +- modules/core/include/opencv2/core/cuda_devptrs.hpp | 18 --- modules/core/src/cuda/matrix_operations.cu | 20 +-- modules/core/src/gpumat.cpp | 12 +- modules/gpu/src/arithm.cpp | 6 +- modules/gpu/src/bgfg_gmg.cpp | 6 +- modules/gpu/src/bgfg_mog.cpp | 12 +- modules/gpu/src/bilateral_filter.cpp | 4 +- modules/gpu/src/blend.cpp | 4 +- modules/gpu/src/brute_force_matcher.cpp | 14 +- modules/gpu/src/calib3d.cpp | 4 +- modules/gpu/src/cascadeclassifier.cpp | 6 +- modules/gpu/src/color.cpp | 162 ++++++++++----------- modules/gpu/src/cuda/NV12ToARGB.cu | 4 +- modules/gpu/src/cuda/bf_knnmatch.cu | 20 +-- modules/gpu/src/cuda/bf_match.cu | 18 +-- modules/gpu/src/cuda/bf_radius_match.cu | 14 +- modules/gpu/src/cuda/bgfg_gmg.cu | 8 +- modules/gpu/src/cuda/bgfg_mog.cu | 10 +- modules/gpu/src/cuda/bilateral_filter.cu | 12 +- modules/gpu/src/cuda/blend.cu | 6 +- modules/gpu/src/cuda/calib3d.cu | 16 +- modules/gpu/src/cuda/canny.cu | 16 +- modules/gpu/src/cuda/ccomponetns.cu | 10 +- modules/gpu/src/cuda/clahe.cu | 14 +- modules/gpu/src/cuda/color.cu | 12 +- modules/gpu/src/cuda/column_filter.h | 10 +- modules/gpu/src/cuda/copy_make_border.cu | 8 +- modules/gpu/src/cuda/debayer.cu | 16 +- modules/gpu/src/cuda/disp_bilateral_filter.cu | 8 +- modules/gpu/src/cuda/element_operations.cu | 72 ++++----- modules/gpu/src/cuda/fast.cu | 6 +- modules/gpu/src/cuda/fgd_bgfg.cu | 14 +- modules/gpu/src/cuda/gftt.cu | 6 +- modules/gpu/src/cuda/global_motion.cu | 4 +- modules/gpu/src/cuda/hist.cu | 12 +- modules/gpu/src/cuda/hog.cu | 12 +- modules/gpu/src/cuda/hough.cu | 14 +- modules/gpu/src/cuda/imgproc.cu | 14 +- modules/gpu/src/cuda/integral_image.cu | 4 +- modules/gpu/src/cuda/internal_shared.hpp | 2 +- modules/gpu/src/cuda/lbp.cu | 6 +- modules/gpu/src/cuda/lbp.hpp | 6 +- modules/gpu/src/cuda/match_template.cu | 8 +- modules/gpu/src/cuda/mathfunc.cu | 6 +- modules/gpu/src/cuda/matrix_reductions.cu | 42 +++--- modules/gpu/src/cuda/nlm.cu | 24 +-- modules/gpu/src/cuda/optflowbm.cu | 10 +- modules/gpu/src/cuda/optical_flow.cu | 4 +- modules/gpu/src/cuda/optical_flow_farneback.cu | 8 +- modules/gpu/src/cuda/orb.cu | 8 +- modules/gpu/src/cuda/pyr_down.cu | 14 +- modules/gpu/src/cuda/pyr_up.cu | 14 +- modules/gpu/src/cuda/pyrlk.cu | 14 +- modules/gpu/src/cuda/remap.cu | 18 +-- modules/gpu/src/cuda/resize.cu | 20 +-- modules/gpu/src/cuda/rgb_to_yv12.cu | 6 +- modules/gpu/src/cuda/row_filter.h | 10 +- modules/gpu/src/cuda/split_merge.cu | 6 +- modules/gpu/src/cuda/stereobm.cu | 6 +- modules/gpu/src/cuda/stereobp.cu | 12 +- modules/gpu/src/cuda/stereocsbp.cu | 16 +- modules/gpu/src/cuda/texture_binder.hpp | 2 +- modules/gpu/src/cuda/tvl1flow.cu | 8 +- modules/gpu/src/cuda/warp.cu | 18 +-- modules/gpu/src/cvt_color_internal.h | 2 +- modules/gpu/src/denoising.cpp | 16 +- modules/gpu/src/fast.cpp | 6 +- modules/gpu/src/filtering.cpp | 6 +- modules/gpu/src/gftt.cpp | 4 +- modules/gpu/src/global_motion.cpp | 6 +- modules/gpu/src/graphcuts.cpp | 12 +- modules/gpu/src/hog.cpp | 4 +- modules/gpu/src/hough.cpp | 42 +++--- modules/gpu/src/imgproc.cpp | 68 ++++----- modules/gpu/src/match_template.cpp | 4 +- modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu | 6 +- modules/gpu/src/nvidia/NCVHaarObjectDetection.cu | 8 +- modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu | 8 +- modules/gpu/src/nvidia/core/NCVPyramid.cu | 6 +- modules/gpu/src/optical_flow.cpp | 4 +- modules/gpu/src/optical_flow_farneback.cpp | 38 ++--- modules/gpu/src/orb.cpp | 12 +- modules/gpu/src/pyramids.cpp | 14 +- modules/gpu/src/remap.cpp | 4 +- modules/gpu/src/resize.cpp | 4 +- modules/gpu/src/split_merge.cpp | 6 +- modules/gpu/src/stereobm.cpp | 4 +- modules/gpu/src/stereobp.cpp | 4 +- modules/gpu/src/stereocsbp.cpp | 4 +- modules/gpu/src/video_reader.cpp | 4 +- modules/gpu/src/video_writer.cpp | 4 +- modules/gpu/src/warp.cpp | 10 +- modules/nonfree/src/cuda/surf.cu | 24 +-- modules/nonfree/src/cuda/vibe.cu | 6 +- modules/nonfree/src/precomp.hpp | 2 +- modules/nonfree/src/surf_gpu.cpp | 4 +- modules/nonfree/src/vibe_gpu.cpp | 6 +- modules/softcascade/src/cuda/channels.cu | 2 +- modules/softcascade/src/cuda/icf-sc.cu | 2 +- modules/softcascade/src/cuda_invoker.hpp | 2 +- modules/softcascade/src/detector_cuda.cpp | 48 +++--- modules/superres/src/btv_l1_gpu.cpp | 16 +- modules/superres/src/cuda/btv_l1_gpu.cu | 20 +-- 136 files changed, 769 insertions(+), 787 deletions(-) rename modules/core/include/opencv2/core/{device => cuda}/block.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/border_interpolate.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/color.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/common.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/datamov_utils.hpp (97%) rename modules/core/include/opencv2/core/{device => cuda}/detail/color_detail.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/detail/reduce.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/detail/reduce_key_val.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/detail/transform_detail.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/detail/type_traits_detail.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/detail/vec_distance_detail.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/dynamic_smem.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/emulation.hpp (97%) rename modules/core/include/opencv2/core/{device => cuda}/filters.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/funcattrib.hpp (96%) rename modules/core/include/opencv2/core/{device => cuda}/functional.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/limits.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/reduce.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/saturate_cast.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/scan.hpp (96%) rename modules/core/include/opencv2/core/{device => cuda}/simd_functions.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/static_check.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/transform.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/type_traits.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/utility.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/vec_distance.hpp (98%) rename modules/core/include/opencv2/core/{device => cuda}/vec_math.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/vec_traits.hpp (99%) rename modules/core/include/opencv2/core/{device => cuda}/warp.hpp (97%) rename modules/core/include/opencv2/core/{device => cuda}/warp_reduce.hpp (96%) rename modules/core/include/opencv2/core/{device => cuda}/warp_shuffle.hpp (99%) diff --git a/cmake/OpenCVModule.cmake b/cmake/OpenCVModule.cmake index 1ce9c89..fe31b70 100644 --- a/cmake/OpenCVModule.cmake +++ b/cmake/OpenCVModule.cmake @@ -432,16 +432,16 @@ macro(ocv_glob_module_sources) file(GLOB lib_hdrs "include/opencv2/*.hpp" "include/opencv2/${name}/*.hpp" "include/opencv2/${name}/*.h") file(GLOB lib_hdrs_detail "include/opencv2/${name}/detail/*.hpp" "include/opencv2/${name}/detail/*.h") - file(GLOB lib_device_srcs "src/cuda/*.cu") - set(device_objs "") - set(lib_device_hdrs "") + file(GLOB lib_cuda_srcs "src/cuda/*.cu") + set(cuda_objs "") + set(lib_cuda_hdrs "") - if (HAVE_CUDA AND lib_device_srcs) + if(HAVE_CUDA AND lib_cuda_srcs) ocv_include_directories(${CUDA_INCLUDE_DIRS}) - file(GLOB lib_device_hdrs "src/cuda/*.hpp") + file(GLOB lib_cuda_hdrs "src/cuda/*.hpp") - ocv_cuda_compile(device_objs ${lib_device_srcs} ${lib_device_hdrs}) - source_group("Src\\Cuda" FILES ${lib_device_srcs} ${lib_device_hdrs}) + ocv_cuda_compile(cuda_objs ${lib_cuda_srcs} ${lib_cuda_hdrs}) + source_group("Src\\Cuda" FILES ${lib_cuda_srcs} ${lib_cuda_hdrs}) endif() file(GLOB cl_kernels "src/opencl/*.cl") @@ -457,7 +457,7 @@ macro(ocv_glob_module_sources) endif() ocv_set_module_sources(${ARGN} HEADERS ${lib_hdrs} ${lib_hdrs_detail} - SOURCES ${lib_srcs} ${lib_int_hdrs} ${device_objs} ${lib_device_srcs} ${lib_device_hdrs}) + SOURCES ${lib_srcs} ${lib_int_hdrs} ${cuda_objs} ${lib_cuda_srcs} ${lib_cuda_hdrs}) source_group("Src" FILES ${lib_srcs} ${lib_int_hdrs}) source_group("Include" FILES ${lib_hdrs}) diff --git a/modules/core/CMakeLists.txt b/modules/core/CMakeLists.txt index 6b5c3ac..f2b130d 100644 --- a/modules/core/CMakeLists.txt +++ b/modules/core/CMakeLists.txt @@ -6,10 +6,10 @@ if(HAVE_CUDA) ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef) endif() -file(GLOB lib_device_hdrs "include/opencv2/${name}/device/*.hpp" "include/opencv2/${name}/device/*.h") -file(GLOB lib_device_hdrs_detail "include/opencv2/${name}/device/detail/*.hpp" "include/opencv2/${name}/device/detail/*.h") +file(GLOB lib_cuda_hdrs "include/opencv2/${name}/cuda/*.hpp" "include/opencv2/${name}/cuda/*.h") +file(GLOB lib_cuda_hdrs_detail "include/opencv2/${name}/cuda/detail/*.hpp" "include/opencv2/${name}/cuda/detail/*.h") ocv_glob_module_sources(SOURCES "${opencv_core_BINARY_DIR}/version_string.inc" - HEADERS ${lib_device_hdrs} ${lib_device_hdrs_detail}) + HEADERS ${lib_cuda_hdrs} ${lib_cuda_hdrs_detail}) ocv_create_module() ocv_add_precompiled_headers(${the_module}) diff --git a/modules/core/include/opencv2/core/device/block.hpp b/modules/core/include/opencv2/core/cuda/block.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/block.hpp rename to modules/core/include/opencv2/core/cuda/block.hpp index 86ce205..b2bcb9f 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { struct Block { diff --git a/modules/core/include/opencv2/core/device/border_interpolate.hpp b/modules/core/include/opencv2/core/cuda/border_interpolate.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/border_interpolate.hpp rename to modules/core/include/opencv2/core/cuda/border_interpolate.hpp index 2ec9743..3854e84 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { ////////////////////////////////////////////////////////////// // BrdConstant @@ -709,6 +709,6 @@ namespace cv { namespace gpu { namespace device const int width; const D val; }; -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_BORDER_INTERPOLATE_HPP__ diff --git a/modules/core/include/opencv2/core/device/color.hpp b/modules/core/include/opencv2/core/cuda/color.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/color.hpp rename to modules/core/include/opencv2/core/cuda/color.hpp index c087d17..c625308 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { // 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 device OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lbgra, 4, 4, false, 0) #undef OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_BORDER_INTERPOLATE_HPP__ diff --git a/modules/core/include/opencv2/core/device/common.hpp b/modules/core/include/opencv2/core/cuda/common.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/common.hpp rename to modules/core/include/opencv2/core/cuda/common.hpp index 64d82c8..680ec49 100644 --- a/modules/core/include/opencv2/core/device/common.hpp +++ b/modules/core/include/opencv2/core/cuda/common.hpp @@ -92,7 +92,7 @@ namespace cv { namespace gpu return (total + grain - 1) / grain; } - namespace device + namespace cuda { using cv::gpu::divUp; diff --git a/modules/core/include/opencv2/core/device/datamov_utils.hpp b/modules/core/include/opencv2/core/cuda/datamov_utils.hpp similarity index 97% rename from modules/core/include/opencv2/core/device/datamov_utils.hpp rename to modules/core/include/opencv2/core/cuda/datamov_utils.hpp index a3f62fb..d2aadf9 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 @@ -100,6 +100,6 @@ namespace cv { namespace gpu { namespace device #undef OPENCV_GPU_ASM_PTR #endif // __CUDA_ARCH__ >= 200 -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_DATAMOV_UTILS_HPP__ diff --git a/modules/core/include/opencv2/core/device/detail/color_detail.hpp b/modules/core/include/opencv2/core/cuda/detail/color_detail.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/detail/color_detail.hpp rename to modules/core/include/opencv2/core/cuda/detail/color_detail.hpp index d02027f..5853475 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { #ifndef CV_DESCALE #define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n)) @@ -149,7 +149,7 @@ namespace cv { namespace gpu { namespace device #define OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(name, scn, bidx, green_bits) \ struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2RGB5x5 functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(name, dcn, bidx, green_bits) \ struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB5x52RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_GRAY2RGB_TRAITS(name, dcn) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::Gray2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_GRAY2RGB5x5_TRAITS(name, green_bits) \ struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::Gray2RGB5x5 functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB5x52GRAY_TRAITS(name, green_bits) \ struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB5x52Gray functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS(name, scn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2Gray functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2YUV functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::YUV2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2YCrCb functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::YCrCb2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2XYZ functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::XYZ2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2HSV functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template struct name ## _full_traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2HSV functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template <> struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2HSV functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template <> struct name ## _full_traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2HSV functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::HSV2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template struct name ## _full_traits \ { \ - typedef ::cv::gpu::device::color_detail::HSV2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template <> struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::HSV2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template <> struct name ## _full_traits \ { \ - typedef ::cv::gpu::device::color_detail::HSV2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2HLS functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template struct name ## _full_traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2HLS functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template <> struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2HLS functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template <> struct name ## _full_traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2HLS functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(name, scn, dcn, bidx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::HLS2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template struct name ## _full_traits \ { \ - typedef ::cv::gpu::device::color_detail::HLS2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template <> struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::HLS2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device }; \ template <> struct name ## _full_traits \ { \ - typedef ::cv::gpu::device::color_detail::HLS2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(name, scn, dcn, srgb, blueIdx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2Lab functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::Lab2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(name, scn, dcn, srgb, blueIdx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::RGB2Luv functor_type; \ + typedef ::cv::gpu::cuda::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 device #define OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \ template struct name ## _traits \ { \ - typedef ::cv::gpu::device::color_detail::Luv2RGB functor_type; \ + typedef ::cv::gpu::cuda::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 device #undef CV_DESCALE -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_COLOR_DETAIL_HPP__ diff --git a/modules/core/include/opencv2/core/device/detail/reduce.hpp b/modules/core/include/opencv2/core/cuda/detail/reduce.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/detail/reduce.hpp rename to modules/core/include/opencv2/core/cuda/detail/reduce.hpp index 091a160..6866016 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { namespace reduce_detail { diff --git a/modules/core/include/opencv2/core/device/detail/reduce_key_val.hpp b/modules/core/include/opencv2/core/cuda/detail/reduce_key_val.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/detail/reduce_key_val.hpp rename to modules/core/include/opencv2/core/cuda/detail/reduce_key_val.hpp index a84e0c2..dde37dc 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { namespace reduce_key_val_detail { diff --git a/modules/core/include/opencv2/core/device/detail/transform_detail.hpp b/modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/detail/transform_detail.hpp rename to modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp index 10da593..0ed8d64 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { namespace transform_detail { @@ -390,6 +390,6 @@ namespace cv { namespace gpu { namespace device } }; } // namespace transform_detail -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_TRANSFORM_DETAIL_HPP__ diff --git a/modules/core/include/opencv2/core/device/detail/type_traits_detail.hpp b/modules/core/include/opencv2/core/cuda/detail/type_traits_detail.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/detail/type_traits_detail.hpp rename to modules/core/include/opencv2/core/cuda/detail/type_traits_detail.hpp index 97ff00d..ee92184 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { namespace type_traits_detail { @@ -182,6 +182,6 @@ namespace cv { namespace gpu { namespace device enum { value = 1 }; }; } // namespace type_traits_detail -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_TYPE_TRAITS_DETAIL_HPP__ diff --git a/modules/core/include/opencv2/core/device/detail/vec_distance_detail.hpp b/modules/core/include/opencv2/core/cuda/detail/vec_distance_detail.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/detail/vec_distance_detail.hpp rename to modules/core/include/opencv2/core/cuda/detail/vec_distance_detail.hpp index 78ab556..435e69c 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { namespace vec_distance_detail { @@ -112,6 +112,6 @@ namespace cv { namespace gpu { namespace device } }; } // namespace vec_distance_detail -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_VEC_DISTANCE_DETAIL_HPP__ diff --git a/modules/core/include/opencv2/core/device/dynamic_smem.hpp b/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/dynamic_smem.hpp rename to modules/core/include/opencv2/core/cuda/dynamic_smem.hpp index cf431d9..41f9ecc 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { template struct DynamicSharedMem { diff --git a/modules/core/include/opencv2/core/device/emulation.hpp b/modules/core/include/opencv2/core/cuda/emulation.hpp similarity index 97% rename from modules/core/include/opencv2/core/device/emulation.hpp rename to modules/core/include/opencv2/core/cuda/emulation.hpp index bf47bc5..e2c7437 100644 --- a/modules/core/include/opencv2/core/device/emulation.hpp +++ b/modules/core/include/opencv2/core/cuda/emulation.hpp @@ -45,7 +45,7 @@ #include "warp_reduce.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { struct Emulation { @@ -133,6 +133,6 @@ namespace cv { namespace gpu { namespace device } }; }; -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* OPENCV_GPU_EMULATION_HPP_ */ diff --git a/modules/core/include/opencv2/core/device/filters.hpp b/modules/core/include/opencv2/core/cuda/filters.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/filters.hpp rename to modules/core/include/opencv2/core/cuda/filters.hpp index d193969..607aecf 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { template struct PointFilter { @@ -273,6 +273,6 @@ namespace cv { namespace gpu { namespace device float scale_x, scale_y; int width, haight; }; -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_FILTERS_HPP__ diff --git a/modules/core/include/opencv2/core/device/funcattrib.hpp b/modules/core/include/opencv2/core/cuda/funcattrib.hpp similarity index 96% rename from modules/core/include/opencv2/core/device/funcattrib.hpp rename to modules/core/include/opencv2/core/cuda/funcattrib.hpp index 2ed7980..e97b5be 100644 --- a/modules/core/include/opencv2/core/device/funcattrib.hpp +++ b/modules/core/include/opencv2/core/cuda/funcattrib.hpp @@ -45,7 +45,7 @@ #include -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template void printFuncAttrib(Func& func) @@ -66,6 +66,6 @@ namespace cv { namespace gpu { namespace device printf("\n"); fflush(stdout); } -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_ */ diff --git a/modules/core/include/opencv2/core/device/functional.hpp b/modules/core/include/opencv2/core/cuda/functional.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/functional.hpp rename to modules/core/include/opencv2/core/cuda/functional.hpp index 6064e8e..5ac90eb 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { // Function Objects template struct unary_function : public std::unary_function {}; @@ -786,6 +786,6 @@ namespace cv { namespace gpu { namespace device #define OPENCV_GPU_TRANSFORM_FUNCTOR_TRAITS(type) \ template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type > -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_FUNCTIONAL_HPP__ diff --git a/modules/core/include/opencv2/core/device/limits.hpp b/modules/core/include/opencv2/core/cuda/limits.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/limits.hpp rename to modules/core/include/opencv2/core/cuda/limits.hpp index b040f19..fc5e7b2 100644 --- a/modules/core/include/opencv2/core/device/limits.hpp +++ b/modules/core/include/opencv2/core/cuda/limits.hpp @@ -46,7 +46,7 @@ #include #include "common.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct numeric_limits { @@ -230,6 +230,6 @@ namespace cv { namespace gpu { namespace device __device__ __forceinline__ static type signaling_NaN(); static const bool is_signed = true; }; -}}} // namespace cv { namespace gpu { namespace device { +}}} // namespace cv { namespace gpu { namespace cuda { #endif // __OPENCV_GPU_LIMITS_GPU_HPP__ diff --git a/modules/core/include/opencv2/core/device/reduce.hpp b/modules/core/include/opencv2/core/cuda/reduce.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/reduce.hpp rename to modules/core/include/opencv2/core/cuda/reduce.hpp index 2161b06..ca16a20 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { template __device__ __forceinline__ void reduce(volatile T* smem, T& val, unsigned int tid, const Op& op) diff --git a/modules/core/include/opencv2/core/device/saturate_cast.hpp b/modules/core/include/opencv2/core/cuda/saturate_cast.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/saturate_cast.hpp rename to modules/core/include/opencv2/core/cuda/saturate_cast.hpp index 7a2799f..f80a43a 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { 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/device/scan.hpp b/modules/core/include/opencv2/core/cuda/scan.hpp similarity index 96% rename from modules/core/include/opencv2/core/device/scan.hpp rename to modules/core/include/opencv2/core/cuda/scan.hpp index 777754c..0493ee3 100644 --- a/modules/core/include/opencv2/core/device/scan.hpp +++ b/modules/core/include/opencv2/core/cuda/scan.hpp @@ -43,12 +43,12 @@ #ifndef __OPENCV_GPU_SCAN_HPP__ #define __OPENCV_GPU_SCAN_HPP__ -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/utility.hpp" -#include "opencv2/core/device/warp.hpp" -#include "opencv2/core/device/warp_shuffle.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/warp.hpp" +#include "opencv2/core/cuda/warp_shuffle.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { enum ScanKind { EXCLUSIVE = 0, INCLUSIVE = 1 }; @@ -174,13 +174,13 @@ namespace cv { namespace gpu { namespace device __device__ T warpScanInclusive(T idata, volatile T* s_Data, unsigned int tid) { #if __CUDA_ARCH__ >= 300 - const unsigned int laneId = cv::gpu::device::Warp::laneId(); + const unsigned int laneId = cv::gpu::cuda::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::device::shfl_up(idata, i); + const T n = cv::gpu::cuda::shfl_up(idata, i); if (laneId >= i) idata += n; } diff --git a/modules/core/include/opencv2/core/device/simd_functions.hpp b/modules/core/include/opencv2/core/cuda/simd_functions.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/simd_functions.hpp rename to modules/core/include/opencv2/core/cuda/simd_functions.hpp index b0377e5..b6d2fe1 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { // 2 diff --git a/modules/core/include/opencv2/core/device/static_check.hpp b/modules/core/include/opencv2/core/cuda/static_check.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/static_check.hpp rename to modules/core/include/opencv2/core/cuda/static_check.hpp index e77691b..db1bb8e 100644 --- a/modules/core/include/opencv2/core/device/static_check.hpp +++ b/modules/core/include/opencv2/core/cuda/static_check.hpp @@ -51,7 +51,7 @@ namespace cv { namespace gpu { - namespace device + namespace cuda { template struct Static {}; diff --git a/modules/core/include/opencv2/core/device/transform.hpp b/modules/core/include/opencv2/core/cuda/transform.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/transform.hpp rename to modules/core/include/opencv2/core/cuda/transform.hpp index 636caac..bf9b68b 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { 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/device/type_traits.hpp b/modules/core/include/opencv2/core/cuda/type_traits.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/type_traits.hpp rename to modules/core/include/opencv2/core/cuda/type_traits.hpp index 1b36acc..0fb5fda 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { template struct IsSimpleParameter { diff --git a/modules/core/include/opencv2/core/device/utility.hpp b/modules/core/include/opencv2/core/cuda/utility.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/utility.hpp rename to modules/core/include/opencv2/core/cuda/utility.hpp index 83eaaa2..7997385 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { #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 device return false; } -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_UTILITY_HPP__ diff --git a/modules/core/include/opencv2/core/device/vec_distance.hpp b/modules/core/include/opencv2/core/cuda/vec_distance.hpp similarity index 98% rename from modules/core/include/opencv2/core/device/vec_distance.hpp rename to modules/core/include/opencv2/core/cuda/vec_distance.hpp index d5b4bb2..5521001 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { template struct L1Dist { @@ -219,6 +219,6 @@ namespace cv { namespace gpu { namespace device U vec1Vals[MAX_LEN / THREAD_DIM]; }; -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_VEC_DISTANCE_HPP__ diff --git a/modules/core/include/opencv2/core/device/vec_math.hpp b/modules/core/include/opencv2/core/cuda/vec_math.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/vec_math.hpp rename to modules/core/include/opencv2/core/cuda/vec_math.hpp index 1c46dc0..238c71e 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { namespace vec_math_detail { @@ -325,6 +325,6 @@ namespace cv { namespace gpu { namespace device #undef OPENCV_GPU_IMPLEMENT_VEC_BINOP #undef OPENCV_GPU_IMPLEMENT_VEC_OP #undef OPENCV_GPU_IMPLEMENT_VEC_INT_OP -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_VECMATH_HPP__ diff --git a/modules/core/include/opencv2/core/device/vec_traits.hpp b/modules/core/include/opencv2/core/cuda/vec_traits.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/vec_traits.hpp rename to modules/core/include/opencv2/core/cuda/vec_traits.hpp index 8d179c8..9cdf64b 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { template struct TypeVec; @@ -275,6 +275,6 @@ namespace cv { namespace gpu { namespace device 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 device +}}} // namespace cv { namespace gpu { namespace cuda #endif // __OPENCV_GPU_VEC_TRAITS_HPP__ diff --git a/modules/core/include/opencv2/core/device/warp.hpp b/modules/core/include/opencv2/core/cuda/warp.hpp similarity index 97% rename from modules/core/include/opencv2/core/device/warp.hpp rename to modules/core/include/opencv2/core/cuda/warp.hpp index 0f1dc79..14cf6cb 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { struct Warp { @@ -126,6 +126,6 @@ namespace cv { namespace gpu { namespace device *t = value; } }; -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* __OPENCV_GPU_DEVICE_WARP_HPP__ */ diff --git a/modules/core/include/opencv2/core/device/warp_reduce.hpp b/modules/core/include/opencv2/core/cuda/warp_reduce.hpp similarity index 96% rename from modules/core/include/opencv2/core/device/warp_reduce.hpp rename to modules/core/include/opencv2/core/cuda/warp_reduce.hpp index d4e64c4..7d4d838 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { template __device__ __forceinline__ T warp_reduce(volatile T *ptr , const unsigned int tid = threadIdx.x) @@ -63,6 +63,6 @@ namespace cv { namespace gpu { namespace device return ptr[tid - lane]; } -}}} // namespace cv { namespace gpu { namespace device { +}}} // namespace cv { namespace gpu { namespace cuda { #endif /* OPENCV_GPU_WARP_REDUCE_HPP__ */ diff --git a/modules/core/include/opencv2/core/device/warp_shuffle.hpp b/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp similarity index 99% rename from modules/core/include/opencv2/core/device/warp_shuffle.hpp rename to modules/core/include/opencv2/core/cuda/warp_shuffle.hpp index 8b4479a..954b618 100644 --- a/modules/core/include/opencv2/core/device/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 device +namespace cv { namespace gpu { namespace cuda { template __device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize) diff --git a/modules/core/include/opencv2/core/cuda_devptrs.hpp b/modules/core/include/opencv2/core/cuda_devptrs.hpp index 5af5ab8..9e0ba11 100644 --- a/modules/core/include/opencv2/core/cuda_devptrs.hpp +++ b/modules/core/include/opencv2/core/cuda_devptrs.hpp @@ -148,24 +148,6 @@ namespace cv typedef DevMem2Db DevMem2D; typedef DevMem2D_ DevMem2Df; typedef DevMem2D_ DevMem2Di; - -//#undef __CV_GPU_DEPR_BEFORE__ -//#undef __CV_GPU_DEPR_AFTER__ - - namespace device - { - using cv::gpu::PtrSz; - using cv::gpu::PtrStep; - using cv::gpu::PtrStepSz; - - using cv::gpu::PtrStepSzb; - using cv::gpu::PtrStepSzf; - using cv::gpu::PtrStepSzi; - - using cv::gpu::PtrStepb; - using cv::gpu::PtrStepf; - using cv::gpu::PtrStepi; - } } } diff --git a/modules/core/src/cuda/matrix_operations.cu b/modules/core/src/cuda/matrix_operations.cu index dd67c54..2cb8218 100644 --- a/modules/core/src/cuda/matrix_operations.cu +++ b/modules/core/src/cuda/matrix_operations.cu @@ -40,12 +40,12 @@ // //M*/ -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/transform.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/type_traits.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/transform.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/type_traits.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { void writeScalar(const uchar*); void writeScalar(const schar*); @@ -58,7 +58,7 @@ namespace cv { namespace gpu { namespace device void convert_gpu(PtrStepSzb, int, PtrStepSzb, int, double, double, cudaStream_t); }}} -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct shift_and_sizeof; template <> struct shift_and_sizeof { enum { shift = 0 }; }; @@ -76,9 +76,9 @@ namespace cv { namespace gpu { namespace device template void copyToWithMask(PtrStepSzb src, PtrStepSzb dst, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream) { if (colorMask) - cv::gpu::device::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMask(mask), stream); + cv::gpu::cuda::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMask(mask), stream); else - cv::gpu::device::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMaskChannels(mask, cn), stream); + cv::gpu::cuda::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 device cudaSafeCall( cudaSetDoubleForDevice(&alpha) ); cudaSafeCall( cudaSetDoubleForDevice(&beta) ); Convertor op(static_cast(alpha), static_cast(beta)); - cv::gpu::device::transform((PtrStepSz)src, (PtrStepSz)dst, op, WithOutMask(), stream); + cv::gpu::cuda::transform((PtrStepSz)src, (PtrStepSz)dst, op, WithOutMask(), stream); } #if defined __clang__ @@ -379,4 +379,4 @@ namespace cv { namespace gpu { namespace device #if defined __clang__ # pragma clang diagnostic pop #endif -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda diff --git a/modules/core/src/gpumat.cpp b/modules/core/src/gpumat.cpp index 09e1562..12252b5 100644 --- a/modules/core/src/gpumat.cpp +++ b/modules/core/src/gpumat.cpp @@ -870,7 +870,7 @@ namespace #else // HAVE_CUDA -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { void copyToWithMask_gpu(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream); @@ -888,13 +888,13 @@ namespace template void kernelSetCaller(GpuMat& src, Scalar s, cudaStream_t stream) { Scalar_ sf = s; - cv::gpu::device::set_to_gpu(src, sf.val, src.channels(), stream); + cv::gpu::cuda::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::device::set_to_gpu(src, sf.val, mask, src.channels(), stream); + cv::gpu::cuda::set_to_gpu(src, sf.val, mask, src.channels(), stream); } } @@ -918,17 +918,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::device::copyToWithMask_gpu(src.reshape(1), dst.reshape(1), src.elemSize1(), src.channels(), mask.reshape(1), mask.channels() != 1, stream); + cv::gpu::cuda::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::device::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), 1.0, 0.0, 0); + cv::gpu::cuda::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::device::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), alpha, beta, stream); + cv::gpu::cuda::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 ad0b894..cefc29f 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 device +namespace cv { namespace gpu { namespace cuda { 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::device::mathfunc; + using namespace ::cv::gpu::cuda::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::device::mathfunc; + using namespace ::cv::gpu::cuda::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 46f3ba2..b474823 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 device { +namespace cv { namespace gpu { namespace cuda { 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::device::bgfg_gmg; + using namespace cv::gpu::cuda::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::device::bgfg_gmg; + using namespace cv::gpu::cuda::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 05b4ffb..13db079 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 device +namespace cv { namespace gpu { namespace cuda { 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::device::mog; + using namespace cv::gpu::cuda::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::device::mog; + using namespace cv::gpu::cuda::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::device::mog; + using namespace cv::gpu::cuda::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::device::mog; + using namespace cv::gpu::cuda::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::device::mog; + using namespace cv::gpu::cuda::mog; backgroundImage.create(frameSize_, frameType_); diff --git a/modules/gpu/src/bilateral_filter.cpp b/modules/gpu/src/bilateral_filter.cpp index c60954d..41640a5 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 device +namespace cv { namespace gpu { namespace cuda { namespace disp_bilateral_filter { @@ -65,7 +65,7 @@ namespace cv { namespace gpu { namespace device } }}} -using namespace ::cv::gpu::device::disp_bilateral_filter; +using namespace ::cv::gpu::cuda::disp_bilateral_filter; namespace { diff --git a/modules/gpu/src/blend.cpp b/modules/gpu/src/blend.cpp index 1f4cf7d..0819338 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 device +namespace cv { namespace gpu { namespace cuda { namespace blend { @@ -62,7 +62,7 @@ namespace cv { namespace gpu { namespace device } }}} -using namespace ::cv::gpu::device::blend; +using namespace ::cv::gpu::cuda::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 b57b13f..1f9c11c 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 device +namespace cv { namespace gpu { namespace cuda { 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::device::bf_match; + using namespace cv::gpu::cuda::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::device::bf_match; + using namespace cv::gpu::cuda::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::device::bf_knnmatch; + using namespace cv::gpu::cuda::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::device::bf_knnmatch; + using namespace cv::gpu::cuda::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::device::bf_radius_match; + using namespace cv::gpu::cuda::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::device::bf_radius_match; + using namespace cv::gpu::cuda::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 ab7f63f..78408c7 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 device +namespace cv { namespace gpu { namespace cuda { namespace transform_points { @@ -78,7 +78,7 @@ namespace cv { namespace gpu { namespace device } }}} -using namespace ::cv::gpu::device; +using namespace ::cv::gpu::cuda; namespace { diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp index 10108e9..3ba8368 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 device +namespace cv { namespace gpu { namespace cuda { namespace lbp { @@ -441,7 +441,7 @@ public: acc += level.sFrame.width + 1; } - device::lbp::classifyPyramid(image.cols, image.rows, NxM.width - 1, NxM.height - 1, iniScale, scaleFactor, total, stage_mat, stage_mat.cols / sizeof(Stage), nodes_mat, + 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, leaves_mat, subsets_mat, features_mat, subsetSize, candidates, dclassified.ptr(), integral); } @@ -449,7 +449,7 @@ public: return 0; cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); - device::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified.ptr()); + cuda::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified.ptr()); cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) ); cudaSafeCall( cudaDeviceSynchronize() ); diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp index 76793d5..9f6ca25 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_nog #include "cvt_color_internal.h" namespace cv { namespace gpu { - namespace device + namespace cuda { 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::device; +using namespace ::cv::gpu::cuda; namespace { @@ -77,7 +77,7 @@ namespace void bgr_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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); - device::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream)); + cuda::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream)); } void gray_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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); - device::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream)); + cuda::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); - device::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream)); + cuda::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream)); } void rgb_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) { - using namespace cv::gpu::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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::device; + using namespace cv::gpu::cuda; 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) - device::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); + cuda::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); else - device::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); + cuda::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); - device::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); + cuda::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 45e44a5..4110dbc 100644 --- a/modules/gpu/src/cuda/NV12ToARGB.cu +++ b/modules/gpu/src/cuda/NV12ToARGB.cu @@ -49,9 +49,9 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device { +namespace cv { namespace gpu { namespace cuda { 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 2177bb2..629ac69 100644 --- a/modules/gpu/src/cuda/bf_knnmatch.cu +++ b/modules/gpu/src/cuda/bf_knnmatch.cu @@ -42,15 +42,15 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/utility.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/vec_distance.hpp" -#include "opencv2/core/device/datamov_utils.hpp" -#include "opencv2/core/device/warp_shuffle.hpp" - -namespace cv { namespace gpu { namespace device +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/vec_distance.hpp" +#include "opencv2/core/cuda/datamov_utils.hpp" +#include "opencv2/core/cuda/warp_shuffle.hpp" + +namespace cv { namespace gpu { namespace cuda { namespace bf_knnmatch { @@ -1249,7 +1249,7 @@ namespace cv { namespace gpu { namespace device //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 device { +}}} // namespace cv { namespace gpu { namespace cuda { #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/bf_match.cu b/modules/gpu/src/cuda/bf_match.cu index c3efa74..1e7cf8a 100644 --- a/modules/gpu/src/cuda/bf_match.cu +++ b/modules/gpu/src/cuda/bf_match.cu @@ -42,14 +42,14 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/utility.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/vec_distance.hpp" -#include "opencv2/core/device/datamov_utils.hpp" - -namespace cv { namespace gpu { namespace device +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/vec_distance.hpp" +#include "opencv2/core/cuda/datamov_utils.hpp" + +namespace cv { namespace gpu { namespace cuda { namespace bf_match { @@ -768,7 +768,7 @@ namespace cv { namespace gpu { namespace device //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 device { +}}} // namespace cv { namespace gpu { namespace cuda { #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/bf_radius_match.cu b/modules/gpu/src/cuda/bf_radius_match.cu index d5ba42f..1c08577 100644 --- a/modules/gpu/src/cuda/bf_radius_match.cu +++ b/modules/gpu/src/cuda/bf_radius_match.cu @@ -42,13 +42,13 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/utility.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/vec_distance.hpp" -#include "opencv2/core/device/datamov_utils.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/vec_distance.hpp" +#include "opencv2/core/cuda/datamov_utils.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace bf_radius_match { @@ -457,7 +457,7 @@ namespace cv { namespace gpu { namespace device //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 device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/bgfg_gmg.cu b/modules/gpu/src/cuda/bgfg_gmg.cu index a943815..0047fe5 100644 --- a/modules/gpu/src/cuda/bgfg_gmg.cu +++ b/modules/gpu/src/cuda/bgfg_gmg.cu @@ -42,11 +42,11 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/limits.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/limits.hpp" -namespace cv { namespace gpu { namespace device { +namespace cv { namespace gpu { namespace cuda { 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 57ab1f2..6a514f7 100644 --- a/modules/gpu/src/cuda/bgfg_mog.cu +++ b/modules/gpu/src/cuda/bgfg_mog.cu @@ -42,12 +42,12 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/limits.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/limits.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace mog { diff --git a/modules/gpu/src/cuda/bilateral_filter.cu b/modules/gpu/src/cuda/bilateral_filter.cu index 64a2a30..17c7c1d 100644 --- a/modules/gpu/src/cuda/bilateral_filter.cu +++ b/modules/gpu/src/cuda/bilateral_filter.cu @@ -42,10 +42,10 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/border_interpolate.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" using namespace cv::gpu; @@ -55,7 +55,7 @@ typedef unsigned short ushort; ////////////////////////////////////////////////////////////////////////////////// /// Bilateral filtering -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -163,7 +163,7 @@ namespace cv { namespace gpu { namespace device #define OCV_INSTANTIATE_BILATERAL_FILTER(T) \ - template void cv::gpu::device::imgproc::bilateral_filter_gpu(const PtrStepSzb&, PtrStepSzb, int, float, float, int, cudaStream_t); + template void cv::gpu::cuda::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 db49e6c..b4ecfbb 100644 --- a/modules/gpu/src/cuda/blend.cu +++ b/modules/gpu/src/cuda/blend.cu @@ -42,9 +42,9 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace blend { @@ -115,7 +115,7 @@ namespace cv { namespace gpu { namespace device cudaSafeCall(cudaDeviceSynchronize()); } } // namespace blend -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/calib3d.cu b/modules/gpu/src/cuda/calib3d.cu index 2c122fd..5d12405 100644 --- a/modules/gpu/src/cuda/calib3d.cu +++ b/modules/gpu/src/cuda/calib3d.cu @@ -42,12 +42,12 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/transform.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/reduce.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/transform.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/reduce.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { #define SOLVE_PNP_RANSAC_MAX_NUM_ITERS 200 @@ -79,7 +79,7 @@ namespace cv { namespace gpu { namespace device cudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3)); cudaSafeCall(cudaMemcpyToSymbol(crot2, rot + 6, sizeof(float) * 3)); cudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3)); - cv::gpu::device::transform(src, dst, TransformOp(), WithOutMask(), stream); + cv::gpu::cuda::transform(src, dst, TransformOp(), WithOutMask(), stream); } } // namespace transform_points @@ -120,7 +120,7 @@ namespace cv { namespace gpu { namespace device cudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3)); cudaSafeCall(cudaMemcpyToSymbol(cproj0, proj, sizeof(float) * 3)); cudaSafeCall(cudaMemcpyToSymbol(cproj1, proj + 3, sizeof(float) * 3)); - cv::gpu::device::transform(src, dst, ProjectOp(), WithOutMask(), stream); + cv::gpu::cuda::transform(src, dst, ProjectOp(), WithOutMask(), stream); } } // namespace project_points @@ -187,7 +187,7 @@ namespace cv { namespace gpu { namespace device cudaSafeCall( cudaDeviceSynchronize() ); } } // namespace solvepnp_ransac -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/canny.cu b/modules/gpu/src/cuda/canny.cu index 2b7bcf7..151f234 100644 --- a/modules/gpu/src/cuda/canny.cu +++ b/modules/gpu/src/cuda/canny.cu @@ -44,14 +44,14 @@ #include #include //std::swap -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/emulation.hpp" -#include "opencv2/core/device/transform.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/utility.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/emulation.hpp" +#include "opencv2/core/cuda/transform.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/utility.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace canny { @@ -77,7 +77,7 @@ namespace canny }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits : DefaultTransformFunctorTraits { @@ -475,7 +475,7 @@ namespace canny }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits : DefaultTransformFunctorTraits { diff --git a/modules/gpu/src/cuda/ccomponetns.cu b/modules/gpu/src/cuda/ccomponetns.cu index 5519efe..2e52ff2 100644 --- a/modules/gpu/src/cuda/ccomponetns.cu +++ b/modules/gpu/src/cuda/ccomponetns.cu @@ -42,15 +42,15 @@ #if !defined CUDA_DISABLER -#include -#include -#include -#include +#include +#include +#include +#include #include #include -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace ccl { diff --git a/modules/gpu/src/cuda/clahe.cu b/modules/gpu/src/cuda/clahe.cu index 35ce9c1..c5ecbf4 100644 --- a/modules/gpu/src/cuda/clahe.cu +++ b/modules/gpu/src/cuda/clahe.cu @@ -42,15 +42,15 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/emulation.hpp" -#include "opencv2/core/device/scan.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/saturate_cast.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/emulation.hpp" +#include "opencv2/core/cuda/scan.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace clahe { diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu index 61794e2..a388735 100644 --- a/modules/gpu/src/cuda/color.cu +++ b/modules/gpu/src/cuda/color.cu @@ -42,12 +42,12 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/transform.hpp" -#include "opencv2/core/device/color.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/transform.hpp" +#include "opencv2/core/cuda/color.hpp" #include "cvt_color_internal.h" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { OPENCV_GPU_TRANSFORM_FUNCTOR_TRAITS(bgra_to_rgba_traits::functor_type) { @@ -229,7 +229,7 @@ namespace cv { namespace gpu { namespace device 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::device::transform((PtrStepSz)src, (PtrStepSz)dst, functor, WithOutMask(), stream); \ + cv::gpu::cuda::transform((PtrStepSz)src, (PtrStepSz)dst, functor, WithOutMask(), stream); \ } #define OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(name) \ @@ -456,6 +456,6 @@ namespace cv { namespace gpu { namespace device #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F_FULL -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/column_filter.h b/modules/gpu/src/cuda/column_filter.h index d28ceaf..0027810 100644 --- a/modules/gpu/src/cuda/column_filter.h +++ b/modules/gpu/src/cuda/column_filter.h @@ -40,13 +40,13 @@ // //M*/ -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/border_interpolate.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace column_filter { diff --git a/modules/gpu/src/cuda/copy_make_border.cu b/modules/gpu/src/cuda/copy_make_border.cu index 3797c09..4354465 100644 --- a/modules/gpu/src/cuda/copy_make_border.cu +++ b/modules/gpu/src/cuda/copy_make_border.cu @@ -42,10 +42,10 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/border_interpolate.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -126,6 +126,6 @@ namespace cv { namespace gpu { namespace device 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 device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/debayer.cu b/modules/gpu/src/cuda/debayer.cu index 468c12f..a079bd5 100644 --- a/modules/gpu/src/cuda/debayer.cu +++ b/modules/gpu/src/cuda/debayer.cu @@ -42,14 +42,14 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/color.hpp" -#include "opencv2/core/device/saturate_cast.hpp" - -namespace cv { namespace gpu { namespace device +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/color.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" + +namespace cv { namespace gpu { namespace cuda { template struct Bayer2BGR; diff --git a/modules/gpu/src/cuda/disp_bilateral_filter.cu b/modules/gpu/src/cuda/disp_bilateral_filter.cu index 9a01b8d..c3edff3 100644 --- a/modules/gpu/src/cuda/disp_bilateral_filter.cu +++ b/modules/gpu/src/cuda/disp_bilateral_filter.cu @@ -42,10 +42,10 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/limits.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/limits.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace disp_bilateral_filter { @@ -218,6 +218,6 @@ namespace cv { namespace gpu { namespace device 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 device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu index 06003e9..1d3b7ca 100644 --- a/modules/gpu/src/cuda/element_operations.cu +++ b/modules/gpu/src/cuda/element_operations.cu @@ -42,16 +42,16 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/transform.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/simd_functions.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/transform.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/simd_functions.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace arithm { @@ -193,7 +193,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits< arithm::VAdd4 > : arithm::ArithmFuncTraits { @@ -308,7 +308,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< arithm::AddScalar > : arithm::ArithmFuncTraits { @@ -428,7 +428,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits< arithm::VSub4 > : arithm::ArithmFuncTraits { @@ -657,7 +657,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits : arithm::ArithmFuncTraits { @@ -774,7 +774,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< arithm::MulScalar > : arithm::ArithmFuncTraits { @@ -925,7 +925,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits : arithm::ArithmFuncTraits { @@ -1111,7 +1111,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< arithm::DivInv > : arithm::ArithmFuncTraits { @@ -1240,7 +1240,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits< arithm::VAbsDiff4 > : arithm::ArithmFuncTraits { @@ -1305,7 +1305,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< arithm::AbsDiffScalar > : arithm::ArithmFuncTraits { @@ -1334,7 +1334,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////// // absMat -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< abs_func > : arithm::ArithmFuncTraits { @@ -1375,7 +1375,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< arithm::Sqr > : arithm::ArithmFuncTraits { @@ -1402,7 +1402,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////// // sqrtMat -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< sqrt_func > : arithm::ArithmFuncTraits { @@ -1429,7 +1429,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////// // logMat -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< log_func > : arithm::ArithmFuncTraits { @@ -1471,7 +1471,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< arithm::Exp > : arithm::ArithmFuncTraits { @@ -1554,7 +1554,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits< arithm::VCmpEq4 > : arithm::ArithmFuncTraits { @@ -1716,7 +1716,7 @@ namespace arithm #undef TYPE_VEC } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< arithm::CmpScalar > : arithm::ArithmFuncTraits { @@ -1875,7 +1875,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////////////////// // bitMat -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< bit_not > : arithm::ArithmFuncTraits { @@ -1948,7 +1948,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////////////////// // bitScalar -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< binder2nd< bit_and > > : arithm::ArithmFuncTraits { @@ -1967,17 +1967,17 @@ namespace arithm { template void bitScalarAnd(PtrStepSzb src1, uint src2, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::device::bind2nd(bit_and(), src2), WithOutMask(), stream); + transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::cuda::bind2nd(bit_and(), src2), WithOutMask(), stream); } template void bitScalarOr(PtrStepSzb src1, uint src2, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::device::bind2nd(bit_or(), src2), WithOutMask(), stream); + transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::cuda::bind2nd(bit_or(), src2), WithOutMask(), stream); } template void bitScalarXor(PtrStepSzb src1, uint src2, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::device::bind2nd(bit_xor(), src2), WithOutMask(), stream); + transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::cuda::bind2nd(bit_xor(), src2), WithOutMask(), stream); } template void bitScalarAnd(PtrStepSzb src1, uint src2, PtrStepSzb dst, cudaStream_t stream); @@ -2026,7 +2026,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits< arithm::VMin4 > : arithm::ArithmFuncTraits { @@ -2076,7 +2076,7 @@ namespace arithm template void minScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::device::bind2nd(minimum(), src2), WithOutMask(), stream); + transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::cuda::bind2nd(minimum(), src2), WithOutMask(), stream); } template void minScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); @@ -2118,7 +2118,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits< arithm::VMax4 > : arithm::ArithmFuncTraits { @@ -2168,7 +2168,7 @@ namespace arithm template void maxScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream) { - transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::device::bind2nd(maximum(), src2), WithOutMask(), stream); + transform((PtrStepSz) src1, (PtrStepSz) dst, cv::gpu::cuda::bind2nd(maximum(), src2), WithOutMask(), stream); } template void maxScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); @@ -2183,7 +2183,7 @@ namespace arithm ////////////////////////////////////////////////////////////////////////// // threshold -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< thresh_binary_func > : arithm::ArithmFuncTraits { @@ -2297,7 +2297,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct TransformFunctorTraits< arithm::PowOp > : arithm::ArithmFuncTraits { @@ -2372,7 +2372,7 @@ namespace arithm }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template struct AddWeightedTraits : DefaultTransformFunctorTraits< arithm::AddWeighted > { diff --git a/modules/gpu/src/cuda/fast.cu b/modules/gpu/src/cuda/fast.cu index 57a7a34..7ac7968 100644 --- a/modules/gpu/src/cuda/fast.cu +++ b/modules/gpu/src/cuda/fast.cu @@ -42,10 +42,10 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/utility.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/utility.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace fast { diff --git a/modules/gpu/src/cuda/fgd_bgfg.cu b/modules/gpu/src/cuda/fgd_bgfg.cu index ebc7cf5..b0da49f 100644 --- a/modules/gpu/src/cuda/fgd_bgfg.cu +++ b/modules/gpu/src/cuda/fgd_bgfg.cu @@ -42,16 +42,16 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/utility.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/functional.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/functional.hpp" #include "fgd_bgfg_common.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace bgfg { diff --git a/modules/gpu/src/cuda/gftt.cu b/modules/gpu/src/cuda/gftt.cu index 593053d..cc4467b 100644 --- a/modules/gpu/src/cuda/gftt.cu +++ b/modules/gpu/src/cuda/gftt.cu @@ -45,10 +45,10 @@ #include #include -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/utility.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/utility.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace gfft { diff --git a/modules/gpu/src/cuda/global_motion.cu b/modules/gpu/src/cuda/global_motion.cu index 79a19dd..f8d95d4 100644 --- a/modules/gpu/src/cuda/global_motion.cu +++ b/modules/gpu/src/cuda/global_motion.cu @@ -45,9 +45,9 @@ #include #include #include -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device { namespace globmotion { +namespace cv { namespace gpu { namespace cuda { namespace globmotion { __constant__ float cml[9]; __constant__ float cmr[9]; diff --git a/modules/gpu/src/cuda/hist.cu b/modules/gpu/src/cuda/hist.cu index 2e67a40..5663572 100644 --- a/modules/gpu/src/cuda/hist.cu +++ b/modules/gpu/src/cuda/hist.cu @@ -42,13 +42,13 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/emulation.hpp" -#include "opencv2/core/device/transform.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/emulation.hpp" +#include "opencv2/core/cuda/transform.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace hist { @@ -127,7 +127,7 @@ namespace hist }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { template <> struct TransformFunctorTraits : DefaultTransformFunctorTraits { diff --git a/modules/gpu/src/cuda/hog.cu b/modules/gpu/src/cuda/hog.cu index 34cc80c..6d02007 100644 --- a/modules/gpu/src/cuda/hog.cu +++ b/modules/gpu/src/cuda/hog.cu @@ -42,12 +42,12 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/warp_shuffle.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/warp_shuffle.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { // Other values are not supported #define CELL_WIDTH 8 @@ -808,7 +808,7 @@ namespace cv { namespace gpu { namespace device void resize_8UC1(const PtrStepSzb& src, PtrStepSzb dst) { resize_for_hog (src, dst, resize8UC1_tex); } void resize_8UC4(const PtrStepSzb& src, PtrStepSzb dst) { resize_for_hog(src, dst, resize8UC4_tex); } } // namespace hog -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/hough.cu b/modules/gpu/src/cuda/hough.cu index 1f885fb..d8f8c89 100644 --- a/modules/gpu/src/cuda/hough.cu +++ b/modules/gpu/src/cuda/hough.cu @@ -45,13 +45,13 @@ #include #include -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/emulation.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/dynamic_smem.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/emulation.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/dynamic_smem.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace hough { @@ -1424,7 +1424,7 @@ namespace cv { namespace gpu { namespace device cudaSafeCall( cudaDeviceSynchronize() ); thrust::device_ptr sizesPtr(sizes); - thrust::transform(sizesPtr, sizesPtr + levels + 1, sizesPtr, device::bind2nd(device::minimum(), maxSize)); + thrust::transform(sizesPtr, sizesPtr + levels + 1, sizesPtr, cuda::bind2nd(cuda::minimum(), maxSize)); } void GHT_Guil_Full_buildTemplFeatureList_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount, diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index 0f59a77..0fa52d0 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -42,14 +42,14 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/border_interpolate.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" #include "internal_shared.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -1002,7 +1002,7 @@ namespace cv { namespace gpu { namespace device template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); } // namespace imgproc -}}} // namespace cv { namespace gpu { namespace device { +}}} // namespace cv { namespace gpu { namespace cuda { #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/integral_image.cu b/modules/gpu/src/cuda/integral_image.cu index 9ef3166..3759d04 100644 --- a/modules/gpu/src/cuda/integral_image.cu +++ b/modules/gpu/src/cuda/integral_image.cu @@ -42,9 +42,9 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { diff --git a/modules/gpu/src/cuda/internal_shared.hpp b/modules/gpu/src/cuda/internal_shared.hpp index 6e84611..4cdf816 100644 --- a/modules/gpu/src/cuda/internal_shared.hpp +++ b/modules/gpu/src/cuda/internal_shared.hpp @@ -48,7 +48,7 @@ #include "NPP_staging.hpp" #include "opencv2/gpu/devmem2d.hpp" #include "safe_call.hpp" -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" namespace cv { namespace gpu { diff --git a/modules/gpu/src/cuda/lbp.cu b/modules/gpu/src/cuda/lbp.cu index 87dfbfc..cc63bd7 100644 --- a/modules/gpu/src/cuda/lbp.cu +++ b/modules/gpu/src/cuda/lbp.cu @@ -43,10 +43,10 @@ #if !defined CUDA_DISABLER #include "lbp.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/saturate_cast.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace lbp { diff --git a/modules/gpu/src/cuda/lbp.hpp b/modules/gpu/src/cuda/lbp.hpp index 9cfebd5..1b8ffcd 100644 --- a/modules/gpu/src/cuda/lbp.hpp +++ b/modules/gpu/src/cuda/lbp.hpp @@ -43,10 +43,10 @@ #ifndef __OPENCV_GPU_DEVICE_LBP_HPP_ #define __OPENCV_GPU_DEVICE_LBP_HPP_ -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/emulation.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/emulation.hpp" -namespace cv { namespace gpu { namespace device { +namespace cv { namespace gpu { namespace cuda { namespace lbp { diff --git a/modules/gpu/src/cuda/match_template.cu b/modules/gpu/src/cuda/match_template.cu index 9ad3338..95e7c38 100644 --- a/modules/gpu/src/cuda/match_template.cu +++ b/modules/gpu/src/cuda/match_template.cu @@ -42,10 +42,10 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_math.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_math.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace match_template { @@ -910,7 +910,7 @@ namespace cv { namespace gpu { namespace device cudaSafeCall( cudaDeviceSynchronize() ); } } //namespace match_template -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu index 01ebf3d..ed91811 100644 --- a/modules/gpu/src/cuda/mathfunc.cu +++ b/modules/gpu/src/cuda/mathfunc.cu @@ -42,9 +42,9 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace mathfunc { @@ -212,6 +212,6 @@ namespace cv { namespace gpu { namespace device callers[mag.data == 0](mag, angle, x, y, angleInDegrees, stream); } } // namespace mathfunc -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/matrix_reductions.cu b/modules/gpu/src/cuda/matrix_reductions.cu index d0edea5..0c1f719 100644 --- a/modules/gpu/src/cuda/matrix_reductions.cu +++ b/modules/gpu/src/cuda/matrix_reductions.cu @@ -42,18 +42,18 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/utility.hpp" -#include "opencv2/core/device/type_traits.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/type_traits.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace detail { @@ -205,7 +205,7 @@ namespace detail template static __device__ __forceinline__ thrust::tuple smem_tuple(R* smem) { - return cv::gpu::device::smem_tuple(smem, smem + BLOCK_SIZE); + return cv::gpu::cuda::smem_tuple(smem, smem + BLOCK_SIZE); } template @@ -225,7 +225,7 @@ namespace detail template static __device__ __forceinline__ thrust::tuple smem_tuple(R* smem) { - return cv::gpu::device::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE); + return cv::gpu::cuda::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE); } template @@ -245,7 +245,7 @@ namespace detail template static __device__ __forceinline__ thrust::tuple smem_tuple(R* smem) { - return cv::gpu::device::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE, smem + 3 * BLOCK_SIZE); + return cv::gpu::cuda::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE, smem + 3 * BLOCK_SIZE); } template @@ -340,7 +340,7 @@ namespace sum { sum = tid < gridDim.x * gridDim.y ? result[tid] : VecTraits::all(0); - device::reduce(detail::Unroll::template smem_tuple(smem), detail::Unroll::tie(sum), tid, detail::Unroll::op(plus())); + cuda::reduce(detail::Unroll::template smem_tuple(smem), detail::Unroll::tie(sum), tid, detail::Unroll::op(plus())); if (tid == 0) { @@ -383,7 +383,7 @@ namespace sum } } - device::reduce(detail::Unroll::template smem_tuple(smem), detail::Unroll::tie(sum), tid, detail::Unroll::op(plus())); + cuda::reduce(detail::Unroll::template smem_tuple(smem), detail::Unroll::tie(sum), tid, detail::Unroll::op(plus())); GlobalReduce::run(sum, result, tid, bid, smem); } @@ -642,7 +642,7 @@ namespace minMax const minimum minOp; const maximum maxOp; - device::reduce(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax), tid, thrust::make_tuple(minOp, maxOp)); + cuda::reduce(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax), tid, thrust::make_tuple(minOp, maxOp)); if (tid == 0) { @@ -690,7 +690,7 @@ namespace minMax } } - device::reduce(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax), tid, thrust::make_tuple(minOp, maxOp)); + cuda::reduce(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax), tid, thrust::make_tuple(minOp, maxOp)); GlobalReduce::run(mymin, mymax, minval, maxval, tid, bid, sminval, smaxval); } @@ -994,7 +994,7 @@ namespace countNonZero } } - device::reduce(scount, mycount, tid, plus()); + cuda::reduce(scount, mycount, tid, plus()); #if __CUDA_ARCH__ >= 200 if (tid == 0) @@ -1019,7 +1019,7 @@ namespace countNonZero { mycount = tid < gridDim.x * gridDim.y ? count[tid] : 0; - device::reduce(scount, mycount, tid, plus()); + cuda::reduce(scount, mycount, tid, plus()); if (tid == 0) { @@ -1217,7 +1217,7 @@ namespace reduce volatile S* srow = smem + threadIdx.y * 16; myVal = srow[threadIdx.x]; - device::reduce<16>(srow, myVal, threadIdx.x, op); + cuda::reduce<16>(srow, myVal, threadIdx.x, op); if (threadIdx.x == 0) srow[0] = myVal; @@ -1301,7 +1301,7 @@ namespace reduce for (int x = threadIdx.x; x < src.cols; x += BLOCK_SIZE) myVal = op(myVal, saturate_cast(srcRow[x])); - device::reduce(detail::Unroll::template smem_tuple(smem), detail::Unroll::tie(myVal), threadIdx.x, detail::Unroll::op(op)); + cuda::reduce(detail::Unroll::template smem_tuple(smem), detail::Unroll::tie(myVal), threadIdx.x, detail::Unroll::op(op)); if (threadIdx.x == 0) dst[y] = saturate_cast(op.result(myVal, src.cols)); diff --git a/modules/gpu/src/cuda/nlm.cu b/modules/gpu/src/cuda/nlm.cu index 6fb5dba..fae2495 100644 --- a/modules/gpu/src/cuda/nlm.cu +++ b/modules/gpu/src/cuda/nlm.cu @@ -42,12 +42,12 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/border_interpolate.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" using namespace cv::gpu; @@ -57,7 +57,7 @@ typedef unsigned short ushort; ////////////////////////////////////////////////////////////////////////////////// //// Non Local Means Denosing -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -179,7 +179,7 @@ namespace cv { namespace gpu { namespace device ////////////////////////////////////////////////////////////////////////////////// //// Non Local Means Denosing (fast approximate version) -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -190,7 +190,7 @@ namespace cv { namespace gpu { namespace device template static __device__ __forceinline__ thrust::tuple smem_tuple(float* smem) { - return cv::gpu::device::smem_tuple(smem, smem + BLOCK_SIZE); + return cv::gpu::cuda::smem_tuple(smem, smem + BLOCK_SIZE); } static __device__ __forceinline__ thrust::tuple tie(float& val1, float& val2) @@ -209,7 +209,7 @@ namespace cv { namespace gpu { namespace device template static __device__ __forceinline__ thrust::tuple smem_tuple(float* smem) { - return cv::gpu::device::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE); + return cv::gpu::cuda::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE); } static __device__ __forceinline__ thrust::tuple tie(float& val1, float2& val2) @@ -228,7 +228,7 @@ namespace cv { namespace gpu { namespace device template static __device__ __forceinline__ thrust::tuple smem_tuple(float* smem) { - return cv::gpu::device::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE, smem + 3 * BLOCK_SIZE); + return cv::gpu::cuda::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE, smem + 3 * BLOCK_SIZE); } static __device__ __forceinline__ thrust::tuple tie(float& val1, float3& val2) @@ -247,7 +247,7 @@ namespace cv { namespace gpu { namespace device template static __device__ __forceinline__ thrust::tuple smem_tuple(float* smem) { - return cv::gpu::device::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE, smem + 3 * BLOCK_SIZE, smem + 4 * BLOCK_SIZE); + return cv::gpu::cuda::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE, smem + 3 * BLOCK_SIZE, smem + 4 * BLOCK_SIZE); } static __device__ __forceinline__ thrust::tuple tie(float& val1, float4& val2) diff --git a/modules/gpu/src/cuda/optflowbm.cu b/modules/gpu/src/cuda/optflowbm.cu index 7cc4d2f..9adf210 100644 --- a/modules/gpu/src/cuda/optflowbm.cu +++ b/modules/gpu/src/cuda/optflowbm.cu @@ -42,13 +42,13 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/reduce.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/reduce.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace optflowbm { diff --git a/modules/gpu/src/cuda/optical_flow.cu b/modules/gpu/src/cuda/optical_flow.cu index a80d485..7219b2f 100644 --- a/modules/gpu/src/cuda/optical_flow.cu +++ b/modules/gpu/src/cuda/optical_flow.cu @@ -42,9 +42,9 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace optical_flow { diff --git a/modules/gpu/src/cuda/optical_flow_farneback.cu b/modules/gpu/src/cuda/optical_flow_farneback.cu index f9f50c8..df1662e 100644 --- a/modules/gpu/src/cuda/optical_flow_farneback.cu +++ b/modules/gpu/src/cuda/optical_flow_farneback.cu @@ -42,8 +42,8 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/border_interpolate.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" #define tx threadIdx.x #define ty threadIdx.y @@ -55,7 +55,7 @@ #define BORDER_SIZE 5 #define MAX_KSIZE_HALF 100 -namespace cv { namespace gpu { namespace device { namespace optflow_farneback +namespace cv { namespace gpu { namespace cuda { namespace optflow_farneback { __constant__ float c_g[8]; __constant__ float c_xg[8]; @@ -641,7 +641,7 @@ namespace cv { namespace gpu { namespace device { namespace optflow_farneback callers[borderMode](src, ksizeHalf, dst, stream); } -}}}} // namespace cv { namespace gpu { namespace device { namespace optflow_farneback +}}}} // namespace cv { namespace gpu { namespace cuda { namespace optflow_farneback #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/orb.cu b/modules/gpu/src/cuda/orb.cu index 9c4e0ff..95a9899 100644 --- a/modules/gpu/src/cuda/orb.cu +++ b/modules/gpu/src/cuda/orb.cu @@ -45,11 +45,11 @@ #include #include -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/functional.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/functional.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace orb { diff --git a/modules/gpu/src/cuda/pyr_down.cu b/modules/gpu/src/cuda/pyr_down.cu index 4244250..ce90818 100644 --- a/modules/gpu/src/cuda/pyr_down.cu +++ b/modules/gpu/src/cuda/pyr_down.cu @@ -42,13 +42,13 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/border_interpolate.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/saturate_cast.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -222,7 +222,7 @@ namespace cv { namespace gpu { namespace device template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); } // namespace imgproc -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/pyr_up.cu b/modules/gpu/src/cuda/pyr_up.cu index 52ef35a..821871c 100644 --- a/modules/gpu/src/cuda/pyr_up.cu +++ b/modules/gpu/src/cuda/pyr_up.cu @@ -42,13 +42,13 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/border_interpolate.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/saturate_cast.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -191,6 +191,6 @@ namespace cv { namespace gpu { namespace device template void pyrUp_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); template void pyrUp_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); } // namespace imgproc -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/pyrlk.cu b/modules/gpu/src/cuda/pyrlk.cu index fc387a7..32d8a12 100644 --- a/modules/gpu/src/cuda/pyrlk.cu +++ b/modules/gpu/src/cuda/pyrlk.cu @@ -42,15 +42,15 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/utility.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/reduce.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/reduce.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace pyrlk { diff --git a/modules/gpu/src/cuda/remap.cu b/modules/gpu/src/cuda/remap.cu index 4589834..bc5692c 100644 --- a/modules/gpu/src/cuda/remap.cu +++ b/modules/gpu/src/cuda/remap.cu @@ -42,14 +42,14 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/border_interpolate.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/filters.hpp" - -namespace cv { namespace gpu { namespace device +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/filters.hpp" + +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -268,7 +268,7 @@ namespace cv { namespace gpu { namespace device template void remap_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, PtrStepSzf xmap, PtrStepSzf ymap, PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20); template void remap_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, PtrStepSzf xmap, PtrStepSzf ymap, PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20); } // namespace imgproc -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/resize.cu b/modules/gpu/src/cuda/resize.cu index 86f8d55..9976d5e 100644 --- a/modules/gpu/src/cuda/resize.cu +++ b/modules/gpu/src/cuda/resize.cu @@ -43,15 +43,15 @@ #if !defined CUDA_DISABLER #include -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/border_interpolate.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/filters.hpp" -#include "opencv2/core/device/scan.hpp" - -namespace cv { namespace gpu { namespace device +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/filters.hpp" +#include "opencv2/core/cuda/scan.hpp" + +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -296,7 +296,7 @@ namespace cv { namespace gpu { namespace device }; } // namespace imgproc -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/rgb_to_yv12.cu b/modules/gpu/src/cuda/rgb_to_yv12.cu index 56800fc..f8f5b95 100644 --- a/modules/gpu/src/cuda/rgb_to_yv12.cu +++ b/modules/gpu/src/cuda/rgb_to_yv12.cu @@ -42,10 +42,10 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/vec_traits.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace video_encoding { diff --git a/modules/gpu/src/cuda/row_filter.h b/modules/gpu/src/cuda/row_filter.h index beaee37..254d8d9 100644 --- a/modules/gpu/src/cuda/row_filter.h +++ b/modules/gpu/src/cuda/row_filter.h @@ -40,13 +40,13 @@ // //M*/ -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/border_interpolate.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; namespace row_filter { diff --git a/modules/gpu/src/cuda/split_merge.cu b/modules/gpu/src/cuda/split_merge.cu index d92ecbc..47b6e34 100644 --- a/modules/gpu/src/cuda/split_merge.cu +++ b/modules/gpu/src/cuda/split_merge.cu @@ -42,9 +42,9 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace split_merge { @@ -505,7 +505,7 @@ namespace cv { namespace gpu { namespace device split_func(src, dst, stream); } } // namespace split_merge -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/stereobm.cu b/modules/gpu/src/cuda/stereobm.cu index 10a63ec..5a4cc17 100644 --- a/modules/gpu/src/cuda/stereobm.cu +++ b/modules/gpu/src/cuda/stereobm.cu @@ -42,9 +42,9 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace stereobm { @@ -534,7 +534,7 @@ namespace cv { namespace gpu { namespace device cudaSafeCall( cudaUnbindTexture (texForTF) ); } } // namespace stereobm -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/stereobp.cu b/modules/gpu/src/cuda/stereobp.cu index 2aa70bf..3fe946c 100644 --- a/modules/gpu/src/cuda/stereobp.cu +++ b/modules/gpu/src/cuda/stereobp.cu @@ -42,11 +42,11 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/limits.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/limits.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace stereobp { @@ -384,7 +384,7 @@ namespace cv { namespace gpu { namespace device template __device__ void message(const T* msg1, const T* msg2, const T* msg3, const T* data, T* dst, size_t msg_disp_step, size_t data_disp_step) { - float minimum = device::numeric_limits::max(); + float minimum = cuda::numeric_limits::max(); for(int i = 0; i < cndisp; ++i) { @@ -533,6 +533,6 @@ namespace cv { namespace gpu { namespace device template void output_gpu(const PtrStepSzb& u, const PtrStepSzb& d, const PtrStepSzb& l, const PtrStepSzb& r, const PtrStepSzb& data, const PtrStepSz& disp, cudaStream_t stream); template void output_gpu(const PtrStepSzb& u, const PtrStepSzb& d, const PtrStepSzb& l, const PtrStepSzb& r, const PtrStepSzb& data, const PtrStepSz& disp, cudaStream_t stream); } // namespace stereobp -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/stereocsbp.cu b/modules/gpu/src/cuda/stereocsbp.cu index 717398e..2af8402 100644 --- a/modules/gpu/src/cuda/stereocsbp.cu +++ b/modules/gpu/src/cuda/stereocsbp.cu @@ -42,13 +42,13 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/functional.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/functional.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace stereocsbp { @@ -146,7 +146,7 @@ namespace cv { namespace gpu { namespace device for(int i = 0; i < nr_plane; i++) { - T minimum = device::numeric_limits::max(); + T minimum = cuda::numeric_limits::max(); int id = 0; for(int d = 0; d < cndisp; d++) { @@ -859,6 +859,6 @@ namespace cv { namespace gpu { namespace device template void compute_disp(const float* u, const float* d, const float* l, const float* r, const float* data_cost_selected, const float* disp_selected, size_t msg_step, const PtrStepSz& disp, int nr_plane, cudaStream_t stream); } // namespace stereocsbp -}}} // namespace cv { namespace gpu { namespace device { +}}} // namespace cv { namespace gpu { namespace cuda { #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/texture_binder.hpp b/modules/gpu/src/cuda/texture_binder.hpp index 801a6d2..e967746 100644 --- a/modules/gpu/src/cuda/texture_binder.hpp +++ b/modules/gpu/src/cuda/texture_binder.hpp @@ -83,7 +83,7 @@ namespace cv }; } - namespace device + namespace cuda { using cv::gpu::TextureBinder; } diff --git a/modules/gpu/src/cuda/tvl1flow.cu b/modules/gpu/src/cuda/tvl1flow.cu index 29470fb..8502b7b 100644 --- a/modules/gpu/src/cuda/tvl1flow.cu +++ b/modules/gpu/src/cuda/tvl1flow.cu @@ -42,12 +42,12 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/border_interpolate.hpp" -#include "opencv2/core/device/limits.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" +#include "opencv2/core/cuda/limits.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; //////////////////////////////////////////////////////////// // centeredGradient diff --git a/modules/gpu/src/cuda/warp.cu b/modules/gpu/src/cuda/warp.cu index 4a00644..1073519 100644 --- a/modules/gpu/src/cuda/warp.cu +++ b/modules/gpu/src/cuda/warp.cu @@ -42,14 +42,14 @@ #if !defined CUDA_DISABLER -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/border_interpolate.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/filters.hpp" - -namespace cv { namespace gpu { namespace device +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/filters.hpp" + +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -383,7 +383,7 @@ namespace cv { namespace gpu { namespace device template void warpPerspective_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[3 * 3], PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20); template void warpPerspective_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[3 * 3], PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20); } // namespace imgproc -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cvt_color_internal.h b/modules/gpu/src/cvt_color_internal.h index 1b7c68f..7e3bab1 100644 --- a/modules/gpu/src/cvt_color_internal.h +++ b/modules/gpu/src/cvt_color_internal.h @@ -43,7 +43,7 @@ #ifndef __cvt_color_internal_h__ #define __cvt_color_internal_h__ -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { #define OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name) \ void name(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); diff --git a/modules/gpu/src/denoising.cpp b/modules/gpu/src/denoising.cpp index 995b809..d3bcae2 100644 --- a/modules/gpu/src/denoising.cpp +++ b/modules/gpu/src/denoising.cpp @@ -59,7 +59,7 @@ void cv::gpu::FastNonLocalMeansDenoising::labMethod( const GpuMat&, GpuMat&, flo ////////////////////////////////////////////////////////////////////////////////// //// Non Local Means Denosing (brute force) -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -73,7 +73,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::bilateralFilter(const GpuMat& src, GpuMat& dst, int kernel_size, float sigma_color, float sigma_spatial, int borderMode, Stream& s) { - using cv::gpu::device::imgproc::bilateral_filter_gpu; + using cv::gpu::cuda::imgproc::bilateral_filter_gpu; typedef void (*func_t)(const PtrStepSzb& src, PtrStepSzb dst, int kernel_size, float sigma_spatial, float sigma_color, int borderMode, cudaStream_t s); @@ -109,7 +109,7 @@ void cv::gpu::bilateralFilter(const GpuMat& src, GpuMat& dst, int kernel_size, f void cv::gpu::nonLocalMeans(const GpuMat& src, GpuMat& dst, float h, int search_window, int block_window, int borderMode, Stream& s) { - using cv::gpu::device::imgproc::nlm_bruteforce_gpu; + using cv::gpu::cuda::imgproc::nlm_bruteforce_gpu; typedef void (*func_t)(const PtrStepSzb& src, PtrStepSzb dst, int search_radius, int block_radius, float h, int borderMode, cudaStream_t stream); static const func_t funcs[4] = { nlm_bruteforce_gpu, nlm_bruteforce_gpu, nlm_bruteforce_gpu, 0/*nlm_bruteforce_gpu,*/ }; @@ -134,7 +134,7 @@ void cv::gpu::nonLocalMeans(const GpuMat& src, GpuMat& dst, float h, int search_ //// Non Local Means Denosing (fast approxinate) -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -163,10 +163,10 @@ void cv::gpu::FastNonLocalMeansDenoising::simpleMethod(const GpuMat& src, GpuMat GpuMat src_hdr = extended_src(Rect(Point2i(border_size, border_size), src.size())); int bcols, brows; - device::imgproc::nln_fast_get_buffer_size(src_hdr, search_window, block_window, bcols, brows); + cuda::imgproc::nln_fast_get_buffer_size(src_hdr, search_window, block_window, bcols, brows); buffer.create(brows, bcols, CV_32S); - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; typedef void (*nlm_fast_t)(const PtrStepSzb&, PtrStepSzb, PtrStepi, int, int, float, cudaStream_t); static const nlm_fast_t funcs[] = { nlm_fast_gpu, nlm_fast_gpu, nlm_fast_gpu, 0}; @@ -183,12 +183,12 @@ void cv::gpu::FastNonLocalMeansDenoising::labMethod( const GpuMat& src, GpuMat& l.create(src.size(), CV_8U); ab.create(src.size(), CV_8UC2); - device::imgproc::fnlm_split_channels(lab, l, ab, StreamAccessor::getStream(s)); + cuda::imgproc::fnlm_split_channels(lab, l, ab, StreamAccessor::getStream(s)); simpleMethod(l, l, h_luminance, search_window, block_window, s); simpleMethod(ab, ab, h_color, search_window, block_window, s); - device::imgproc::fnlm_merge_channels(l, ab, lab, StreamAccessor::getStream(s)); + cuda::imgproc::fnlm_merge_channels(l, ab, lab, StreamAccessor::getStream(s)); cv::gpu::cvtColor(lab, dst, CV_Lab2BGR, 0, s); } diff --git a/modules/gpu/src/fast.cpp b/modules/gpu/src/fast.cpp index 01ddb75..13a719b 100644 --- a/modules/gpu/src/fast.cpp +++ b/modules/gpu/src/fast.cpp @@ -108,7 +108,7 @@ void cv::gpu::FAST_GPU::operator ()(const GpuMat& img, const GpuMat& mask, GpuMa keypoints.cols = getKeyPoints(keypoints); } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace fast { @@ -119,7 +119,7 @@ namespace cv { namespace gpu { namespace device int cv::gpu::FAST_GPU::calcKeyPointsLocation(const GpuMat& img, const GpuMat& mask) { - using namespace cv::gpu::device::fast; + using namespace cv::gpu::cuda::fast; CV_Assert(img.type() == CV_8UC1); CV_Assert(mask.empty() || (mask.type() == CV_8UC1 && mask.size() == img.size())); @@ -142,7 +142,7 @@ int cv::gpu::FAST_GPU::calcKeyPointsLocation(const GpuMat& img, const GpuMat& ma int cv::gpu::FAST_GPU::getKeyPoints(GpuMat& keypoints) { - using namespace cv::gpu::device::fast; + using namespace cv::gpu::cuda::fast; if (count_ == 0) return 0; diff --git a/modules/gpu/src/filtering.cpp b/modules/gpu/src/filtering.cpp index 8905eae..989f735 100644 --- a/modules/gpu/src/filtering.cpp +++ b/modules/gpu/src/filtering.cpp @@ -661,7 +661,7 @@ void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& ke //////////////////////////////////////////////////////////////////////////////////////////////////// // Linear Filter -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -723,7 +723,7 @@ namespace virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; Point ofs; Size wholeSize; @@ -742,7 +742,7 @@ namespace Ptr cv::gpu::getLinearFilter_GPU(int srcType, int dstType, const Mat& kernel, Point anchor, int brd_type) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; int sdepth = CV_MAT_DEPTH(srcType); int scn = CV_MAT_CN(srcType); diff --git a/modules/gpu/src/gftt.cpp b/modules/gpu/src/gftt.cpp index 76aee56..9ed4f77 100644 --- a/modules/gpu/src/gftt.cpp +++ b/modules/gpu/src/gftt.cpp @@ -51,7 +51,7 @@ void cv::gpu::GoodFeaturesToTrackDetector_GPU::operator ()(const GpuMat&, GpuMat #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace gfft { @@ -62,7 +62,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::GoodFeaturesToTrackDetector_GPU::operator ()(const GpuMat& image, GpuMat& corners, const GpuMat& mask) { - using namespace cv::gpu::device::gfft; + using namespace cv::gpu::cuda::gfft; CV_Assert(qualityLevel > 0 && minDistance >= 0 && maxCorners >= 0); CV_Assert(mask.empty() || (mask.type() == CV_8UC1 && mask.size() == image.size())); diff --git a/modules/gpu/src/global_motion.cpp b/modules/gpu/src/global_motion.cpp index f7da2e8..010bbd9 100644 --- a/modules/gpu/src/global_motion.cpp +++ b/modules/gpu/src/global_motion.cpp @@ -53,7 +53,7 @@ void cv::gpu::calcWobbleSuppressionMaps( #else -namespace cv { namespace gpu { namespace device { namespace globmotion { +namespace cv { namespace gpu { namespace cuda { namespace globmotion { int compactPoints(int N, float *points0, float *points1, const uchar *mask); @@ -70,7 +70,7 @@ void cv::gpu::compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask CV_Assert(points0.cols == mask.cols && points1.cols == mask.cols); int npoints = points0.cols; - int remaining = cv::gpu::device::globmotion::compactPoints( + int remaining = cv::gpu::cuda::globmotion::compactPoints( npoints, (float*)points0.data, (float*)points1.data, mask.data); points0 = points0.colRange(0, remaining); @@ -88,7 +88,7 @@ void cv::gpu::calcWobbleSuppressionMaps( mapx.create(size, CV_32F); mapy.create(size, CV_32F); - cv::gpu::device::globmotion::calcWobbleSuppressionMaps( + cv::gpu::cuda::globmotion::calcWobbleSuppressionMaps( left, idx, right, size.width, size.height, ml.ptr(), mr.ptr(), mapx, mapy); } diff --git a/modules/gpu/src/graphcuts.cpp b/modules/gpu/src/graphcuts.cpp index b8c0d4f..0513845 100644 --- a/modules/gpu/src/graphcuts.cpp +++ b/modules/gpu/src/graphcuts.cpp @@ -52,7 +52,7 @@ void cv::gpu::labelComponents(const GpuMat&, GpuMat&, int, Stream&) { throw_nogp #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace ccl { @@ -81,12 +81,12 @@ void cv::gpu::connectivityMask(const GpuMat& image, GpuMat& mask, const cv::Scal static const func_t suppotLookup[8][4] = { // 1, 2, 3, 4 - { device::ccl::computeEdges, 0, device::ccl::computeEdges, device::ccl::computeEdges },// CV_8U + { cuda::ccl::computeEdges, 0, cuda::ccl::computeEdges, cuda::ccl::computeEdges },// CV_8U { 0, 0, 0, 0 },// CV_16U - { device::ccl::computeEdges, 0, device::ccl::computeEdges, device::ccl::computeEdges },// CV_8S + { cuda::ccl::computeEdges, 0, cuda::ccl::computeEdges, cuda::ccl::computeEdges },// CV_8S { 0, 0, 0, 0 },// CV_16S - { device::ccl::computeEdges, 0, 0, 0 },// CV_32S - { device::ccl::computeEdges, 0, 0, 0 },// CV_32F + { cuda::ccl::computeEdges, 0, 0, 0 },// CV_32S + { cuda::ccl::computeEdges, 0, 0, 0 },// CV_32F { 0, 0, 0, 0 },// CV_64F { 0, 0, 0, 0 } // CV_USRTYPE1 }; @@ -112,7 +112,7 @@ void cv::gpu::labelComponents(const GpuMat& mask, GpuMat& components, int flags, components.create(mask.size(), CV_32SC1); cudaStream_t stream = StreamAccessor::getStream(s); - device::ccl::labelComponents(mask, components, flags, stream); + cuda::ccl::labelComponents(mask, components, flags, stream); } namespace diff --git a/modules/gpu/src/hog.cpp b/modules/gpu/src/hog.cpp index 03ca0a5..b14c69b 100644 --- a/modules/gpu/src/hog.cpp +++ b/modules/gpu/src/hog.cpp @@ -62,7 +62,7 @@ void cv::gpu::HOGDescriptor::computeConfidenceMultiScale(const GpuMat&, std::vec #else -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace hog { @@ -102,7 +102,7 @@ namespace cv { namespace gpu { namespace device } }}} -using namespace ::cv::gpu::device; +using namespace ::cv::gpu::cuda; cv::gpu::HOGDescriptor::HOGDescriptor(Size win_size_, Size block_size_, Size block_stride_, Size cell_size_, int nbins_, double win_sigma_, double threshold_L2hys_, bool gamma_correction_, int nlevels_) diff --git a/modules/gpu/src/hough.cpp b/modules/gpu/src/hough.cpp index 610384e..4a66b32 100644 --- a/modules/gpu/src/hough.cpp +++ b/modules/gpu/src/hough.cpp @@ -70,7 +70,7 @@ void cv::gpu::GeneralizedHough_GPU::release() {} #include "opencv2/core/utility.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace hough { @@ -81,7 +81,7 @@ namespace cv { namespace gpu { namespace device ////////////////////////////////////////////////////////// // HoughLines -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace hough { @@ -98,7 +98,7 @@ void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, float rho, float thet void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int threshold, bool doSort, int maxLines) { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(src.type() == CV_8UC1); CV_Assert(src.cols < std::numeric_limits::max()); @@ -161,7 +161,7 @@ void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, Ou ////////////////////////////////////////////////////////// // HoughLinesP -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace hough { @@ -171,7 +171,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::HoughLinesP(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int minLineLength, int maxLineGap, int maxLines) { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert( src.type() == CV_8UC1 ); CV_Assert( src.cols < std::numeric_limits::max() ); @@ -210,7 +210,7 @@ void cv::gpu::HoughLinesP(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, ////////////////////////////////////////////////////////// // HoughCircles -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace hough { @@ -230,7 +230,7 @@ void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, int method, float void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, HoughCirclesBuf& buf, int method, float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles) { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(src.type() == CV_8UC1); CV_Assert(src.cols < std::numeric_limits::max()); @@ -371,7 +371,7 @@ void cv::gpu::HoughCirclesDownload(const GpuMat& d_circles, cv::OutputArray h_ci ////////////////////////////////////////////////////////// // GeneralizedHough -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace hough { @@ -559,7 +559,7 @@ namespace void GHT_Pos::buildEdgePointList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy) { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; typedef int (*func_t)(PtrStepSzb edges, PtrStepSzb dx, PtrStepSzb dy, unsigned int* coordList, float* thetaList); static const func_t funcs[] = @@ -747,7 +747,7 @@ namespace void GHT_Ballard_Pos::processTempl() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(levels > 0); @@ -773,7 +773,7 @@ namespace void GHT_Ballard_Pos::calcHist() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(levels > 0 && r_table.rows == (levels + 1) && r_sizes.cols == (levels + 1)); CV_Assert(dp > 0.0); @@ -796,7 +796,7 @@ namespace void GHT_Ballard_Pos::findPosInHist() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(votesThreshold > 0); @@ -851,7 +851,7 @@ namespace void GHT_Ballard_PosScale::calcHist() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(levels > 0 && r_table.rows == (levels + 1) && r_sizes.cols == (levels + 1)); CV_Assert(dp > 0.0); @@ -879,7 +879,7 @@ namespace void GHT_Ballard_PosScale::findPosInHist() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(votesThreshold > 0); @@ -939,7 +939,7 @@ namespace void GHT_Ballard_PosRotation::calcHist() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(levels > 0 && r_table.rows == (levels + 1) && r_sizes.cols == (levels + 1)); CV_Assert(dp > 0.0); @@ -967,7 +967,7 @@ namespace void GHT_Ballard_PosRotation::findPosInHist() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(votesThreshold > 0); @@ -1146,7 +1146,7 @@ namespace void GHT_Guil_Full::processTempl() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; buildFeatureList(templEdges, templDx, templDy, templFeatures, GHT_Guil_Full_setTemplFeatures, GHT_Guil_Full_buildTemplFeatureList_gpu, @@ -1159,7 +1159,7 @@ namespace void GHT_Guil_Full::processImage() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; CV_Assert(levels > 0); CV_Assert(templFeatures.sizes.cols == levels + 1); @@ -1271,7 +1271,7 @@ namespace void GHT_Guil_Full::calcOrientation() { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; const double iAngleStep = 1.0 / angleStep; const int angleRange = cvCeil((maxAngle - minAngle) * iAngleStep); @@ -1295,7 +1295,7 @@ namespace void GHT_Guil_Full::calcScale(double angle) { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; const double iScaleStep = 1.0 / scaleStep; const int scaleRange = cvCeil((maxScale - minScale) * iScaleStep); @@ -1319,7 +1319,7 @@ namespace void GHT_Guil_Full::calcPosition(double angle, int angleVotes, double scale, int scaleVotes) { - using namespace cv::gpu::device::hough; + using namespace cv::gpu::cuda::hough; hist.setTo(Scalar::all(0)); GHT_Guil_Full_calcPHist_gpu(templFeatures.sizes.ptr(), imageFeatures.sizes.ptr(0), diff --git a/modules/gpu/src/imgproc.cpp b/modules/gpu/src/imgproc.cpp index 916dde0..ee2887c 100644 --- a/modules/gpu/src/imgproc.cpp +++ b/modules/gpu/src/imgproc.cpp @@ -98,7 +98,7 @@ cv::Ptr cv::gpu::createCLAHE(double, cv::Size) { throw_nogpu(); //////////////////////////////////////////////////////////////////////// // meanShiftFiltering_GPU -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -108,7 +108,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria, Stream& stream) { - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; if( src.empty() ) CV_Error( CV_StsBadArg, "The input image is empty" ); @@ -134,7 +134,7 @@ void cv::gpu::meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, //////////////////////////////////////////////////////////////////////// // meanShiftProc_GPU -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -144,7 +144,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int sp, int sr, TermCriteria criteria, Stream& stream) { - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; if( src.empty() ) CV_Error( CV_StsBadArg, "The input image is empty" ); @@ -171,7 +171,7 @@ void cv::gpu::meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int //////////////////////////////////////////////////////////////////////// // drawColorDisp -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -185,7 +185,7 @@ namespace template void drawColorDisp_caller(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream) { - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; dst.create(src.size(), CV_8UC4); @@ -207,7 +207,7 @@ void cv::gpu::drawColorDisp(const GpuMat& src, GpuMat& dst, int ndisp, Stream& s //////////////////////////////////////////////////////////////////////// // reprojectImageTo3D -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -218,7 +218,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q, int dst_cn, Stream& stream) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; typedef void (*func_t)(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); static const func_t funcs[2][4] = @@ -239,7 +239,7 @@ void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q, //////////////////////////////////////////////////////////////////////// // copyMakeBorder -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -251,7 +251,7 @@ namespace { template void copyMakeBorder_caller(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream) { - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; Scalar_ val(saturate_cast(value[0]), saturate_cast(value[1]), saturate_cast(value[2]), saturate_cast(value[3])); @@ -348,7 +348,7 @@ void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom ////////////////////////////////////////////////////////////////////////////// // buildWarpPlaneMaps -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -362,7 +362,7 @@ void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, cons float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream) { (void)src_size; - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); @@ -375,14 +375,14 @@ void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, cons map_x.create(dst_roi.size(), CV_32F); map_y.create(dst_roi.size(), CV_32F); - device::imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), + cuda::imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), T.ptr(), scale, StreamAccessor::getStream(stream)); } ////////////////////////////////////////////////////////////////////////////// // buildWarpCylyndricalMaps -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -396,7 +396,7 @@ void cv::gpu::buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat &K GpuMat& map_x, GpuMat& map_y, Stream& stream) { (void)src_size; - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); @@ -408,14 +408,14 @@ void cv::gpu::buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat &K map_x.create(dst_roi.size(), CV_32F); map_y.create(dst_roi.size(), CV_32F); - device::imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), scale, StreamAccessor::getStream(stream)); + cuda::imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), scale, StreamAccessor::getStream(stream)); } ////////////////////////////////////////////////////////////////////////////// // buildWarpSphericalMaps -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -429,7 +429,7 @@ void cv::gpu::buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K, GpuMat& map_x, GpuMat& map_y, Stream& stream) { (void)src_size; - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); @@ -441,7 +441,7 @@ void cv::gpu::buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K, map_x.create(dst_roi.size(), CV_32F); map_y.create(dst_roi.size(), CV_32F); - device::imgproc::buildWarpSphericalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), scale, StreamAccessor::getStream(stream)); + cuda::imgproc::buildWarpSphericalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), scale, StreamAccessor::getStream(stream)); } //////////////////////////////////////////////////////////////////////// @@ -532,7 +532,7 @@ void cv::gpu::integral(const GpuMat& src, GpuMat& sum, Stream& s) integralBuffered(src, sum, buffer, s); } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -556,7 +556,7 @@ void cv::gpu::integralBuffered(const GpuMat& src, GpuMat& sum, GpuMat& buffer, S { ensureSizeIsEnough(((src.rows + 7) / 8) * 8, ((src.cols + 63) / 64) * 64, CV_32SC1, buffer); - cv::gpu::device::imgproc::shfl_integral_gpu(src, buffer, stream); + cv::gpu::cuda::imgproc::shfl_integral_gpu(src, buffer, stream); sum.create(src.rows + 1, src.cols + 1, CV_32SC1); if (s) @@ -631,7 +631,7 @@ void cv::gpu::sqrIntegral(const GpuMat& src, GpuMat& sqsum, Stream& s) ////////////////////////////////////////////////////////////////////////////// // columnSum -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -641,13 +641,13 @@ namespace cv { namespace gpu { namespace device void cv::gpu::columnSum(const GpuMat& src, GpuMat& dst) { - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; CV_Assert(src.type() == CV_32F); dst.create(src.size(), CV_32F); - device::imgproc::columnSum_32F(src, dst); + cuda::imgproc::columnSum_32F(src, dst); } void cv::gpu::rectStdDev(const GpuMat& src, const GpuMat& sqr, GpuMat& dst, const Rect& rect, Stream& s) @@ -1021,7 +1021,7 @@ void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, GpuMat& //////////////////////////////////////////////////////////////////////// // cornerHarris & minEgenVal -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -1074,7 +1074,7 @@ void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& D void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, GpuMat& buf, int blockSize, int ksize, double k, int borderType, Stream& stream) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; CV_Assert(borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT); @@ -1102,7 +1102,7 @@ void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuM void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, GpuMat& buf, int blockSize, int ksize, int borderType, Stream& stream) { - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; CV_Assert(borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT); @@ -1119,7 +1119,7 @@ void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuM ////////////////////////////////////////////////////////////////////////////// // mulSpectrums -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -1132,11 +1132,11 @@ namespace cv { namespace gpu { namespace device void cv::gpu::mulSpectrums(const GpuMat& a, const GpuMat& b, GpuMat& c, int flags, bool conjB, Stream& stream) { (void)flags; - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; typedef void (*Caller)(const PtrStep, const PtrStep, PtrStepSz, cudaStream_t stream); - static Caller callers[] = { device::imgproc::mulSpectrums, device::imgproc::mulSpectrums_CONJ }; + static Caller callers[] = { cuda::imgproc::mulSpectrums, cuda::imgproc::mulSpectrums_CONJ }; CV_Assert(a.type() == b.type() && a.type() == CV_32FC2); CV_Assert(a.size() == b.size()); @@ -1150,7 +1150,7 @@ void cv::gpu::mulSpectrums(const GpuMat& a, const GpuMat& b, GpuMat& c, int flag ////////////////////////////////////////////////////////////////////////////// // mulAndScaleSpectrums -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -1163,10 +1163,10 @@ namespace cv { namespace gpu { namespace device void cv::gpu::mulAndScaleSpectrums(const GpuMat& a, const GpuMat& b, GpuMat& c, int flags, float scale, bool conjB, Stream& stream) { (void)flags; - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; typedef void (*Caller)(const PtrStep, const PtrStep, float scale, PtrStepSz, cudaStream_t stream); - static Caller callers[] = { device::imgproc::mulAndScaleSpectrums, device::imgproc::mulAndScaleSpectrums_CONJ }; + static Caller callers[] = { cuda::imgproc::mulAndScaleSpectrums, cuda::imgproc::mulAndScaleSpectrums_CONJ }; CV_Assert(a.type() == b.type() && a.type() == CV_32FC2); CV_Assert(a.size() == b.size()); @@ -1335,7 +1335,7 @@ void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, bool ccorr, ConvolveBuf& buf, Stream& stream) { - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; #ifndef HAVE_CUFFT throw_nogpu(); diff --git a/modules/gpu/src/match_template.cpp b/modules/gpu/src/match_template.cpp index 26dcd4b..7115b00 100644 --- a/modules/gpu/src/match_template.cpp +++ b/modules/gpu/src/match_template.cpp @@ -51,7 +51,7 @@ void cv::gpu::matchTemplate(const GpuMat&, const GpuMat&, GpuMat&, int, Stream&) #else -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace match_template { @@ -137,7 +137,7 @@ namespace cv { namespace gpu { namespace device } }}} -using namespace ::cv::gpu::device::match_template; +using namespace ::cv::gpu::cuda::match_template; namespace { diff --git a/modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu b/modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu index 881d126..986bfd5 100644 --- a/modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu +++ b/modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu @@ -65,7 +65,7 @@ #include "NPP_staging/NPP_staging.hpp" #include "NCVBroxOpticalFlow.hpp" -#include "opencv2/core/device/utility.hpp" +#include "opencv2/core/cuda/utility.hpp" typedef NCVVectorAlloc FloatVector; @@ -1141,8 +1141,8 @@ NCVStatus NCVBroxOpticalFlow(const NCVBroxOpticalFlowDescriptor desc, ScaleVector(ptrVNew->ptr(), ptrVNew->ptr(), 1.0f/scale_factor, ns * nh, stream); ncvAssertCUDALastErrorReturn((int)NCV_CUDA_ERROR); - cv::gpu::device::swap(ptrU, ptrUNew); - cv::gpu::device::swap(ptrV, ptrVNew); + cv::gpu::cuda::swap(ptrU, ptrUNew); + cv::gpu::cuda::swap(ptrV, ptrVNew); } scale /= scale_factor; } diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu b/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu index 4ff632b..fae2b88 100644 --- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu +++ b/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu @@ -66,8 +66,8 @@ #include "NPP_staging/NPP_staging.hpp" #include "NCVRuntimeTemplates.hpp" #include "NCVHaarObjectDetection.hpp" -#include "opencv2/core/device/warp.hpp" -#include "opencv2/core/device/warp_shuffle.hpp" +#include "opencv2/core/cuda/warp.hpp" +#include "opencv2/core/cuda/warp_shuffle.hpp" //============================================================================== @@ -85,13 +85,13 @@ NCV_CT_ASSERT(K_WARP_SIZE == 32); //this is required for the manual unroll of th __device__ Ncv32u warpScanInclusive(Ncv32u idata, volatile Ncv32u *s_Data) { #if __CUDA_ARCH__ >= 300 - const unsigned int laneId = cv::gpu::device::Warp::laneId(); + const unsigned int laneId = cv::gpu::cuda::Warp::laneId(); // scan on shuffl functions #pragma unroll for (int i = 1; i <= (K_WARP_SIZE / 2); i *= 2) { - const Ncv32u n = cv::gpu::device::shfl_up(idata, i); + const Ncv32u n = cv::gpu::cuda::shfl_up(idata, i); if (laneId >= i) idata += n; } diff --git a/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu b/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu index 8d89817..2830dca 100644 --- a/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu +++ b/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu @@ -45,8 +45,8 @@ #include #include #include "NPP_staging.hpp" -#include "opencv2/core/device/warp.hpp" -#include "opencv2/core/device/warp_shuffle.hpp" +#include "opencv2/core/cuda/warp.hpp" +#include "opencv2/core/cuda/warp_shuffle.hpp" texture tex8u; @@ -95,13 +95,13 @@ template inline __device__ T warpScanInclusive(T idata, volatile T *s_Data) { #if __CUDA_ARCH__ >= 300 - const unsigned int laneId = cv::gpu::device::Warp::laneId(); + const unsigned int laneId = cv::gpu::cuda::Warp::laneId(); // scan on shuffl functions #pragma unroll for (int i = 1; i <= (K_WARP_SIZE / 2); i *= 2) { - const T n = cv::gpu::device::shfl_up(idata, i); + const T n = cv::gpu::cuda::shfl_up(idata, i); if (laneId >= i) idata += n; } diff --git a/modules/gpu/src/nvidia/core/NCVPyramid.cu b/modules/gpu/src/nvidia/core/NCVPyramid.cu index 89d1d4b..6eef5de 100644 --- a/modules/gpu/src/nvidia/core/NCVPyramid.cu +++ b/modules/gpu/src/nvidia/core/NCVPyramid.cu @@ -48,7 +48,7 @@ #include "NCVAlg.hpp" #include "NCVPyramid.hpp" #include "NCVPixelOperations.hpp" -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" template struct __average4_CN {static __host__ __device__ T _average4_CN(const T &p00, const T &p01, const T &p10, const T &p11);}; @@ -204,7 +204,7 @@ __global__ void kernelDownsampleX2(T *d_src, } } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace pyramid { @@ -279,7 +279,7 @@ __global__ void kernelInterpolateFrom1(T *d_srcTop, d_dst_line[j] = outPix; } } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace pyramid { diff --git a/modules/gpu/src/optical_flow.cpp b/modules/gpu/src/optical_flow.cpp index a06bbac..662c7ad 100644 --- a/modules/gpu/src/optical_flow.cpp +++ b/modules/gpu/src/optical_flow.cpp @@ -188,7 +188,7 @@ void cv::gpu::interpolateFrames(const GpuMat& frame0, const GpuMat& frame1, cons cudaSafeCall( cudaDeviceSynchronize() ); } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace optical_flow { @@ -199,7 +199,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::createOpticalFlowNeedleMap(const GpuMat& u, const GpuMat& v, GpuMat& vertex, GpuMat& colors) { - using namespace cv::gpu::device::optical_flow; + using namespace cv::gpu::cuda::optical_flow; CV_Assert(u.type() == CV_32FC1); CV_Assert(v.type() == u.type() && v.size() == u.size()); diff --git a/modules/gpu/src/optical_flow_farneback.cpp b/modules/gpu/src/optical_flow_farneback.cpp index 349ca0c..1f71486 100644 --- a/modules/gpu/src/optical_flow_farneback.cpp +++ b/modules/gpu/src/optical_flow_farneback.cpp @@ -59,7 +59,7 @@ void cv::gpu::FarnebackOpticalFlow::operator ()(const GpuMat&, const GpuMat&, Gp #else -namespace cv { namespace gpu { namespace device { namespace optflow_farneback +namespace cv { namespace gpu { namespace cuda { namespace optflow_farneback { void setPolynomialExpansionConsts( int polyN, const float *g, const float *xg, const float *xxg, @@ -93,7 +93,7 @@ namespace cv { namespace gpu { namespace device { namespace optflow_farneback void gaussianBlur5Gpu_CC11( const PtrStepSzf src, int ksizeHalf, PtrStepSzf dst, int borderType, cudaStream_t stream); -}}}} // namespace cv { namespace gpu { namespace device { namespace optflow_farneback +}}}} // namespace cv { namespace gpu { namespace cuda { namespace optflow_farneback void cv::gpu::FarnebackOpticalFlow::prepareGaussian( @@ -163,7 +163,7 @@ void cv::gpu::FarnebackOpticalFlow::setPolynomialExpansionConsts(int n, double s double ig11, ig03, ig33, ig55; prepareGaussian(n, sigma, g, xg, xxg, ig11, ig03, ig33, ig55); - device::optflow_farneback::setPolynomialExpansionConsts(n, g, xg, xxg, static_cast(ig11), static_cast(ig03), static_cast(ig33), static_cast(ig55)); + cuda::optflow_farneback::setPolynomialExpansionConsts(n, g, xg, xxg, static_cast(ig11), static_cast(ig03), static_cast(ig33), static_cast(ig55)); } @@ -172,17 +172,17 @@ void cv::gpu::FarnebackOpticalFlow::updateFlow_boxFilter( GpuMat& M, GpuMat &bufM, int blockSize, bool updateMatrices, Stream streams[]) { if (deviceSupports(FEATURE_SET_COMPUTE_12)) - device::optflow_farneback::boxFilter5Gpu(M, blockSize/2, bufM, S(streams[0])); + cuda::optflow_farneback::boxFilter5Gpu(M, blockSize/2, bufM, S(streams[0])); else - device::optflow_farneback::boxFilter5Gpu_CC11(M, blockSize/2, bufM, S(streams[0])); + cuda::optflow_farneback::boxFilter5Gpu_CC11(M, blockSize/2, bufM, S(streams[0])); swap(M, bufM); for (int i = 1; i < 5; ++i) streams[i].waitForCompletion(); - device::optflow_farneback::updateFlowGpu(M, flowx, flowy, S(streams[0])); + cuda::optflow_farneback::updateFlowGpu(M, flowx, flowy, S(streams[0])); if (updateMatrices) - device::optflow_farneback::updateMatricesGpu(flowx, flowy, R0, R1, M, S(streams[0])); + cuda::optflow_farneback::updateMatricesGpu(flowx, flowy, R0, R1, M, S(streams[0])); } @@ -191,17 +191,17 @@ void cv::gpu::FarnebackOpticalFlow::updateFlow_gaussianBlur( GpuMat& M, GpuMat &bufM, int blockSize, bool updateMatrices, Stream streams[]) { if (deviceSupports(FEATURE_SET_COMPUTE_12)) - device::optflow_farneback::gaussianBlur5Gpu( + cuda::optflow_farneback::gaussianBlur5Gpu( M, blockSize/2, bufM, BORDER_REPLICATE_GPU, S(streams[0])); else - device::optflow_farneback::gaussianBlur5Gpu_CC11( + cuda::optflow_farneback::gaussianBlur5Gpu_CC11( M, blockSize/2, bufM, BORDER_REPLICATE_GPU, S(streams[0])); swap(M, bufM); - device::optflow_farneback::updateFlowGpu(M, flowx, flowy, S(streams[0])); + cuda::optflow_farneback::updateFlowGpu(M, flowx, flowy, S(streams[0])); if (updateMatrices) - device::optflow_farneback::updateMatricesGpu(flowx, flowy, R0, R1, M, S(streams[0])); + cuda::optflow_farneback::updateMatricesGpu(flowx, flowy, R0, R1, M, S(streams[0])); } @@ -253,7 +253,7 @@ void cv::gpu::FarnebackOpticalFlow::operator ()( } setPolynomialExpansionConsts(polyN, polySigma); - device::optflow_farneback::setUpdateMatricesConsts(); + cuda::optflow_farneback::setUpdateMatricesConsts(); for (int k = numLevelsCropped; k >= 0; k--) { @@ -344,8 +344,8 @@ void cv::gpu::FarnebackOpticalFlow::operator ()( if (fastPyramids) { - device::optflow_farneback::polynomialExpansionGpu(pyramid0_[k], polyN, R[0], S(streams[0])); - device::optflow_farneback::polynomialExpansionGpu(pyramid1_[k], polyN, R[1], S(streams[1])); + cuda::optflow_farneback::polynomialExpansionGpu(pyramid0_[k], polyN, R[0], S(streams[0])); + cuda::optflow_farneback::polynomialExpansionGpu(pyramid1_[k], polyN, R[1], S(streams[1])); } else { @@ -361,11 +361,11 @@ void cv::gpu::FarnebackOpticalFlow::operator ()( }; Mat g = getGaussianKernel(smoothSize, sigma, CV_32F); - device::optflow_farneback::setGaussianBlurKernel(g.ptr(smoothSize/2), smoothSize/2); + cuda::optflow_farneback::setGaussianBlurKernel(g.ptr(smoothSize/2), smoothSize/2); for (int i = 0; i < 2; i++) { - device::optflow_farneback::gaussianBlurGpu( + cuda::optflow_farneback::gaussianBlurGpu( frames_[i], smoothSize/2, blurredFrame[i], BORDER_REFLECT101_GPU, S(streams[i])); #if ENABLE_GPU_RESIZE resize(blurredFrame[i], pyrLevel[i], Size(width, height), INTER_LINEAR, streams[i]); @@ -375,17 +375,17 @@ void cv::gpu::FarnebackOpticalFlow::operator ()( resize(tmp1, tmp2, Size(width, height), INTER_LINEAR); I[i].upload(tmp2); #endif - device::optflow_farneback::polynomialExpansionGpu(pyrLevel[i], polyN, R[i], S(streams[i])); + cuda::optflow_farneback::polynomialExpansionGpu(pyrLevel[i], polyN, R[i], S(streams[i])); } } streams[1].waitForCompletion(); - device::optflow_farneback::updateMatricesGpu(curFlowX, curFlowY, R[0], R[1], M, S(streams[0])); + cuda::optflow_farneback::updateMatricesGpu(curFlowX, curFlowY, R[0], R[1], M, S(streams[0])); if (flags & OPTFLOW_FARNEBACK_GAUSSIAN) { Mat g = getGaussianKernel(winSize, winSize/2*0.3f, CV_32F); - device::optflow_farneback::setGaussianBlurKernel(g.ptr(winSize/2), winSize/2); + cuda::optflow_farneback::setGaussianBlurKernel(g.ptr(winSize/2), winSize/2); } for (int i = 0; i < numIters; i++) { diff --git a/modules/gpu/src/orb.cpp b/modules/gpu/src/orb.cpp index f59e545..46d96a7 100644 --- a/modules/gpu/src/orb.cpp +++ b/modules/gpu/src/orb.cpp @@ -62,7 +62,7 @@ void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat&) { throw_nogpu(); } #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace orb { @@ -431,7 +431,7 @@ cv::gpu::ORB_GPU::ORB_GPU(int nFeatures, float scaleFactor, int nLevels, int edg ++v_0; } CV_Assert(u_max.size() < 32); - cv::gpu::device::orb::loadUMax(&u_max[0], static_cast(u_max.size())); + cv::gpu::cuda::orb::loadUMax(&u_max[0], static_cast(u_max.size())); // Calc pattern const int npoints = 512; @@ -543,7 +543,7 @@ namespace //takes keypoints and culls them by the response void cull(GpuMat& keypoints, int& count, int n_points) { - using namespace cv::gpu::device::orb; + using namespace cv::gpu::cuda::orb; //this is only necessary if the keypoints size is greater than the number of desired points. if (count > n_points) @@ -561,7 +561,7 @@ namespace void cv::gpu::ORB_GPU::computeKeyPointsPyramid() { - using namespace cv::gpu::device::orb; + using namespace cv::gpu::cuda::orb; int half_patch_size = patchSize_ / 2; @@ -604,7 +604,7 @@ void cv::gpu::ORB_GPU::computeKeyPointsPyramid() void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors) { - using namespace cv::gpu::device::orb; + using namespace cv::gpu::cuda::orb; int nAllkeypoints = 0; @@ -644,7 +644,7 @@ void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors) void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat& keypoints) { - using namespace cv::gpu::device::orb; + using namespace cv::gpu::cuda::orb; int nAllkeypoints = 0; diff --git a/modules/gpu/src/pyramids.cpp b/modules/gpu/src/pyramids.cpp index 85fb990..2c8e80a 100644 --- a/modules/gpu/src/pyramids.cpp +++ b/modules/gpu/src/pyramids.cpp @@ -54,7 +54,7 @@ void cv::gpu::ImagePyramid::getLayer(GpuMat&, Size, Stream&) const { throw_nogpu ////////////////////////////////////////////////////////////////////////////// // pyrDown -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -64,7 +64,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::pyrDown(const GpuMat& src, GpuMat& dst, Stream& stream) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -92,7 +92,7 @@ void cv::gpu::pyrDown(const GpuMat& src, GpuMat& dst, Stream& stream) ////////////////////////////////////////////////////////////////////////////// // pyrUp -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -102,7 +102,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::pyrUp(const GpuMat& src, GpuMat& dst, Stream& stream) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -130,7 +130,7 @@ void cv::gpu::pyrUp(const GpuMat& src, GpuMat& dst, Stream& stream) ////////////////////////////////////////////////////////////////////////////// // ImagePyramid -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace pyramid { @@ -141,7 +141,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::ImagePyramid::build(const GpuMat& img, int numLayers, Stream& stream) { - using namespace cv::gpu::device::pyramid; + using namespace cv::gpu::cuda::pyramid; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -189,7 +189,7 @@ void cv::gpu::ImagePyramid::build(const GpuMat& img, int numLayers, Stream& stre void cv::gpu::ImagePyramid::getLayer(GpuMat& outImg, Size outRoi, Stream& stream) const { - using namespace cv::gpu::device::pyramid; + using namespace cv::gpu::cuda::pyramid; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); diff --git a/modules/gpu/src/remap.cpp b/modules/gpu/src/remap.cpp index 4b87286..80ec221 100644 --- a/modules/gpu/src/remap.cpp +++ b/modules/gpu/src/remap.cpp @@ -48,7 +48,7 @@ void cv::gpu::remap(const GpuMat&, GpuMat&, const GpuMat&, const GpuMat&, int, i #else // HAVE_CUDA -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -60,7 +60,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap, int interpolation, int borderMode, Scalar borderValue, Stream& stream) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, PtrStepSzf xmap, PtrStepSzf ymap, PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20); diff --git a/modules/gpu/src/resize.cpp b/modules/gpu/src/resize.cpp index 6850135..c913e45 100644 --- a/modules/gpu/src/resize.cpp +++ b/modules/gpu/src/resize.cpp @@ -59,7 +59,7 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub #else // HAVE_CUDA -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -137,7 +137,7 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub } else { - using namespace ::cv::gpu::device::imgproc; + using namespace ::cv::gpu::cuda::imgproc; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float fx, float fy, PtrStepSzb dst, int interpolation, cudaStream_t stream); diff --git a/modules/gpu/src/split_merge.cpp b/modules/gpu/src/split_merge.cpp index 61a928c..fea53ad 100644 --- a/modules/gpu/src/split_merge.cpp +++ b/modules/gpu/src/split_merge.cpp @@ -54,7 +54,7 @@ void cv::gpu::split(const GpuMat& /*src*/, std::vector& /*dst*/, Stream& #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace split_merge { @@ -67,7 +67,7 @@ namespace { void merge(const GpuMat* src, size_t n, GpuMat& dst, const cudaStream_t& stream) { - using namespace ::cv::gpu::device::split_merge; + using namespace ::cv::gpu::cuda::split_merge; CV_Assert(src); CV_Assert(n > 0); @@ -112,7 +112,7 @@ namespace void split(const GpuMat& src, GpuMat* dst, const cudaStream_t& stream) { - using namespace ::cv::gpu::device::split_merge; + using namespace ::cv::gpu::cuda::split_merge; CV_Assert(dst); diff --git a/modules/gpu/src/stereobm.cpp b/modules/gpu/src/stereobm.cpp index e633bf0..f3a95cb 100644 --- a/modules/gpu/src/stereobm.cpp +++ b/modules/gpu/src/stereobm.cpp @@ -55,7 +55,7 @@ void cv::gpu::StereoBM_GPU::operator() ( const GpuMat&, const GpuMat&, GpuMat&, #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace stereobm { @@ -101,7 +101,7 @@ namespace { void stereo_bm_gpu_operator( GpuMat& minSSD, GpuMat& leBuf, GpuMat& riBuf, int preset, int ndisp, int winSize, float avergeTexThreshold, const GpuMat& left, const GpuMat& right, GpuMat& disparity, cudaStream_t stream) { - using namespace ::cv::gpu::device::stereobm; + using namespace ::cv::gpu::cuda::stereobm; CV_Assert(left.rows == right.rows && left.cols == right.cols); CV_Assert(left.type() == CV_8UC1); diff --git a/modules/gpu/src/stereobp.cpp b/modules/gpu/src/stereobp.cpp index 2a472e2..bc4ca90 100644 --- a/modules/gpu/src/stereobp.cpp +++ b/modules/gpu/src/stereobp.cpp @@ -58,7 +58,7 @@ void cv::gpu::StereoBeliefPropagation::operator()(const GpuMat&, GpuMat&, Stream #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace stereobp { @@ -78,7 +78,7 @@ namespace cv { namespace gpu { namespace device } }}} -using namespace ::cv::gpu::device::stereobp; +using namespace ::cv::gpu::cuda::stereobp; namespace { diff --git a/modules/gpu/src/stereocsbp.cpp b/modules/gpu/src/stereocsbp.cpp index f3e55be..5c8f39a 100644 --- a/modules/gpu/src/stereocsbp.cpp +++ b/modules/gpu/src/stereocsbp.cpp @@ -57,7 +57,7 @@ void cv::gpu::StereoConstantSpaceBP::operator()(const GpuMat&, const GpuMat&, Gp #else /* !defined (HAVE_CUDA) */ #include "opencv2/core/utility.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace stereocsbp { @@ -89,7 +89,7 @@ namespace cv { namespace gpu { namespace device } }}} -using namespace ::cv::gpu::device::stereocsbp; +using namespace ::cv::gpu::cuda::stereocsbp; namespace { diff --git a/modules/gpu/src/video_reader.cpp b/modules/gpu/src/video_reader.cpp index 118eb95..36afd53 100644 --- a/modules/gpu/src/video_reader.cpp +++ b/modules/gpu/src/video_reader.cpp @@ -128,7 +128,7 @@ cv::gpu::VideoReader_GPU::Impl::~Impl() videoSource_->stop(); } -namespace cv { namespace gpu { namespace device { +namespace cv { namespace gpu { namespace cuda { namespace video_decoding { void loadHueCSC(float hueCSC[9]); @@ -189,7 +189,7 @@ namespace void cudaPostProcessFrame(const cv::gpu::GpuMat& decodedFrame, cv::gpu::GpuMat& interopFrame, int width, int height) { - using namespace cv::gpu::device::video_decoding; + using namespace cv::gpu::cuda::video_decoding; static bool updateCSC = true; static float hueColorSpaceMat[9]; diff --git a/modules/gpu/src/video_writer.cpp b/modules/gpu/src/video_writer.cpp index b134f51..8f4d78e 100644 --- a/modules/gpu/src/video_writer.cpp +++ b/modules/gpu/src/video_writer.cpp @@ -501,7 +501,7 @@ void cv::gpu::VideoWriter_GPU::Impl::createHWEncoder() CV_Assert( err == 0 ); } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace video_encoding { @@ -674,7 +674,7 @@ void cv::gpu::VideoWriter_GPU::Impl::write(const cv::gpu::GpuMat& frame, bool la CV_Assert( res == CUDA_SUCCESS ); if (inputFormat_ == SF_BGR) - cv::gpu::device::video_encoding::YV12_gpu(frame, frame.channels(), videoFrame_); + cv::gpu::cuda::video_encoding::YV12_gpu(frame, frame.channels(), videoFrame_); else { switch (surfaceFormat_) diff --git a/modules/gpu/src/warp.cpp b/modules/gpu/src/warp.cpp index 827d521..874c9d2 100644 --- a/modules/gpu/src/warp.cpp +++ b/modules/gpu/src/warp.cpp @@ -53,7 +53,7 @@ void cv::gpu::buildWarpPerspectiveMaps(const Mat&, bool, Size, GpuMat&, GpuMat&, #else // HAVE_CUDA -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace imgproc { @@ -73,7 +73,7 @@ namespace cv { namespace gpu { namespace device void cv::gpu::buildWarpAffineMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; CV_Assert(M.rows == 2 && M.cols == 3); @@ -97,7 +97,7 @@ void cv::gpu::buildWarpAffineMaps(const Mat& M, bool inverse, Size dsize, GpuMat void cv::gpu::buildWarpPerspectiveMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream) { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; CV_Assert(M.rows == 3 && M.cols == 3); @@ -272,7 +272,7 @@ void cv::gpu::warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsiz } else { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[2 * 3], PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20); @@ -410,7 +410,7 @@ void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size } else { - using namespace cv::gpu::device::imgproc; + using namespace cv::gpu::cuda::imgproc; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[2 * 3], PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20); diff --git a/modules/nonfree/src/cuda/surf.cu b/modules/nonfree/src/cuda/surf.cu index 39de3c3..c78e11e 100644 --- a/modules/nonfree/src/cuda/surf.cu +++ b/modules/nonfree/src/cuda/surf.cu @@ -44,15 +44,15 @@ #ifdef HAVE_OPENCV_GPU -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/limits.hpp" -#include "opencv2/core/device/saturate_cast.hpp" -#include "opencv2/core/device/reduce.hpp" -#include "opencv2/core/device/utility.hpp" -#include "opencv2/core/device/functional.hpp" -#include "opencv2/core/device/filters.hpp" - -namespace cv { namespace gpu { namespace device +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/limits.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/functional.hpp" +#include "opencv2/core/cuda/filters.hpp" + +namespace cv { namespace gpu { namespace cuda { namespace surf { @@ -79,7 +79,7 @@ namespace cv { namespace gpu { namespace device } }}} -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace surf { @@ -626,7 +626,7 @@ namespace cv { namespace gpu { namespace device } plus op; - device::reduce<32>(smem_tuple(s_sumx + threadIdx.y * 32, s_sumy + threadIdx.y * 32), + cuda::reduce<32>(smem_tuple(s_sumx + threadIdx.y * 32, s_sumy + threadIdx.y * 32), thrust::tie(sumx, sumy), threadIdx.x, thrust::make_tuple(op, op)); const float temp_mod = sumx * sumx + sumy * sumy; @@ -955,7 +955,7 @@ namespace cv { namespace gpu { namespace device } } } // namespace surf -}}} // namespace cv { namespace gpu { namespace device +}}} // namespace cv { namespace gpu { namespace cuda #endif /* CUDA_DISABLER */ diff --git a/modules/nonfree/src/cuda/vibe.cu b/modules/nonfree/src/cuda/vibe.cu index 52cbed7..3e1299b 100644 --- a/modules/nonfree/src/cuda/vibe.cu +++ b/modules/nonfree/src/cuda/vibe.cu @@ -44,9 +44,9 @@ #ifdef HAVE_OPENCV_GPU -#include "opencv2/core/device/common.hpp" +#include "opencv2/core/cuda/common.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace vibe { @@ -58,7 +58,7 @@ namespace cv { namespace gpu { namespace device } }}} -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace vibe { diff --git a/modules/nonfree/src/precomp.hpp b/modules/nonfree/src/precomp.hpp index 51d2aba..0d5c288 100644 --- a/modules/nonfree/src/precomp.hpp +++ b/modules/nonfree/src/precomp.hpp @@ -56,7 +56,7 @@ #if defined(HAVE_CUDA) #include "opencv2/core/stream_accessor.hpp" - #include "opencv2/core/device/common.hpp" + #include "opencv2/core/cuda/common.hpp" static inline void throw_nogpu() { CV_Error(CV_StsNotImplemented, "The called functionality is disabled for current build or platform"); } #else diff --git a/modules/nonfree/src/surf_gpu.cpp b/modules/nonfree/src/surf_gpu.cpp index a5f7389..0fe3dff 100644 --- a/modules/nonfree/src/surf_gpu.cpp +++ b/modules/nonfree/src/surf_gpu.cpp @@ -64,7 +64,7 @@ void cv::gpu::SURF_GPU::releaseMemory() { throw_nogpu(); } #else // !defined (HAVE_CUDA) -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace surf { @@ -91,7 +91,7 @@ namespace cv { namespace gpu { namespace device } }}} -using namespace ::cv::gpu::device::surf; +using namespace ::cv::gpu::cuda::surf; namespace { diff --git a/modules/nonfree/src/vibe_gpu.cpp b/modules/nonfree/src/vibe_gpu.cpp index e348627..7d74506 100644 --- a/modules/nonfree/src/vibe_gpu.cpp +++ b/modules/nonfree/src/vibe_gpu.cpp @@ -53,7 +53,7 @@ void cv::gpu::VIBE_GPU::release() {} #else -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { namespace vibe { @@ -84,7 +84,7 @@ cv::gpu::VIBE_GPU::VIBE_GPU(unsigned long rngSeed) : void cv::gpu::VIBE_GPU::initialize(const GpuMat& firstFrame, Stream& s) { - using namespace cv::gpu::device::vibe; + using namespace cv::gpu::cuda::vibe; CV_Assert(firstFrame.type() == CV_8UC1 || firstFrame.type() == CV_8UC3 || firstFrame.type() == CV_8UC4); @@ -112,7 +112,7 @@ void cv::gpu::VIBE_GPU::initialize(const GpuMat& firstFrame, Stream& s) void cv::gpu::VIBE_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, Stream& s) { - using namespace cv::gpu::device::vibe; + using namespace cv::gpu::cuda::vibe; CV_Assert(frame.depth() == CV_8U); diff --git a/modules/softcascade/src/cuda/channels.cu b/modules/softcascade/src/cuda/channels.cu index 782e243..558cbd4 100644 --- a/modules/softcascade/src/cuda/channels.cu +++ b/modules/softcascade/src/cuda/channels.cu @@ -61,7 +61,7 @@ __host__ __device__ __forceinline__ int divUp(int total, int grain) return (total + grain - 1) / grain; } -namespace cv { namespace softcascade { namespace device +namespace cv { namespace softcascade { namespace cuda { typedef unsigned int uint; typedef unsigned short ushort; diff --git a/modules/softcascade/src/cuda/icf-sc.cu b/modules/softcascade/src/cuda/icf-sc.cu index 6bea949..1f6a0bc 100644 --- a/modules/softcascade/src/cuda/icf-sc.cu +++ b/modules/softcascade/src/cuda/icf-sc.cu @@ -62,7 +62,7 @@ static inline void ___cudaSafeCall(cudaError_t err, const char *file, const int #define CV_PI 3.1415926535897932384626433832795 #endif -namespace cv { namespace softcascade { namespace device { +namespace cv { namespace softcascade { namespace cuda { typedef unsigned char uchar; diff --git a/modules/softcascade/src/cuda_invoker.hpp b/modules/softcascade/src/cuda_invoker.hpp index 135f405..c173b6a 100644 --- a/modules/softcascade/src/cuda_invoker.hpp +++ b/modules/softcascade/src/cuda_invoker.hpp @@ -54,7 +54,7 @@ #endif -namespace cv { namespace softcascade { namespace device { +namespace cv { namespace softcascade { namespace cuda { typedef unsigned char uchar; typedef unsigned int uint; diff --git a/modules/softcascade/src/detector_cuda.cpp b/modules/softcascade/src/detector_cuda.cpp index f0d71c6..6f9c907 100644 --- a/modules/softcascade/src/detector_cuda.cpp +++ b/modules/softcascade/src/detector_cuda.cpp @@ -77,7 +77,7 @@ namespace } } -cv::softcascade::device::Level::Level(int idx, const Octave& oct, const float scale, const int w, const int h) +cv::softcascade::cuda::Level::Level(int idx, const Octave& oct, const float scale, const int w, const int h) : octave(idx), step(oct.stages), relScale(scale / oct.scale) { workRect.x = (unsigned char)cvRound(w / (float)oct.shrinkage); @@ -96,7 +96,7 @@ cv::softcascade::device::Level::Level(int idx, const Octave& oct, const float sc } } -namespace cv { namespace softcascade { namespace device { +namespace cv { namespace softcascade { namespace cuda { void fillBins(cv::gpu::PtrStepSzb hogluv, const cv::gpu::PtrStepSzf& nangle, const int fw, const int fh, const int bins, cudaStream_t stream); @@ -153,9 +153,9 @@ struct cv::softcascade::SCascade::Fields FileNode fn = root[SC_OCTAVES]; if (fn.empty()) return 0; - std::vector voctaves; + std::vector voctaves; std::vector vstages; - std::vector vnodes; + std::vector vnodes; std::vector vleaves; FileNodeIterator it = fn.begin(), it_end = fn.end(); @@ -171,7 +171,7 @@ struct cv::softcascade::SCascade::Fields size.x = (unsigned short)cvRound(origWidth * scale); size.y = (unsigned short)cvRound(origHeight * scale); - device::Octave octave(octIndex, nweaks, shrinkage, size, scale); + cuda::Octave octave(octIndex, nweaks, shrinkage, size, scale); CV_Assert(octave.stages > 0); voctaves.push_back(octave); @@ -240,7 +240,7 @@ struct cv::softcascade::SCascade::Fields rect.w = saturate_cast(r.height); unsigned int channel = saturate_cast(feature_channels[featureIdx]); - vnodes.push_back(device::Node(rect, channel, th)); + vnodes.push_back(cuda::Node(rect, channel, th)); } intfns = octfn[SC_LEAF]; @@ -252,13 +252,13 @@ struct cv::softcascade::SCascade::Fields } } - cv::Mat hoctaves(1, (int) (voctaves.size() * sizeof(device::Octave)), CV_8UC1, (uchar*)&(voctaves[0])); + cv::Mat hoctaves(1, (int) (voctaves.size() * sizeof(cuda::Octave)), CV_8UC1, (uchar*)&(voctaves[0])); CV_Assert(!hoctaves.empty()); cv::Mat hstages(cv::Mat(vstages).reshape(1,1)); CV_Assert(!hstages.empty()); - cv::Mat hnodes(1, (int) (vnodes.size() * sizeof(device::Node)), CV_8UC1, (uchar*)&(vnodes[0]) ); + cv::Mat hnodes(1, (int) (vnodes.size() * sizeof(cuda::Node)), CV_8UC1, (uchar*)&(vnodes[0]) ); CV_Assert(!hnodes.empty()); cv::Mat hleaves(cv::Mat(vleaves).reshape(1,1)); @@ -285,7 +285,7 @@ struct cv::softcascade::SCascade::Fields int createLevels(const int fh, const int fw) { - std::vector vlevels; + std::vector vlevels; float logFactor = (::log(maxScale) - ::log(minScale)) / (totals -1); float scale = minScale; @@ -298,7 +298,7 @@ struct cv::softcascade::SCascade::Fields float logScale = ::log(scale); int fit = fitOctave(voctaves, logScale); - device::Level level(fit, voctaves[fit], scale, width, height); + cuda::Level level(fit, voctaves[fit], scale, width, height); if (!width || !height) break; @@ -312,7 +312,7 @@ struct cv::softcascade::SCascade::Fields scale = ::std::min(maxScale, ::expf(::log(scale) + logFactor)); } - cv::Mat hlevels = cv::Mat(1, (int) (vlevels.size() * sizeof(device::Level)), CV_8UC1, (uchar*)&(vlevels[0]) ); + cv::Mat hlevels = cv::Mat(1, (int) (vlevels.size() * sizeof(cuda::Level)), CV_8UC1, (uchar*)&(vlevels[0]) ); CV_Assert(!hlevels.empty()); levels.upload(hlevels); downscales = dcs; @@ -355,8 +355,8 @@ struct cv::softcascade::SCascade::Fields cudaSafeCall( cudaGetLastError()); - device::CascadeInvoker invoker - = device::CascadeInvoker(levels, stages, nodes, leaves); + cuda::CascadeInvoker invoker + = cuda::CascadeInvoker(levels, stages, nodes, leaves); cudaStream_t stream = cv::gpu::StreamAccessor::getStream(s); invoker(mask, hogluv, objects, downscales, stream); @@ -379,19 +379,19 @@ struct cv::softcascade::SCascade::Fields } cudaStream_t stream = cv::gpu::StreamAccessor::getStream(s); - device::suppress(objects, overlaps, ndetections, suppressed, stream); + cuda::suppress(objects, overlaps, ndetections, suppressed, stream); } private: - typedef std::vector::const_iterator octIt_t; - static int fitOctave(const std::vector& octs, const float& logFactor) + typedef std::vector::const_iterator octIt_t; + static int fitOctave(const std::vector& octs, const float& logFactor) { float minAbsLog = FLT_MAX; int res = 0; for (int oct = 0; oct < (int)octs.size(); ++oct) { - const device::Octave& octave =octs[oct]; + const cuda::Octave& octave =octs[oct]; float logOctave = ::log(octave.scale); float logAbsScale = ::fabs(logFactor - logOctave); @@ -452,7 +452,7 @@ public: // cv::gpu::GpuMat collected; - std::vector voctaves; + std::vector voctaves; // DeviceInfo info; @@ -498,7 +498,7 @@ void integral(const cv::gpu::GpuMat& src, cv::gpu::GpuMat& sum, cv::gpu::GpuMat& { ensureSizeIsEnough(((src.rows + 7) / 8) * 8, ((src.cols + 63) / 64) * 64, CV_32SC1, buffer); - cv::softcascade::device::shfl_integral(src, buffer, stream); + cv::softcascade::cuda::shfl_integral(src, buffer, stream); sum.create(src.rows + 1, src.cols + 1, CV_32SC1); if (s) @@ -539,7 +539,7 @@ void cv::softcascade::SCascade::detect(InputArray _image, InputArray _rois, Outp flds.mask.create( rois.cols / shr, rois.rows / shr, rois.type()); - device::shrink(rois, flds.mask); + cuda::shrink(rois, flds.mask); //cv::gpu::transpose(flds.genRoiTmp, flds.mask, s); if (type == CV_8UC3) @@ -607,12 +607,12 @@ struct SeparablePreprocessor : public cv::softcascade::ChannelsProcessor setZero(channels, s); gray.create(bgr.size(), CV_8UC1); - cv::softcascade::device::transform(bgr, gray); //cv::gpu::cvtColor(bgr, gray, CV_BGR2GRAY); - cv::softcascade::device::gray2hog(gray, channels(cv::Rect(0, 0, bgr.cols, bgr.rows * (bins + 1))), bins); + cv::softcascade::cuda::transform(bgr, gray); //cv::gpu::cvtColor(bgr, gray, CV_BGR2GRAY); + cv::softcascade::cuda::gray2hog(gray, channels(cv::Rect(0, 0, bgr.cols, bgr.rows * (bins + 1))), bins); cv::gpu::GpuMat luv(channels, cv::Rect(0, bgr.rows * (bins + 1), bgr.cols, bgr.rows * 3)); - cv::softcascade::device::bgr2Luv(bgr, luv); - cv::softcascade::device::shrink(channels, shrunk); + cv::softcascade::cuda::bgr2Luv(bgr, luv); + cv::softcascade::cuda::shrink(channels, shrunk); } private: diff --git a/modules/superres/src/btv_l1_gpu.cpp b/modules/superres/src/btv_l1_gpu.cpp index fffcb23..8fbbf17 100644 --- a/modules/superres/src/btv_l1_gpu.cpp +++ b/modules/superres/src/btv_l1_gpu.cpp @@ -60,7 +60,7 @@ Ptr cv::superres::createSuperResolution_BTVL1_GPU() #else // HAVE_CUDA -namespace btv_l1_device +namespace btv_l1_cuda { void buildMotionMaps(PtrStepSzf forwardMotionX, PtrStepSzf forwardMotionY, PtrStepSzf backwardMotionX, PtrStepSzf bacwardMotionY, @@ -138,7 +138,7 @@ namespace backwardMap.first.create(forwardMotion.first.size(), CV_32FC1); backwardMap.second.create(forwardMotion.first.size(), CV_32FC1); - btv_l1_device::buildMotionMaps(forwardMotion.first, forwardMotion.second, + btv_l1_cuda::buildMotionMaps(forwardMotion.first, forwardMotion.second, backwardMotion.first, backwardMotion.second, forwardMap.first, forwardMap.second, backwardMap.first, backwardMap.second); @@ -149,7 +149,7 @@ namespace typedef void (*func_t)(const PtrStepSzb src, PtrStepSzb dst, int scale, cudaStream_t stream); static const func_t funcs[] = { - 0, btv_l1_device::upscale<1>, 0, btv_l1_device::upscale<3>, btv_l1_device::upscale<4> + 0, btv_l1_cuda::upscale<1>, 0, btv_l1_cuda::upscale<3>, btv_l1_cuda::upscale<4> }; CV_Assert( src.channels() == 1 || src.channels() == 3 || src.channels() == 4 ); @@ -166,7 +166,7 @@ namespace { dst.create(src1.size(), src1.type()); - btv_l1_device::diffSign(src1.reshape(1), src2.reshape(1), dst.reshape(1), StreamAccessor::getStream(stream)); + btv_l1_cuda::diffSign(src1.reshape(1), src2.reshape(1), dst.reshape(1), StreamAccessor::getStream(stream)); } void calcBtvWeights(int btvKernelSize, double alpha, std::vector& btvWeights) @@ -184,7 +184,7 @@ namespace btvWeights[ind] = pow(alpha_f, std::abs(m) + std::abs(l)); } - btv_l1_device::loadBtvWeights(&btvWeights[0], size); + btv_l1_cuda::loadBtvWeights(&btvWeights[0], size); } void calcBtvRegularization(const GpuMat& src, GpuMat& dst, int btvKernelSize) @@ -193,10 +193,10 @@ namespace static const func_t funcs[] = { 0, - btv_l1_device::calcBtvRegularization<1>, + btv_l1_cuda::calcBtvRegularization<1>, 0, - btv_l1_device::calcBtvRegularization<3>, - btv_l1_device::calcBtvRegularization<4> + btv_l1_cuda::calcBtvRegularization<3>, + btv_l1_cuda::calcBtvRegularization<4> }; dst.create(src.size(), src.type()); diff --git a/modules/superres/src/cuda/btv_l1_gpu.cu b/modules/superres/src/cuda/btv_l1_gpu.cu index 2bb93d1..09f67e6 100644 --- a/modules/superres/src/cuda/btv_l1_gpu.cu +++ b/modules/superres/src/cuda/btv_l1_gpu.cu @@ -40,15 +40,15 @@ // //M*/ -#include "opencv2/core/device/common.hpp" -#include "opencv2/core/device/transform.hpp" -#include "opencv2/core/device/vec_traits.hpp" -#include "opencv2/core/device/vec_math.hpp" +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/transform.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" using namespace cv::gpu; -using namespace cv::gpu::device; +using namespace cv::gpu::cuda; -namespace btv_l1_device +namespace btv_l1_cuda { void buildMotionMaps(PtrStepSzf forwardMotionX, PtrStepSzf forwardMotionY, PtrStepSzf backwardMotionX, PtrStepSzf bacwardMotionY, @@ -64,7 +64,7 @@ namespace btv_l1_device template void calcBtvRegularization(PtrStepSzb src, PtrStepSzb dst, int ksize); } -namespace btv_l1_device +namespace btv_l1_cuda { __global__ void buildMotionMapsKernel(const PtrStepSzf forwardMotionX, const PtrStepf forwardMotionY, PtrStepf backwardMotionX, PtrStepf backwardMotionY, @@ -169,16 +169,16 @@ namespace btv_l1_device }; } -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace cuda { - template <> struct TransformFunctorTraits : DefaultTransformFunctorTraits + template <> struct TransformFunctorTraits : DefaultTransformFunctorTraits { enum { smart_block_dim_y = 8 }; enum { smart_shift = 4 }; }; }}} -namespace btv_l1_device +namespace btv_l1_cuda { void diffSign(PtrStepSzf src1, PtrStepSzf src2, PtrStepSzf dst, cudaStream_t stream) { -- 2.7.4