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")
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})
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})
#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
{
#include "vec_traits.hpp"
#include "vec_math.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
//////////////////////////////////////////////////////////////
// BrdConstant
const int width;
const D val;
};
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif // __OPENCV_GPU_BORDER_INTERPOLATE_HPP__
#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 <typename T> class ColorSpace1_to_ColorSpace2_traits
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__
return (total + grain - 1) / grain;
}
- namespace device
+ namespace cuda
{
using cv::gpu::divUp;
#include "common.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
#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__
#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))
#define OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2RGB<T, scn, dcn, bidx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2RGB<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(name, scn, bidx, green_bits) \
struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2RGB5x5<scn, bidx, green_bits> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2RGB5x5<scn, bidx, green_bits> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(name, dcn, bidx, green_bits) \
struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB5x52RGB<dcn, bidx, green_bits> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB5x52RGB<dcn, bidx, green_bits> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_GRAY2RGB_TRAITS(name, dcn) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::Gray2RGB<T, dcn> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::Gray2RGB<T, dcn> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_GRAY2RGB5x5_TRAITS(name, green_bits) \
struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::Gray2RGB5x5<green_bits> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::Gray2RGB5x5<green_bits> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB5x52GRAY_TRAITS(name, green_bits) \
struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB5x52Gray<green_bits> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB5x52Gray<green_bits> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS(name, scn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2Gray<T, scn, bidx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2Gray<T, scn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2YUV<T, scn, dcn, bidx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2YUV<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::YUV2RGB<T, scn, dcn, bidx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::YUV2RGB<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2YCrCb<T, scn, dcn, bidx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2YCrCb<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::YCrCb2RGB<T, scn, dcn, bidx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::YCrCb2RGB<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2XYZ<T, scn, dcn, bidx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2XYZ<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::XYZ2RGB<T, scn, dcn, bidx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::XYZ2RGB<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2HSV<T, scn, dcn, bidx, 180> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2HSV<T, scn, dcn, bidx, 180> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <typename T> struct name ## _full_traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2HSV<T, scn, dcn, bidx, 256> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2HSV<T, scn, dcn, bidx, 256> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <> struct name ## _traits<float> \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2HSV<float, scn, dcn, bidx, 360> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2HSV<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <> struct name ## _full_traits<float> \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2HSV<float, scn, dcn, bidx, 360> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2HSV<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::HSV2RGB<T, scn, dcn, bidx, 180> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::HSV2RGB<T, scn, dcn, bidx, 180> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <typename T> struct name ## _full_traits \
{ \
- typedef ::cv::gpu::device::color_detail::HSV2RGB<T, scn, dcn, bidx, 255> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::HSV2RGB<T, scn, dcn, bidx, 255> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <> struct name ## _traits<float> \
{ \
- typedef ::cv::gpu::device::color_detail::HSV2RGB<float, scn, dcn, bidx, 360> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::HSV2RGB<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <> struct name ## _full_traits<float> \
{ \
- typedef ::cv::gpu::device::color_detail::HSV2RGB<float, scn, dcn, bidx, 360> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::HSV2RGB<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2HLS<T, scn, dcn, bidx, 180> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2HLS<T, scn, dcn, bidx, 180> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <typename T> struct name ## _full_traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2HLS<T, scn, dcn, bidx, 256> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2HLS<T, scn, dcn, bidx, 256> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <> struct name ## _traits<float> \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2HLS<float, scn, dcn, bidx, 360> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2HLS<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <> struct name ## _full_traits<float> \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2HLS<float, scn, dcn, bidx, 360> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2HLS<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::HLS2RGB<T, scn, dcn, bidx, 180> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::HLS2RGB<T, scn, dcn, bidx, 180> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <typename T> struct name ## _full_traits \
{ \
- typedef ::cv::gpu::device::color_detail::HLS2RGB<T, scn, dcn, bidx, 255> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::HLS2RGB<T, scn, dcn, bidx, 255> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <> struct name ## _traits<float> \
{ \
- typedef ::cv::gpu::device::color_detail::HLS2RGB<float, scn, dcn, bidx, 360> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::HLS2RGB<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
}; \
template <> struct name ## _full_traits<float> \
{ \
- typedef ::cv::gpu::device::color_detail::HLS2RGB<float, scn, dcn, bidx, 360> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::HLS2RGB<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(name, scn, dcn, srgb, blueIdx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2Lab<T, scn, dcn, srgb, blueIdx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2Lab<T, scn, dcn, srgb, blueIdx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::Lab2RGB<T, scn, dcn, srgb, blueIdx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::Lab2RGB<T, scn, dcn, srgb, blueIdx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(name, scn, dcn, srgb, blueIdx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::RGB2Luv<T, scn, dcn, srgb, blueIdx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::RGB2Luv<T, scn, dcn, srgb, blueIdx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#define OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \
template <typename T> struct name ## _traits \
{ \
- typedef ::cv::gpu::device::color_detail::Luv2RGB<T, scn, dcn, srgb, blueIdx> functor_type; \
+ typedef ::cv::gpu::cuda::color_detail::Luv2RGB<T, scn, dcn, srgb, blueIdx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
#undef CV_DESCALE
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif // __OPENCV_GPU_COLOR_DETAIL_HPP__
#include "../warp.hpp"
#include "../warp_shuffle.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace reduce_detail
{
#include "../warp.hpp"
#include "../warp_shuffle.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace reduce_key_val_detail
{
#include "../vec_traits.hpp"
#include "../functional.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace transform_detail
{
}
};
} // namespace transform_detail
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif // __OPENCV_GPU_TRANSFORM_DETAIL_HPP__
#include "../common.hpp"
#include "../vec_traits.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace type_traits_detail
{
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__
#include "../datamov_utils.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace vec_distance_detail
{
}
};
} // namespace vec_distance_detail
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif // __OPENCV_GPU_VEC_DISTANCE_DETAIL_HPP__
#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<class T> struct DynamicSharedMem
{
#include "warp_reduce.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
struct Emulation
{
}
};
};
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* OPENCV_GPU_EMULATION_HPP_ */
#include "vec_math.hpp"
#include "type_traits.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename Ptr2D> struct PointFilter
{
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__
#include <cstdio>
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template<class Func>
void printFuncAttrib(Func& func)
printf("\n");
fflush(stdout);
}
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_ */
#include "type_traits.hpp"
#include "device_functions.h"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
// Function Objects
template<typename Argument, typename Result> struct unary_function : public std::unary_function<Argument, Result> {};
#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__
#include <limits>
#include "common.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template<class T> struct numeric_limits
{
__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__
#include "detail/reduce.hpp"
#include "detail/reduce_key_val.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <int N, typename T, class Op>
__device__ __forceinline__ void reduce(volatile T* smem, T& val, unsigned int tid, const Op& op)
#include "common.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uchar v) { return _Tp(v); }
template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(schar v) { return _Tp(v); }
#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 };
__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;
}
vmin4(a,b) per-byte unsigned minimum: min(a, b)
*/
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
// 2
namespace cv { namespace gpu
{
- namespace device
+ namespace cuda
{
template<bool expr> struct Static {};
#include "utility.hpp"
#include "detail/transform_detail.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T, typename D, typename UnOp, typename Mask>
static inline void transform(PtrStepSz<T> src, PtrStepSz<D> dst, UnOp op, const Mask& mask, cudaStream_t stream)
#include "detail/type_traits_detail.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct IsSimpleParameter
{
#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)
return false;
}
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif // __OPENCV_GPU_UTILITY_HPP__
#include "functional.hpp"
#include "detail/vec_distance_detail.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct L1Dist
{
U vec1Vals[MAX_LEN / THREAD_DIM];
};
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif // __OPENCV_GPU_VEC_DISTANCE_HPP__
#include "vec_traits.hpp"
#include "functional.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace vec_math_detail
{
#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__
#include "common.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template<typename T, int N> struct TypeVec;
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__
#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
{
*t = value;
}
};
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* __OPENCV_GPU_DEVICE_WARP_HPP__ */
#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 <class T>
__device__ __forceinline__ T warp_reduce(volatile T *ptr , const unsigned int tid = threadIdx.x)
return ptr[tid - lane];
}
-}}} // namespace cv { namespace gpu { namespace device {
+}}} // namespace cv { namespace gpu { namespace cuda {
#endif /* OPENCV_GPU_WARP_REDUCE_HPP__ */
#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 <typename T>
__device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize)
typedef DevMem2Db DevMem2D;
typedef DevMem2D_<float> DevMem2Df;
typedef DevMem2D_<int> 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;
- }
}
}
//
//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*);
void convert_gpu(PtrStepSzb, int, PtrStepSzb, int, double, double, cudaStream_t);
}}}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct shift_and_sizeof;
template <> struct shift_and_sizeof<signed char> { enum { shift = 0 }; };
template <typename T> void copyToWithMask(PtrStepSzb src, PtrStepSzb dst, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream)
{
if (colorMask)
- cv::gpu::device::transform((PtrStepSz<T>)src, (PtrStepSz<T>)dst, identity<T>(), SingleMask(mask), stream);
+ cv::gpu::cuda::transform((PtrStepSz<T>)src, (PtrStepSz<T>)dst, identity<T>(), SingleMask(mask), stream);
else
- cv::gpu::device::transform((PtrStepSz<T>)src, (PtrStepSz<T>)dst, identity<T>(), SingleMaskChannels(mask, cn), stream);
+ cv::gpu::cuda::transform((PtrStepSz<T>)src, (PtrStepSz<T>)dst, identity<T>(), SingleMaskChannels(mask, cn), stream);
}
void copyToWithMask_gpu(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream)
cudaSafeCall( cudaSetDoubleForDevice(&alpha) );
cudaSafeCall( cudaSetDoubleForDevice(&beta) );
Convertor<T, D, S> op(static_cast<S>(alpha), static_cast<S>(beta));
- cv::gpu::device::transform((PtrStepSz<T>)src, (PtrStepSz<D>)dst, op, WithOutMask(), stream);
+ cv::gpu::cuda::transform((PtrStepSz<T>)src, (PtrStepSz<D>)dst, op, WithOutMask(), stream);
}
#if defined __clang__
#if defined __clang__
# pragma clang diagnostic pop
#endif
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#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);
template <typename T> void kernelSetCaller(GpuMat& src, Scalar s, cudaStream_t stream)
{
Scalar_<T> 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 <typename T> void kernelSetCaller(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream)
{
Scalar_<T> 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);
}
}
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)
////////////////////////////////////////////////////////////////////////
// Polar <-> Cart
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace mathfunc
{
{
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);
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);
#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,
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);
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);
#else
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace mog
{
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);
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_);
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);
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;
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_);
#else /* !defined (HAVE_CUDA) */
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace disp_bilateral_filter
{
}
}}}
-using namespace ::cv::gpu::device::disp_bilateral_filter;
+using namespace ::cv::gpu::cuda::disp_bilateral_filter;
namespace
{
#else
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace blend
{
}
}}}
-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)
#else /* !defined (HAVE_CUDA) */
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace bf_match
{
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,
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<PtrStepb>& masks,
const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance,
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,
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<PtrStepb>& masks,
const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance,
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<unsigned int>& nMatches,
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<unsigned int>& nMatches,
#else
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace transform_points
{
}
}}}
-using namespace ::cv::gpu::device;
+using namespace ::cv::gpu::cuda;
namespace
{
cv::Size sWindow;
};
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace lbp
{
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<unsigned int>(), integral);
}
return 0;
cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) );
- device::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified.ptr<unsigned int>());
+ cuda::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified.ptr<unsigned int>());
cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) );
cudaSafeCall( cudaDeviceSynchronize() );
#include "cvt_color_internal.h"
namespace cv { namespace gpu {
- namespace device
+ namespace cuda
{
template <int cn>
void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
}
}}
-using namespace ::cv::gpu::device;
+using namespace ::cv::gpu::cuda;
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);
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);
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);
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);
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);
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);
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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);
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);
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)
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)
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)
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);
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);
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);
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);
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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] =
{
{
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;
}
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;
}
#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);
#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
{
//template void match2Hamming_gpu<short >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, cudaStream_t stream);
template void match2Hamming_gpu<int >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& 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 */
#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
{
//template void matchHamming_gpu<short >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, cudaStream_t stream);
template void matchHamming_gpu<int >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& 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 */
#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
{
//template void matchHamming_gpu<short >(const PtrStepSzb& query, const PtrStepSzb* trains, int n, float maxDistance, const PtrStepSzb* masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, const PtrStepSz<unsigned int>& nMatches, cudaStream_t stream);
template void matchHamming_gpu<int >(const PtrStepSzb& query, const PtrStepSzb* trains, int n, float maxDistance, const PtrStepSzb* masks, const PtrStepSzi& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzf& distance, const PtrStepSz<unsigned int>& nMatches, cudaStream_t stream);
} // namespace bf_radius_match
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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;
#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
{
#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;
//////////////////////////////////////////////////////////////////////////////////
/// Bilateral filtering
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
#define OCV_INSTANTIATE_BILATERAL_FILTER(T) \
- template void cv::gpu::device::imgproc::bilateral_filter_gpu<T>(const PtrStepSzb&, PtrStepSzb, int, float, float, int, cudaStream_t);
+ template void cv::gpu::cuda::imgproc::bilateral_filter_gpu<T>(const PtrStepSzb&, PtrStepSzb, int, float, float, int, cudaStream_t);
OCV_INSTANTIATE_BILATERAL_FILTER(uchar)
//OCV_INSTANTIATE_BILATERAL_FILTER(uchar2)
#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
{
cudaSafeCall(cudaDeviceSynchronize());
}
} // namespace blend
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
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
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
cudaSafeCall( cudaDeviceSynchronize() );
}
} // namespace solvepnp_ransac
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#include <utility>
#include <algorithm>//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
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits<canny::L1> : DefaultTransformFunctorTraits<canny::L1>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits<canny::GetEdges> : DefaultTransformFunctorTraits<canny::GetEdges>
{
#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/emulation.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/emulation.hpp>
#include <iostream>
#include <stdio.h>
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace ccl
{
#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
{
#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<uchar>::functor_type)
{
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_t>)src, (PtrStepSz<dst_t>)dst, functor, WithOutMask(), stream); \
+ cv::gpu::cuda::transform((PtrStepSz<src_t>)src, (PtrStepSz<dst_t>)dst, functor, WithOutMask(), stream); \
}
#define OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(name) \
#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 */
//
//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
{
#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
{
template void copyMakeBorder_gpu<float, 3>(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderMode, const float* borderValue, cudaStream_t stream);
template void copyMakeBorder_gpu<float, 4>(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 */
#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 <typename T> struct Bayer2BGR;
#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
{
template void disp_bilateral_filter<uchar>(PtrStepSz<uchar> disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream);
template void disp_bilateral_filter<short>(PtrStepSz<short> 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 */
#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
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits< arithm::VAdd4 > : arithm::ArithmFuncTraits<sizeof(uint), sizeof(uint)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T, typename S, typename D> struct TransformFunctorTraits< arithm::AddScalar<T, S, D> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(D)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits< arithm::VSub4 > : arithm::ArithmFuncTraits<sizeof(uint), sizeof(uint)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits<arithm::Mul_8uc4_32f> : arithm::ArithmFuncTraits<sizeof(uint), sizeof(uint)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T, typename S, typename D> struct TransformFunctorTraits< arithm::MulScalar<T, S, D> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(D)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits<arithm::Div_8uc4_32f> : arithm::ArithmFuncTraits<sizeof(uint), sizeof(uint)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T, typename S, typename D> struct TransformFunctorTraits< arithm::DivInv<T, S, D> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(D)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits< arithm::VAbsDiff4 > : arithm::ArithmFuncTraits<sizeof(uint), sizeof(uint)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T, typename S> struct TransformFunctorTraits< arithm::AbsDiffScalar<T, S> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
//////////////////////////////////////////////////////////////////////////
// absMat
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< abs_func<T> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< arithm::Sqr<T> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
//////////////////////////////////////////////////////////////////////////
// sqrtMat
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< sqrt_func<T> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
//////////////////////////////////////////////////////////////////////////
// logMat
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< log_func<T> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< arithm::Exp<T> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits< arithm::VCmpEq4 > : arithm::ArithmFuncTraits<sizeof(uint), sizeof(uint)>
{
#undef TYPE_VEC
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <class Op, typename T> struct TransformFunctorTraits< arithm::CmpScalar<Op, T, 1> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(uchar)>
{
//////////////////////////////////////////////////////////////////////////////////////
// bitMat
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< bit_not<T> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
//////////////////////////////////////////////////////////////////////////////////////
// bitScalar
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< binder2nd< bit_and<T> > > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
{
template <typename T> void bitScalarAnd(PtrStepSzb src1, uint src2, PtrStepSzb dst, cudaStream_t stream)
{
- transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::device::bind2nd(bit_and<T>(), src2), WithOutMask(), stream);
+ transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::cuda::bind2nd(bit_and<T>(), src2), WithOutMask(), stream);
}
template <typename T> void bitScalarOr(PtrStepSzb src1, uint src2, PtrStepSzb dst, cudaStream_t stream)
{
- transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::device::bind2nd(bit_or<T>(), src2), WithOutMask(), stream);
+ transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::cuda::bind2nd(bit_or<T>(), src2), WithOutMask(), stream);
}
template <typename T> void bitScalarXor(PtrStepSzb src1, uint src2, PtrStepSzb dst, cudaStream_t stream)
{
- transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::device::bind2nd(bit_xor<T>(), src2), WithOutMask(), stream);
+ transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::cuda::bind2nd(bit_xor<T>(), src2), WithOutMask(), stream);
}
template void bitScalarAnd<uchar>(PtrStepSzb src1, uint src2, PtrStepSzb dst, cudaStream_t stream);
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits< arithm::VMin4 > : arithm::ArithmFuncTraits<sizeof(uint), sizeof(uint)>
{
template <typename T> void minScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream)
{
- transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::device::bind2nd(minimum<T>(), src2), WithOutMask(), stream);
+ transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::cuda::bind2nd(minimum<T>(), src2), WithOutMask(), stream);
}
template void minScalar<uchar >(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream);
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits< arithm::VMax4 > : arithm::ArithmFuncTraits<sizeof(uint), sizeof(uint)>
{
template <typename T> void maxScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream)
{
- transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::device::bind2nd(maximum<T>(), src2), WithOutMask(), stream);
+ transform((PtrStepSz<T>) src1, (PtrStepSz<T>) dst, cv::gpu::cuda::bind2nd(maximum<T>(), src2), WithOutMask(), stream);
}
template void maxScalar<uchar >(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream);
//////////////////////////////////////////////////////////////////////////
// threshold
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< thresh_binary_func<T> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T> struct TransformFunctorTraits< arithm::PowOp<T> > : arithm::ArithmFuncTraits<sizeof(T), sizeof(T)>
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <typename T1, typename T2, typename D, size_t src1_size, size_t src2_size, size_t dst_size> struct AddWeightedTraits : DefaultTransformFunctorTraits< arithm::AddWeighted<T1, T2, D> >
{
#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
{
#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
{
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
-#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
{
#include <thrust/device_ptr.h>
#include <thrust/remove.h>
#include <thrust/functional.h>
-#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];
#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
{
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
template <> struct TransformFunctorTraits<hist::EqualizeHist> : DefaultTransformFunctorTraits<hist::EqualizeHist>
{
#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
void resize_8UC1(const PtrStepSzb& src, PtrStepSzb dst) { resize_for_hog<uchar> (src, dst, resize8UC1_tex); }
void resize_8UC4(const PtrStepSzb& src, PtrStepSzb dst) { resize_for_hog<uchar4>(src, dst, resize8UC4_tex); }
} // namespace hog
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
-#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
{
cudaSafeCall( cudaDeviceSynchronize() );
thrust::device_ptr<int> sizesPtr(sizes);
- thrust::transform(sizesPtr, sizesPtr + levels + 1, sizesPtr, device::bind2nd(device::minimum<int>(), maxSize));
+ thrust::transform(sizesPtr, sizesPtr + levels + 1, sizesPtr, cuda::bind2nd(cuda::minimum<int>(), maxSize));
}
void GHT_Guil_Full_buildTemplFeatureList_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
#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
{
template void filter2D_gpu<float, float>(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<float4, float4>(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 */
#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
{
#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
{
#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
{
#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 {
#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
{
cudaSafeCall( cudaDeviceSynchronize() );
}
} //namespace match_template
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
{
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 */
#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
{
template <int BLOCK_SIZE, typename R>
static __device__ __forceinline__ thrust::tuple<volatile R*, volatile R*> 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 <typename R>
template <int BLOCK_SIZE, typename R>
static __device__ __forceinline__ thrust::tuple<volatile R*, volatile R*, volatile R*> 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 <typename R>
template <int BLOCK_SIZE, typename R>
static __device__ __forceinline__ thrust::tuple<volatile R*, volatile R*, volatile R*, volatile R*> 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 <typename R>
{
sum = tid < gridDim.x * gridDim.y ? result[tid] : VecTraits<result_type>::all(0);
- device::reduce<BLOCK_SIZE>(detail::Unroll<cn>::template smem_tuple<BLOCK_SIZE>(smem), detail::Unroll<cn>::tie(sum), tid, detail::Unroll<cn>::op(plus<R>()));
+ cuda::reduce<BLOCK_SIZE>(detail::Unroll<cn>::template smem_tuple<BLOCK_SIZE>(smem), detail::Unroll<cn>::tie(sum), tid, detail::Unroll<cn>::op(plus<R>()));
if (tid == 0)
{
}
}
- device::reduce<BLOCK_SIZE>(detail::Unroll<cn>::template smem_tuple<BLOCK_SIZE>(smem), detail::Unroll<cn>::tie(sum), tid, detail::Unroll<cn>::op(plus<R>()));
+ cuda::reduce<BLOCK_SIZE>(detail::Unroll<cn>::template smem_tuple<BLOCK_SIZE>(smem), detail::Unroll<cn>::tie(sum), tid, detail::Unroll<cn>::op(plus<R>()));
GlobalReduce<BLOCK_SIZE, R, cn>::run(sum, result, tid, bid, smem);
}
const minimum<R> minOp;
const maximum<R> maxOp;
- device::reduce<BLOCK_SIZE>(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax), tid, thrust::make_tuple(minOp, maxOp));
+ cuda::reduce<BLOCK_SIZE>(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax), tid, thrust::make_tuple(minOp, maxOp));
if (tid == 0)
{
}
}
- device::reduce<BLOCK_SIZE>(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax), tid, thrust::make_tuple(minOp, maxOp));
+ cuda::reduce<BLOCK_SIZE>(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax), tid, thrust::make_tuple(minOp, maxOp));
GlobalReduce<BLOCK_SIZE, R>::run(mymin, mymax, minval, maxval, tid, bid, sminval, smaxval);
}
}
}
- device::reduce<BLOCK_SIZE>(scount, mycount, tid, plus<unsigned int>());
+ cuda::reduce<BLOCK_SIZE>(scount, mycount, tid, plus<unsigned int>());
#if __CUDA_ARCH__ >= 200
if (tid == 0)
{
mycount = tid < gridDim.x * gridDim.y ? count[tid] : 0;
- device::reduce<BLOCK_SIZE>(scount, mycount, tid, plus<unsigned int>());
+ cuda::reduce<BLOCK_SIZE>(scount, mycount, tid, plus<unsigned int>());
if (tid == 0)
{
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;
for (int x = threadIdx.x; x < src.cols; x += BLOCK_SIZE)
myVal = op(myVal, saturate_cast<work_type>(srcRow[x]));
- device::reduce<BLOCK_SIZE>(detail::Unroll<cn>::template smem_tuple<BLOCK_SIZE>(smem), detail::Unroll<cn>::tie(myVal), threadIdx.x, detail::Unroll<cn>::op(op));
+ cuda::reduce<BLOCK_SIZE>(detail::Unroll<cn>::template smem_tuple<BLOCK_SIZE>(smem), detail::Unroll<cn>::tie(myVal), threadIdx.x, detail::Unroll<cn>::op(op));
if (threadIdx.x == 0)
dst[y] = saturate_cast<dst_type>(op.result(myVal, src.cols));
#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;
//////////////////////////////////////////////////////////////////////////////////
//// Non Local Means Denosing
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
//////////////////////////////////////////////////////////////////////////////////
//// Non Local Means Denosing (fast approximate version)
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
template <int BLOCK_SIZE>
static __device__ __forceinline__ thrust::tuple<volatile float*, volatile float*> 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<float&, float&> tie(float& val1, float& val2)
template <int BLOCK_SIZE>
static __device__ __forceinline__ thrust::tuple<volatile float*, volatile float*, volatile float*> 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<float&, float&, float&> tie(float& val1, float2& val2)
template <int BLOCK_SIZE>
static __device__ __forceinline__ thrust::tuple<volatile float*, volatile float*, volatile float*, volatile float*> 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<float&, float&, float&, float&> tie(float& val1, float3& val2)
template <int BLOCK_SIZE>
static __device__ __forceinline__ thrust::tuple<volatile float*, volatile float*, volatile float*, volatile float*, volatile float*> 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<float&, float&, float&, float&, float&> tie(float& val1, float4& val2)
#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
{
#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
{
#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
#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];
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 */
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
-#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
{
#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
{
template void pyrDown_gpu<float3>(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream);
template void pyrDown_gpu<float4>(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream);
} // namespace imgproc
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
{
template void pyrUp_gpu<float3>(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream);
template void pyrUp_gpu<float4>(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream);
} // namespace imgproc
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
{
#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
{
template void remap_gpu<float3>(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<float4>(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 */
#if !defined CUDA_DISABLER
#include <cfloat>
-#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
{
};
} // namespace imgproc
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
{
//
//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
{
#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
{
split_func(src, dst, stream);
}
} // namespace split_merge
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
{
cudaSafeCall( cudaUnbindTexture (texForTF) );
}
} // namespace stereobm
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
{
template <typename T>
__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<float>::max();
+ float minimum = cuda::numeric_limits<float>::max();
for(int i = 0; i < cndisp; ++i)
{
template void output_gpu<short>(const PtrStepSzb& u, const PtrStepSzb& d, const PtrStepSzb& l, const PtrStepSzb& r, const PtrStepSzb& data, const PtrStepSz<short>& disp, cudaStream_t stream);
template void output_gpu<float>(const PtrStepSzb& u, const PtrStepSzb& d, const PtrStepSzb& l, const PtrStepSzb& r, const PtrStepSzb& data, const PtrStepSz<short>& disp, cudaStream_t stream);
} // namespace stereobp
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
{
for(int i = 0; i < nr_plane; i++)
{
- T minimum = device::numeric_limits<T>::max();
+ T minimum = cuda::numeric_limits<T>::max();
int id = 0;
for(int d = 0; d < cndisp; d++)
{
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<short>& disp, int nr_plane, cudaStream_t stream);
} // namespace stereocsbp
-}}} // namespace cv { namespace gpu { namespace device {
+}}} // namespace cv { namespace gpu { namespace cuda {
#endif /* CUDA_DISABLER */
};
}
- namespace device
+ namespace cuda
{
using cv::gpu::TextureBinder;
}
#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
#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
{
template void warpPerspective_gpu<float3>(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<float4>(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 */
#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);
//////////////////////////////////////////////////////////////////////////////////
//// Non Local Means Denosing (brute force)
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
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<uchar>, nlm_bruteforce_gpu<uchar2>, nlm_bruteforce_gpu<uchar3>, 0/*nlm_bruteforce_gpu<uchar4>,*/ };
//// Non Local Means Denosing (fast approxinate)
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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<uchar>, nlm_fast_gpu<uchar2>, nlm_fast_gpu<uchar3>, 0};
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);
}
keypoints.cols = getKeyPoints(keypoints);
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace fast
{
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()));
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;
////////////////////////////////////////////////////////////////////////////////////////////////////
// Linear Filter
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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;
Ptr<BaseFilter_GPU> 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);
#else /* !defined (HAVE_CUDA) */
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace gfft
{
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()));
#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);
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);
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<float>(), mr.ptr<float>(), mapx, mapy);
}
#else /* !defined (HAVE_CUDA) */
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace ccl
{
static const func_t suppotLookup[8][4] =
{ // 1, 2, 3, 4
- { device::ccl::computeEdges<uchar>, 0, device::ccl::computeEdges<uchar3>, device::ccl::computeEdges<uchar4> },// CV_8U
+ { cuda::ccl::computeEdges<uchar>, 0, cuda::ccl::computeEdges<uchar3>, cuda::ccl::computeEdges<uchar4> },// CV_8U
{ 0, 0, 0, 0 },// CV_16U
- { device::ccl::computeEdges<ushort>, 0, device::ccl::computeEdges<ushort3>, device::ccl::computeEdges<ushort4> },// CV_8S
+ { cuda::ccl::computeEdges<ushort>, 0, cuda::ccl::computeEdges<ushort3>, cuda::ccl::computeEdges<ushort4> },// CV_8S
{ 0, 0, 0, 0 },// CV_16S
- { device::ccl::computeEdges<int>, 0, 0, 0 },// CV_32S
- { device::ccl::computeEdges<float>, 0, 0, 0 },// CV_32F
+ { cuda::ccl::computeEdges<int>, 0, 0, 0 },// CV_32S
+ { cuda::ccl::computeEdges<float>, 0, 0, 0 },// CV_32F
{ 0, 0, 0, 0 },// CV_64F
{ 0, 0, 0, 0 } // CV_USRTYPE1
};
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
#else
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace hog
{
}
}}}
-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_)
#include "opencv2/core/utility.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace hough
{
//////////////////////////////////////////////////////////
// HoughLines
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace hough
{
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<unsigned short>::max());
//////////////////////////////////////////////////////////
// HoughLinesP
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace hough
{
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<unsigned short>::max() );
//////////////////////////////////////////////////////////
// HoughCircles
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace hough
{
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<unsigned short>::max());
//////////////////////////////////////////////////////////
// GeneralizedHough
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace hough
{
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[] =
void GHT_Ballard_Pos::processTempl()
{
- using namespace cv::gpu::device::hough;
+ using namespace cv::gpu::cuda::hough;
CV_Assert(levels > 0);
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);
void GHT_Ballard_Pos::findPosInHist()
{
- using namespace cv::gpu::device::hough;
+ using namespace cv::gpu::cuda::hough;
CV_Assert(votesThreshold > 0);
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);
void GHT_Ballard_PosScale::findPosInHist()
{
- using namespace cv::gpu::device::hough;
+ using namespace cv::gpu::cuda::hough;
CV_Assert(votesThreshold > 0);
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);
void GHT_Ballard_PosRotation::findPosInHist()
{
- using namespace cv::gpu::device::hough;
+ using namespace cv::gpu::cuda::hough;
CV_Assert(votesThreshold > 0);
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,
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);
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);
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);
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<int>(), imageFeatures.sizes.ptr<int>(0),
////////////////////////////////////////////////////////////////////////
// meanShiftFiltering_GPU
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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" );
////////////////////////////////////////////////////////////////////////
// meanShiftProc_GPU
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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" );
////////////////////////////////////////////////////////////////////////
// drawColorDisp
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
template <typename T>
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);
////////////////////////////////////////////////////////////////////////
// reprojectImageTo3D
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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] =
////////////////////////////////////////////////////////////////////////
// copyMakeBorder
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
{
template <typename T, int cn> 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_<T> val(saturate_cast<T>(value[0]), saturate_cast<T>(value[1]), saturate_cast<T>(value[2]), saturate_cast<T>(value[3]));
//////////////////////////////////////////////////////////////////////////////
// buildWarpPlaneMaps
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
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<float>(), R_Kinv.ptr<float>(),
+ cuda::imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(),
T.ptr<float>(), scale, StreamAccessor::getStream(stream));
}
//////////////////////////////////////////////////////////////////////////////
// buildWarpCylyndricalMaps
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
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<float>(), R_Kinv.ptr<float>(), scale, StreamAccessor::getStream(stream));
+ cuda::imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(), scale, StreamAccessor::getStream(stream));
}
//////////////////////////////////////////////////////////////////////////////
// buildWarpSphericalMaps
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
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<float>(), R_Kinv.ptr<float>(), scale, StreamAccessor::getStream(stream));
+ cuda::imgproc::buildWarpSphericalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(), scale, StreamAccessor::getStream(stream));
}
////////////////////////////////////////////////////////////////////////
integralBuffered(src, sum, buffer, s);
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
{
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)
//////////////////////////////////////////////////////////////////////////////
// columnSum
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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)
////////////////////////////////////////////////////////////////////////
// cornerHarris & minEgenVal
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
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);
//////////////////////////////////////////////////////////////////////////////
// mulSpectrums
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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<cufftComplex>, const PtrStep<cufftComplex>, PtrStepSz<cufftComplex>, 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());
//////////////////////////////////////////////////////////////////////////////
// mulAndScaleSpectrums
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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<cufftComplex>, const PtrStep<cufftComplex>, float scale, PtrStepSz<cufftComplex>, 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());
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();
#else
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace match_template
{
}
}}}
-using namespace ::cv::gpu::device::match_template;
+using namespace ::cv::gpu::cuda::match_template;
namespace
{
#include "NPP_staging/NPP_staging.hpp"
#include "NCVBroxOpticalFlow.hpp"
-#include "opencv2/core/device/utility.hpp"
+#include "opencv2/core/cuda/utility.hpp"
typedef NCVVectorAlloc<Ncv32f> FloatVector;
ScaleVector(ptrVNew->ptr(), ptrVNew->ptr(), 1.0f/scale_factor, ns * nh, stream);
ncvAssertCUDALastErrorReturn((int)NCV_CUDA_ERROR);
- cv::gpu::device::swap<FloatVector*>(ptrU, ptrUNew);
- cv::gpu::device::swap<FloatVector*>(ptrV, ptrVNew);
+ cv::gpu::cuda::swap<FloatVector*>(ptrU, ptrUNew);
+ cv::gpu::cuda::swap<FloatVector*>(ptrV, ptrVNew);
}
scale /= scale_factor;
}
#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"
//==============================================================================
__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;
}
#include <vector>
#include <cuda_runtime.h>
#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<Ncv8u, 1, cudaReadModeElementType> tex8u;
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;
}
#include "NCVAlg.hpp"
#include "NCVPyramid.hpp"
#include "NCVPixelOperations.hpp"
-#include "opencv2/core/device/common.hpp"
+#include "opencv2/core/cuda/common.hpp"
template<typename T, Ncv32u CN> struct __average4_CN {static __host__ __device__ T _average4_CN(const T &p00, const T &p01, const T &p10, const T &p11);};
}
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace pyramid
{
d_dst_line[j] = outPix;
}
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace pyramid
{
cudaSafeCall( cudaDeviceSynchronize() );
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace optical_flow
{
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());
#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,
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(
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<float>(ig11), static_cast<float>(ig03), static_cast<float>(ig33), static_cast<float>(ig55));
+ cuda::optflow_farneback::setPolynomialExpansionConsts(n, g, xg, xxg, static_cast<float>(ig11), static_cast<float>(ig03), static_cast<float>(ig33), static_cast<float>(ig55));
}
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]));
}
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]));
}
}
setPolynomialExpansionConsts(polyN, polySigma);
- device::optflow_farneback::setUpdateMatricesConsts();
+ cuda::optflow_farneback::setUpdateMatricesConsts();
for (int k = numLevelsCropped; k >= 0; k--)
{
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
{
};
Mat g = getGaussianKernel(smoothSize, sigma, CV_32F);
- device::optflow_farneback::setGaussianBlurKernel(g.ptr<float>(smoothSize/2), smoothSize/2);
+ cuda::optflow_farneback::setGaussianBlurKernel(g.ptr<float>(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]);
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<float>(winSize/2), winSize/2);
+ cuda::optflow_farneback::setGaussianBlurKernel(g.ptr<float>(winSize/2), winSize/2);
}
for (int i = 0; i < numIters; i++)
{
#else /* !defined (HAVE_CUDA) */
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace orb
{
++v_0;
}
CV_Assert(u_max.size() < 32);
- cv::gpu::device::orb::loadUMax(&u_max[0], static_cast<int>(u_max.size()));
+ cv::gpu::cuda::orb::loadUMax(&u_max[0], static_cast<int>(u_max.size()));
// Calc pattern
const int npoints = 512;
//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)
void cv::gpu::ORB_GPU::computeKeyPointsPyramid()
{
- using namespace cv::gpu::device::orb;
+ using namespace cv::gpu::cuda::orb;
int half_patch_size = patchSize_ / 2;
void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors)
{
- using namespace cv::gpu::device::orb;
+ using namespace cv::gpu::cuda::orb;
int nAllkeypoints = 0;
void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat& keypoints)
{
- using namespace cv::gpu::device::orb;
+ using namespace cv::gpu::cuda::orb;
int nAllkeypoints = 0;
//////////////////////////////////////////////////////////////////////////////
// pyrDown
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
//////////////////////////////////////////////////////////////////////////////
// pyrUp
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
//////////////////////////////////////////////////////////////////////////////
// ImagePyramid
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace pyramid
{
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);
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);
#else // HAVE_CUDA
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
#else // HAVE_CUDA
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
}
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);
#else /* !defined (HAVE_CUDA) */
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace split_merge
{
{
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);
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);
#else /* !defined (HAVE_CUDA) */
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace stereobm
{
{
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);
#else /* !defined (HAVE_CUDA) */
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace stereobp
{
}
}}}
-using namespace ::cv::gpu::device::stereobp;
+using namespace ::cv::gpu::cuda::stereobp;
namespace
{
#else /* !defined (HAVE_CUDA) */
#include "opencv2/core/utility.hpp"
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace stereocsbp
{
}
}}}
-using namespace ::cv::gpu::device::stereocsbp;
+using namespace ::cv::gpu::cuda::stereocsbp;
namespace
{
videoSource_->stop();
}
-namespace cv { namespace gpu { namespace device {
+namespace cv { namespace gpu { namespace cuda {
namespace video_decoding
{
void loadHueCSC(float hueCSC[9]);
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];
CV_Assert( err == 0 );
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace video_encoding
{
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_)
#else // HAVE_CUDA
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace imgproc
{
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);
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);
}
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);
}
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);
#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
{
}
}}}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace surf
{
}
plus<float> 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;
}
}
} // namespace surf
-}}} // namespace cv { namespace gpu { namespace device
+}}} // namespace cv { namespace gpu { namespace cuda
#endif /* CUDA_DISABLER */
#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
{
}
}}}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace vibe
{
#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
#else // !defined (HAVE_CUDA)
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace surf
{
}
}}}
-using namespace ::cv::gpu::device::surf;
+using namespace ::cv::gpu::cuda::surf;
namespace
{
#else
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
namespace vibe
{
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);
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);
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;
#define CV_PI 3.1415926535897932384626433832795
#endif
-namespace cv { namespace softcascade { namespace device {
+namespace cv { namespace softcascade { namespace cuda {
typedef unsigned char uchar;
#endif
-namespace cv { namespace softcascade { namespace device {
+namespace cv { namespace softcascade { namespace cuda {
typedef unsigned char uchar;
typedef unsigned int uint;
}
}
-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);
}
}
-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);
FileNode fn = root[SC_OCTAVES];
if (fn.empty()) return 0;
- std::vector<device::Octave> voctaves;
+ std::vector<cuda::Octave> voctaves;
std::vector<float> vstages;
- std::vector<device::Node> vnodes;
+ std::vector<cuda::Node> vnodes;
std::vector<float> vleaves;
FileNodeIterator it = fn.begin(), it_end = fn.end();
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);
rect.w = saturate_cast<uchar>(r.height);
unsigned int channel = saturate_cast<unsigned int>(feature_channels[featureIdx]);
- vnodes.push_back(device::Node(rect, channel, th));
+ vnodes.push_back(cuda::Node(rect, channel, th));
}
intfns = octfn[SC_LEAF];
}
}
- 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));
int createLevels(const int fh, const int fw)
{
- std::vector<device::Level> vlevels;
+ std::vector<cuda::Level> vlevels;
float logFactor = (::log(maxScale) - ::log(minScale)) / (totals -1);
float scale = minScale;
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;
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;
cudaSafeCall( cudaGetLastError());
- device::CascadeInvoker<device::GK107PolicyX4> invoker
- = device::CascadeInvoker<device::GK107PolicyX4>(levels, stages, nodes, leaves);
+ cuda::CascadeInvoker<cuda::GK107PolicyX4> invoker
+ = cuda::CascadeInvoker<cuda::GK107PolicyX4>(levels, stages, nodes, leaves);
cudaStream_t stream = cv::gpu::StreamAccessor::getStream(s);
invoker(mask, hogluv, objects, downscales, stream);
}
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<device::Octave>::const_iterator octIt_t;
- static int fitOctave(const std::vector<device::Octave>& octs, const float& logFactor)
+ typedef std::vector<cuda::Octave>::const_iterator octIt_t;
+ static int fitOctave(const std::vector<cuda::Octave>& 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);
// cv::gpu::GpuMat collected;
- std::vector<device::Octave> voctaves;
+ std::vector<cuda::Octave> voctaves;
// DeviceInfo info;
{
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)
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)
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:
#else // HAVE_CUDA
-namespace btv_l1_device
+namespace btv_l1_cuda
{
void buildMotionMaps(PtrStepSzf forwardMotionX, PtrStepSzf forwardMotionY,
PtrStepSzf backwardMotionX, PtrStepSzf bacwardMotionY,
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);
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 );
{
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<float>& btvWeights)
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)
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());
//
//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,
template <int cn> 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,
};
}
-namespace cv { namespace gpu { namespace device
+namespace cv { namespace gpu { namespace cuda
{
- template <> struct TransformFunctorTraits<btv_l1_device::DiffSign> : DefaultTransformFunctorTraits<btv_l1_device::DiffSign>
+ template <> struct TransformFunctorTraits<btv_l1_cuda::DiffSign> : DefaultTransformFunctorTraits<btv_l1_cuda::DiffSign>
{
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)
{
projects / platform / upstream / opencv.git / commitdiff