diff --git a/modules/core/include/opencv2/core/cuda.hpp b/modules/core/include/opencv2/core/cuda.hpp index 15d526e802..ef554ecf0b 100644 --- a/modules/core/include/opencv2/core/cuda.hpp +++ b/modules/core/include/opencv2/core/cuda.hpp @@ -51,9 +51,20 @@ #include "opencv2/core.hpp" #include "opencv2/core/cuda_types.hpp" +/** + @defgroup cuda CUDA-accelerated Computer Vision + @{ + @defgroup cudacore Core part + @{ + @defgroup cudacore_init Initalization and Information + @defgroup cudacore_struct Data Structures + @} + @} + */ + namespace cv { namespace cuda { -//! @addtogroup cuda_struct +//! @addtogroup cudacore_struct //! @{ //////////////////////////////// GpuMat /////////////////////////////// @@ -514,11 +525,11 @@ private: friend struct EventAccessor; }; -//! @} cuda_struct +//! @} cudacore_struct //////////////////////////////// Initialization & Info //////////////////////// -//! @addtogroup cuda_init +//! @addtogroup cudacore_init //! @{ /** @brief Returns the number of installed CUDA-enabled devices. @@ -813,7 +824,7 @@ private: CV_EXPORTS void printCudaDeviceInfo(int device); CV_EXPORTS void printShortCudaDeviceInfo(int device); -//! @} cuda_init +//! @} cudacore_init }} // namespace cv { namespace cuda { diff --git a/modules/core/include/opencv2/core/cuda/block.hpp b/modules/core/include/opencv2/core/cuda/block.hpp index 8377adf213..0c6f0636bb 100644 --- a/modules/core/include/opencv2/core/cuda/block.hpp +++ b/modules/core/include/opencv2/core/cuda/block.hpp @@ -43,11 +43,14 @@ #ifndef __OPENCV_CUDA_DEVICE_BLOCK_HPP__ #define __OPENCV_CUDA_DEVICE_BLOCK_HPP__ +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ struct Block { static __device__ __forceinline__ unsigned int id() @@ -201,7 +204,8 @@ namespace cv { namespace cuda { namespace device } } }; -//!@} }}} +//! @endcond + #endif /* __OPENCV_CUDA_DEVICE_BLOCK_HPP__ */ diff --git a/modules/core/include/opencv2/core/cuda/border_interpolate.hpp b/modules/core/include/opencv2/core/cuda/border_interpolate.hpp index 3d639b4716..ba7266918c 100644 --- a/modules/core/include/opencv2/core/cuda/border_interpolate.hpp +++ b/modules/core/include/opencv2/core/cuda/border_interpolate.hpp @@ -47,11 +47,14 @@ #include "vec_traits.hpp" #include "vec_math.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ - ////////////////////////////////////////////////////////////// // BrdConstant @@ -712,7 +715,8 @@ namespace cv { namespace cuda { namespace device int width; D val; }; -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/color.hpp b/modules/core/include/opencv2/core/cuda/color.hpp index 524d2e2a47..6faf8c9c5c 100644 --- a/modules/core/include/opencv2/core/cuda/color.hpp +++ b/modules/core/include/opencv2/core/cuda/color.hpp @@ -45,10 +45,14 @@ #include "detail/color_detail.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ // All OPENCV_CUDA_IMPLEMENT_*_TRAITS(ColorSpace1_to_ColorSpace2, ...) macros implements // template class ColorSpace1_to_ColorSpace2_traits // { @@ -298,7 +302,8 @@ namespace cv { namespace cuda { namespace device OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lbgra, 4, 4, false, 0) #undef OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/common.hpp b/modules/core/include/opencv2/core/cuda/common.hpp index 022b13ebb6..b93c3efdcc 100644 --- a/modules/core/include/opencv2/core/cuda/common.hpp +++ b/modules/core/include/opencv2/core/cuda/common.hpp @@ -48,6 +48,11 @@ #include "opencv2/core/cvdef.h" #include "opencv2/core/base.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED #ifndef CV_PI_F #ifndef CV_PI @@ -58,14 +63,11 @@ #endif namespace cv { namespace cuda { -//! @addtogroup cuda -//! @{ static inline void checkCudaError(cudaError_t err, const char* file, const int line, const char* func) { if (cudaSuccess != err) cv::error(cv::Error::GpuApiCallError, cudaGetErrorString(err), func, file, line); } -//! @} }} #ifndef cudaSafeCall @@ -74,8 +76,6 @@ namespace cv { namespace cuda { namespace cv { namespace cuda { -//! @addtogroup cuda -//! @{ template static inline bool isAligned(const T* ptr, size_t size) { return reinterpret_cast(ptr) % size == 0; @@ -85,15 +85,12 @@ namespace cv { namespace cuda { return step % size == 0; } -//! @} }} namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ __host__ __device__ __forceinline__ int divUp(int total, int grain) { return (total + grain - 1) / grain; @@ -104,8 +101,9 @@ namespace cv { namespace cuda cudaChannelFormatDesc desc = cudaCreateChannelDesc(); cudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) ); } -//! @} } }} +//! @endcond + #endif // __OPENCV_CUDA_COMMON_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/datamov_utils.hpp b/modules/core/include/opencv2/core/cuda/datamov_utils.hpp index 4d4035a3a3..bb02cf92df 100644 --- a/modules/core/include/opencv2/core/cuda/datamov_utils.hpp +++ b/modules/core/include/opencv2/core/cuda/datamov_utils.hpp @@ -45,11 +45,14 @@ #include "common.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ - #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 // for Fermi memory space is detected automatically @@ -103,7 +106,8 @@ namespace cv { namespace cuda { namespace device #undef OPENCV_CUDA_ASM_PTR #endif // __CUDA_ARCH__ >= 200 -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif // __OPENCV_CUDA_DATAMOV_UTILS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp b/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp index 8d71532ae7..34884638a9 100644 --- a/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp +++ b/modules/core/include/opencv2/core/cuda/dynamic_smem.hpp @@ -43,10 +43,14 @@ #ifndef __OPENCV_CUDA_DYNAMIC_SMEM_HPP__ #define __OPENCV_CUDA_DYNAMIC_SMEM_HPP__ +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template struct DynamicSharedMem { __device__ __forceinline__ operator T*() @@ -77,7 +81,8 @@ namespace cv { namespace cuda { namespace device return (double*)__smem_d; } }; -//! @} }}} +//! @endcond + #endif // __OPENCV_CUDA_DYNAMIC_SMEM_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/emulation.hpp b/modules/core/include/opencv2/core/cuda/emulation.hpp index 3063266b77..d3468651f6 100644 --- a/modules/core/include/opencv2/core/cuda/emulation.hpp +++ b/modules/core/include/opencv2/core/cuda/emulation.hpp @@ -46,10 +46,14 @@ #include "common.hpp" #include "warp_reduce.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ struct Emulation { @@ -258,7 +262,8 @@ namespace cv { namespace cuda { namespace device } }; }; //struct Emulation -//!@} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif /* OPENCV_CUDA_EMULATION_HPP_ */ diff --git a/modules/core/include/opencv2/core/cuda/filters.hpp b/modules/core/include/opencv2/core/cuda/filters.hpp index 8bec9fe261..9adc00c444 100644 --- a/modules/core/include/opencv2/core/cuda/filters.hpp +++ b/modules/core/include/opencv2/core/cuda/filters.hpp @@ -48,10 +48,14 @@ #include "vec_math.hpp" #include "type_traits.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template struct PointFilter { typedef typename Ptr2D::elem_type elem_type; @@ -275,7 +279,8 @@ namespace cv { namespace cuda { namespace device float scale_x, scale_y; int width, haight; }; -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif // __OPENCV_CUDA_FILTERS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/funcattrib.hpp b/modules/core/include/opencv2/core/cuda/funcattrib.hpp index fa2039ab2b..fbb236b9e6 100644 --- a/modules/core/include/opencv2/core/cuda/funcattrib.hpp +++ b/modules/core/include/opencv2/core/cuda/funcattrib.hpp @@ -45,10 +45,14 @@ #include +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template void printFuncAttrib(Func& func) { @@ -68,7 +72,8 @@ namespace cv { namespace cuda { namespace device printf("\n"); fflush(stdout); } -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif /* __OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP_ */ diff --git a/modules/core/include/opencv2/core/cuda/functional.hpp b/modules/core/include/opencv2/core/cuda/functional.hpp index 3060bbc680..ed3943da45 100644 --- a/modules/core/include/opencv2/core/cuda/functional.hpp +++ b/modules/core/include/opencv2/core/cuda/functional.hpp @@ -49,10 +49,14 @@ #include "type_traits.hpp" #include "device_functions.h" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ // Function Objects template struct unary_function : public std::unary_function {}; template struct binary_function : public std::binary_function {}; @@ -786,7 +790,8 @@ namespace cv { namespace cuda { namespace device #define OPENCV_CUDA_TRANSFORM_FUNCTOR_TRAITS(type) \ template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type > -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif // __OPENCV_CUDA_FUNCTIONAL_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/limits.hpp b/modules/core/include/opencv2/core/cuda/limits.hpp index e22d99d00b..b98bdf2294 100644 --- a/modules/core/include/opencv2/core/cuda/limits.hpp +++ b/modules/core/include/opencv2/core/cuda/limits.hpp @@ -47,10 +47,14 @@ #include #include "common.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template struct numeric_limits; template <> struct numeric_limits @@ -117,7 +121,8 @@ template <> struct numeric_limits __device__ __forceinline__ static double epsilon() { return DBL_EPSILON; } static const bool is_signed = true; }; -//! @} }}} // namespace cv { namespace cuda { namespace cudev { +//! @endcond + #endif // __OPENCV_CUDA_LIMITS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/reduce.hpp b/modules/core/include/opencv2/core/cuda/reduce.hpp index 0f18b1e404..3133c9a89e 100644 --- a/modules/core/include/opencv2/core/cuda/reduce.hpp +++ b/modules/core/include/opencv2/core/cuda/reduce.hpp @@ -47,10 +47,14 @@ #include "detail/reduce.hpp" #include "detail/reduce_key_val.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template __device__ __forceinline__ void reduce(volatile T* smem, T& val, unsigned int tid, const Op& op) { @@ -194,7 +198,8 @@ namespace cv { namespace cuda { namespace device { return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3, (volatile T4*) t4, (volatile T5*) t5, (volatile T6*) t6, (volatile T7*) t7, (volatile T8*) t8, (volatile T9*) t9); } -//! @} }}} +//! @endcond + #endif // __OPENCV_CUDA_UTILITY_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/saturate_cast.hpp b/modules/core/include/opencv2/core/cuda/saturate_cast.hpp index 46927a5097..e7633c76a3 100644 --- a/modules/core/include/opencv2/core/cuda/saturate_cast.hpp +++ b/modules/core/include/opencv2/core/cuda/saturate_cast.hpp @@ -45,10 +45,14 @@ #include "common.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template __device__ __forceinline__ _Tp saturate_cast(uchar v) { return _Tp(v); } template __device__ __forceinline__ _Tp saturate_cast(schar v) { return _Tp(v); } template __device__ __forceinline__ _Tp saturate_cast(ushort v) { return _Tp(v); } @@ -281,7 +285,8 @@ namespace cv { namespace cuda { namespace device return saturate_cast((float)v); #endif } -//! @} }}} +//! @endcond + #endif /* __OPENCV_CUDA_SATURATE_CAST_HPP__ */ diff --git a/modules/core/include/opencv2/core/cuda/scan.hpp b/modules/core/include/opencv2/core/cuda/scan.hpp index e1dbb0059f..687abb508b 100644 --- a/modules/core/include/opencv2/core/cuda/scan.hpp +++ b/modules/core/include/opencv2/core/cuda/scan.hpp @@ -48,10 +48,14 @@ #include "opencv2/core/cuda/warp.hpp" #include "opencv2/core/cuda/warp_shuffle.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ enum ScanKind { EXCLUSIVE = 0, INCLUSIVE = 1 }; template struct WarpScan @@ -247,7 +251,8 @@ namespace cv { namespace cuda { namespace device return warpScanInclusive(idata, s_Data, tid); } } -//! @} }}} +//! @endcond + #endif // __OPENCV_CUDA_SCAN_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/simd_functions.hpp b/modules/core/include/opencv2/core/cuda/simd_functions.hpp index d6b7435419..b9e0041a0a 100644 --- a/modules/core/include/opencv2/core/cuda/simd_functions.hpp +++ b/modules/core/include/opencv2/core/cuda/simd_functions.hpp @@ -76,57 +76,13 @@ #include "common.hpp" /** @file - This header file contains inline functions that implement intra-word SIMD - operations, that are hardware accelerated on sm_3x (Kepler) GPUs. Efficient - emulation code paths are provided for earlier architectures (sm_1x, sm_2x) - to make the code portable across all GPUs supported by CUDA. The following - functions are currently implemented: + * @deprecated Use @ref cudev instead. + */ - vadd2(a,b) per-halfword unsigned addition, with wrap-around: a + b - vsub2(a,b) per-halfword unsigned subtraction, with wrap-around: a - b - vabsdiff2(a,b) per-halfword unsigned absolute difference: |a - b| - vavg2(a,b) per-halfword unsigned average: (a + b) / 2 - vavrg2(a,b) per-halfword unsigned rounded average: (a + b + 1) / 2 - vseteq2(a,b) per-halfword unsigned comparison: a == b ? 1 : 0 - vcmpeq2(a,b) per-halfword unsigned comparison: a == b ? 0xffff : 0 - vsetge2(a,b) per-halfword unsigned comparison: a >= b ? 1 : 0 - vcmpge2(a,b) per-halfword unsigned comparison: a >= b ? 0xffff : 0 - vsetgt2(a,b) per-halfword unsigned comparison: a > b ? 1 : 0 - vcmpgt2(a,b) per-halfword unsigned comparison: a > b ? 0xffff : 0 - vsetle2(a,b) per-halfword unsigned comparison: a <= b ? 1 : 0 - vcmple2(a,b) per-halfword unsigned comparison: a <= b ? 0xffff : 0 - vsetlt2(a,b) per-halfword unsigned comparison: a < b ? 1 : 0 - vcmplt2(a,b) per-halfword unsigned comparison: a < b ? 0xffff : 0 - vsetne2(a,b) per-halfword unsigned comparison: a != b ? 1 : 0 - vcmpne2(a,b) per-halfword unsigned comparison: a != b ? 0xffff : 0 - vmax2(a,b) per-halfword unsigned maximum: max(a, b) - vmin2(a,b) per-halfword unsigned minimum: min(a, b) - - vadd4(a,b) per-byte unsigned addition, with wrap-around: a + b - vsub4(a,b) per-byte unsigned subtraction, with wrap-around: a - b - vabsdiff4(a,b) per-byte unsigned absolute difference: |a - b| - vavg4(a,b) per-byte unsigned average: (a + b) / 2 - vavrg4(a,b) per-byte unsigned rounded average: (a + b + 1) / 2 - vseteq4(a,b) per-byte unsigned comparison: a == b ? 1 : 0 - vcmpeq4(a,b) per-byte unsigned comparison: a == b ? 0xff : 0 - vsetge4(a,b) per-byte unsigned comparison: a >= b ? 1 : 0 - vcmpge4(a,b) per-byte unsigned comparison: a >= b ? 0xff : 0 - vsetgt4(a,b) per-byte unsigned comparison: a > b ? 1 : 0 - vcmpgt4(a,b) per-byte unsigned comparison: a > b ? 0xff : 0 - vsetle4(a,b) per-byte unsigned comparison: a <= b ? 1 : 0 - vcmple4(a,b) per-byte unsigned comparison: a <= b ? 0xff : 0 - vsetlt4(a,b) per-byte unsigned comparison: a < b ? 1 : 0 - vcmplt4(a,b) per-byte unsigned comparison: a < b ? 0xff : 0 - vsetne4(a,b) per-byte unsigned comparison: a != b ? 1: 0 - vcmpne4(a,b) per-byte unsigned comparison: a != b ? 0xff: 0 - vmax4(a,b) per-byte unsigned maximum: max(a, b) - vmin4(a,b) per-byte unsigned minimum: min(a, b) -*/ +//! @cond IGNORED namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ // 2 static __device__ __forceinline__ unsigned int vadd2(unsigned int a, unsigned int b) @@ -906,7 +862,8 @@ namespace cv { namespace cuda { namespace device return r; } -//! @} }}} +//! @endcond + #endif // __OPENCV_CUDA_SIMD_FUNCTIONS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/transform.hpp b/modules/core/include/opencv2/core/cuda/transform.hpp index bc86bd1e86..08a313df36 100644 --- a/modules/core/include/opencv2/core/cuda/transform.hpp +++ b/modules/core/include/opencv2/core/cuda/transform.hpp @@ -47,10 +47,14 @@ #include "utility.hpp" #include "detail/transform_detail.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template static inline void transform(PtrStepSz src, PtrStepSz dst, UnOp op, const Mask& mask, cudaStream_t stream) { @@ -64,7 +68,8 @@ namespace cv { namespace cuda { namespace device typedef TransformFunctorTraits ft; transform_detail::TransformDispatcher::cn == 1 && VecTraits::cn == 1 && VecTraits::cn == 1 && ft::smart_shift != 1>::call(src1, src2, dst, op, mask, stream); } -//! @} }}} +//! @endcond + #endif // __OPENCV_CUDA_TRANSFORM_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/type_traits.hpp b/modules/core/include/opencv2/core/cuda/type_traits.hpp index a9e9180698..f2471ebab0 100644 --- a/modules/core/include/opencv2/core/cuda/type_traits.hpp +++ b/modules/core/include/opencv2/core/cuda/type_traits.hpp @@ -45,10 +45,14 @@ #include "detail/type_traits_detail.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template struct IsSimpleParameter { enum {value = type_traits_detail::IsIntegral::value || type_traits_detail::IsFloat::value || @@ -79,7 +83,8 @@ namespace cv { namespace cuda { namespace device typedef typename type_traits_detail::Select::value, T, typename type_traits_detail::AddParameterType::type>::type ParameterType; }; -//! @} }}} +//! @endcond + #endif // __OPENCV_CUDA_TYPE_TRAITS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/utility.hpp b/modules/core/include/opencv2/core/cuda/utility.hpp index ad6c7094ac..ed604712a7 100644 --- a/modules/core/include/opencv2/core/cuda/utility.hpp +++ b/modules/core/include/opencv2/core/cuda/utility.hpp @@ -46,10 +46,14 @@ #include "saturate_cast.hpp" #include "datamov_utils.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ #define OPENCV_CUDA_LOG_WARP_SIZE (5) #define OPENCV_CUDA_WARP_SIZE (1 << OPENCV_CUDA_LOG_WARP_SIZE) #define OPENCV_CUDA_LOG_MEM_BANKS ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla @@ -210,7 +214,8 @@ namespace cv { namespace cuda { namespace device return false; } -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif // __OPENCV_CUDA_UTILITY_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/vec_distance.hpp b/modules/core/include/opencv2/core/cuda/vec_distance.hpp index 4bad198e0c..013b747a4d 100644 --- a/modules/core/include/opencv2/core/cuda/vec_distance.hpp +++ b/modules/core/include/opencv2/core/cuda/vec_distance.hpp @@ -47,10 +47,14 @@ #include "functional.hpp" #include "detail/vec_distance_detail.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template struct L1Dist { typedef int value_type; @@ -221,7 +225,8 @@ namespace cv { namespace cuda { namespace device U vec1Vals[MAX_LEN / THREAD_DIM]; }; -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif // __OPENCV_CUDA_VEC_DISTANCE_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/vec_math.hpp b/modules/core/include/opencv2/core/cuda/vec_math.hpp index ac4a493d68..8595fb8d0d 100644 --- a/modules/core/include/opencv2/core/cuda/vec_math.hpp +++ b/modules/core/include/opencv2/core/cuda/vec_math.hpp @@ -46,12 +46,15 @@ #include "vec_traits.hpp" #include "saturate_cast.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ - // saturate_cast namespace vec_math_detail @@ -920,8 +923,8 @@ CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double) #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC -//! @} - }}} // namespace cv { namespace cuda { namespace device +//! @endcond + #endif // __OPENCV_CUDA_VECMATH_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/vec_traits.hpp b/modules/core/include/opencv2/core/cuda/vec_traits.hpp index 6369112ffa..905e37f582 100644 --- a/modules/core/include/opencv2/core/cuda/vec_traits.hpp +++ b/modules/core/include/opencv2/core/cuda/vec_traits.hpp @@ -45,10 +45,14 @@ #include "common.hpp" +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template struct TypeVec; struct __align__(8) uchar8 @@ -277,7 +281,8 @@ namespace cv { namespace cuda { namespace device static __device__ __host__ __forceinline__ char8 make(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7) {return make_char8(a0, a1, a2, a3, a4, a5, a6, a7);} static __device__ __host__ __forceinline__ char8 make(const schar* v) {return make_char8(v[0], v[1], v[2], v[3], v[4], v[5], v[6], v[7]);} }; -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif // __OPENCV_CUDA_VEC_TRAITS_HPP__ diff --git a/modules/core/include/opencv2/core/cuda/warp.hpp b/modules/core/include/opencv2/core/cuda/warp.hpp index 67abd05728..d93afe79c5 100644 --- a/modules/core/include/opencv2/core/cuda/warp.hpp +++ b/modules/core/include/opencv2/core/cuda/warp.hpp @@ -43,10 +43,14 @@ #ifndef __OPENCV_CUDA_DEVICE_WARP_HPP__ #define __OPENCV_CUDA_DEVICE_WARP_HPP__ +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ struct Warp { enum @@ -128,7 +132,8 @@ namespace cv { namespace cuda { namespace device *t = value; } }; -//! @} }}} // namespace cv { namespace cuda { namespace cudev +//! @endcond + #endif /* __OPENCV_CUDA_DEVICE_WARP_HPP__ */ diff --git a/modules/core/include/opencv2/core/cuda/warp_reduce.hpp b/modules/core/include/opencv2/core/cuda/warp_reduce.hpp index 6df04bb028..530303d249 100644 --- a/modules/core/include/opencv2/core/cuda/warp_reduce.hpp +++ b/modules/core/include/opencv2/core/cuda/warp_reduce.hpp @@ -43,10 +43,14 @@ #ifndef OPENCV_CUDA_WARP_REDUCE_HPP__ #define OPENCV_CUDA_WARP_REDUCE_HPP__ +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template __device__ __forceinline__ T warp_reduce(volatile T *ptr , const unsigned int tid = threadIdx.x) { @@ -65,7 +69,8 @@ namespace cv { namespace cuda { namespace device return ptr[tid - lane]; } -//! @} }}} // namespace cv { namespace cuda { namespace cudev { +//! @endcond + #endif /* OPENCV_CUDA_WARP_REDUCE_HPP__ */ diff --git a/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp b/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp index f614946730..5cf42ec41d 100644 --- a/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp +++ b/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp @@ -43,10 +43,14 @@ #ifndef __OPENCV_CUDA_WARP_SHUFFLE_HPP__ #define __OPENCV_CUDA_WARP_SHUFFLE_HPP__ +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + namespace cv { namespace cuda { namespace device { -//! @addtogroup cuda -//! @{ template __device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize) { @@ -142,7 +146,8 @@ namespace cv { namespace cuda { namespace device return 0.0; #endif } -//! @} }}} +//! @endcond + #endif // __OPENCV_CUDA_WARP_SHUFFLE_HPP__ diff --git a/modules/core/include/opencv2/core/cuda_stream_accessor.hpp b/modules/core/include/opencv2/core/cuda_stream_accessor.hpp index 66aaf56c52..dd6589bcb6 100644 --- a/modules/core/include/opencv2/core/cuda_stream_accessor.hpp +++ b/modules/core/include/opencv2/core/cuda_stream_accessor.hpp @@ -47,10 +47,9 @@ # error cuda_stream_accessor.hpp header must be compiled as C++ #endif -// This is only header file that depends on Cuda. All other headers are independent. -// So if you use OpenCV binaries you do noot need to install Cuda Toolkit. -// But of you wanna use CUDA by yourself, may get cuda stream instance using the class below. -// In this case you have to install Cuda Toolkit. +/** @file cuda_stream_accessor.hpp + * This is only header file that depends on CUDA Runtime API. All other headers are independent. + */ #include #include "opencv2/core/cvdef.h" @@ -60,22 +59,21 @@ namespace cv namespace cuda { -//! @addtogroup cuda_struct +//! @addtogroup cudacore_struct //! @{ class Stream; class Event; /** @brief Class that enables getting cudaStream_t from cuda::Stream - - because it is the only public header that depends on the CUDA Runtime API. Including it - brings a dependency to your code. */ struct StreamAccessor { CV_EXPORTS static cudaStream_t getStream(const Stream& stream); }; + /** @brief Class that enables getting cudaEvent_t from cuda::Event + */ struct EventAccessor { CV_EXPORTS static cudaEvent_t getEvent(const Event& event); diff --git a/modules/core/include/opencv2/core/cuda_types.hpp b/modules/core/include/opencv2/core/cuda_types.hpp index 490086fb0a..8df816e8df 100644 --- a/modules/core/include/opencv2/core/cuda_types.hpp +++ b/modules/core/include/opencv2/core/cuda_types.hpp @@ -47,6 +47,12 @@ # error cuda_types.hpp header must be compiled as C++ #endif +/** @file + * @deprecated Use @ref cudev instead. + */ + +//! @cond IGNORED + #ifdef __CUDACC__ #define __CV_CUDA_HOST_DEVICE__ __host__ __device__ __forceinline__ #else @@ -58,9 +64,6 @@ namespace cv namespace cuda { -//! @addtogroup cuda_struct -//! @{ - // Simple lightweight structures that encapsulates information about an image on device. // It is intended to pass to nvcc-compiled code. GpuMat depends on headers that nvcc can't compile @@ -89,17 +92,11 @@ namespace cv size_t size; }; - /** @brief Structure similar to cuda::PtrStepSz but containing only a pointer and row step. - - Width and height fields are excluded due to performance reasons. The structure is intended - for internal use or for users who write device code. - */ template struct PtrStep : public DevPtr { __CV_CUDA_HOST_DEVICE__ PtrStep() : step(0) {} __CV_CUDA_HOST_DEVICE__ PtrStep(T* data_, size_t step_) : DevPtr(data_), step(step_) {} - //! stride between two consecutive rows in bytes. Step is stored always and everywhere in bytes!!! size_t step; __CV_CUDA_HOST_DEVICE__ T* ptr(int y = 0) { return ( T*)( ( char*)DevPtr::data + y * step); } @@ -109,12 +106,6 @@ namespace cv __CV_CUDA_HOST_DEVICE__ const T& operator ()(int y, int x) const { return ptr(y)[x]; } }; - /** @brief Lightweight class encapsulating pitched memory on a GPU and passed to nvcc-compiled code (CUDA - kernels). - - Typically, it is used internally by OpenCV and by users who write device code. You can call - its members from both host and device code. - */ template struct PtrStepSz : public PtrStep { __CV_CUDA_HOST_DEVICE__ PtrStepSz() : cols(0), rows(0) {} @@ -136,9 +127,9 @@ namespace cv typedef PtrStep PtrStepf; typedef PtrStep PtrStepi; -//! @} - } } +//! @endcond + #endif /* __OPENCV_CORE_CUDA_TYPES_HPP__ */ diff --git a/modules/cuda/include/opencv2/cuda.hpp b/modules/cuda/include/opencv2/cuda.hpp index ac51b87dde..93bb511cd0 100644 --- a/modules/cuda/include/opencv2/cuda.hpp +++ b/modules/cuda/include/opencv2/cuda.hpp @@ -50,15 +50,11 @@ #include "opencv2/core/cuda.hpp" /** -@defgroup cuda CUDA-accelerated Computer Vision - @ref cuda_intro "Introduction page" + @addtogroup cuda @{ - @defgroup cuda_init Initalization and Information - @defgroup cuda_struct Data Structures @defgroup cuda_calib3d Camera Calibration and 3D Reconstruction @defgroup cuda_objdetect Object Detection @} - */ namespace cv { namespace cuda { diff --git a/modules/cudev/include/opencv2/cudev/ptr2d/glob.hpp b/modules/cudev/include/opencv2/cudev/ptr2d/glob.hpp index 3563e56fcc..4296dd44bc 100644 --- a/modules/cudev/include/opencv2/cudev/ptr2d/glob.hpp +++ b/modules/cudev/include/opencv2/cudev/ptr2d/glob.hpp @@ -54,12 +54,19 @@ namespace cv { namespace cudev { //! @addtogroup cudev //! @{ +/** @brief Structure similar to cv::cudev::GlobPtrSz but containing only a pointer and row step. + +Width and height fields are excluded due to performance reasons. The structure is intended +for internal use or for users who write device code. + */ template struct GlobPtr { typedef T value_type; typedef int index_type; T* data; + + //! stride between two consecutive rows in bytes. Step is stored always and everywhere in bytes!!! size_t step; __device__ __forceinline__ T* row(int y) { return ( T*)( ( uchar*)data + y * step); } @@ -69,6 +76,12 @@ template struct GlobPtr __device__ __forceinline__ const T& operator ()(int y, int x) const { return row(y)[x]; } }; +/** @brief Lightweight class encapsulating pitched memory on a GPU and passed to nvcc-compiled code (CUDA +kernels). + +Typically, it is used internally by OpenCV and by users who write device code. You can call +its members from both host and device code. + */ template struct GlobPtrSz : GlobPtr { int rows, cols;