AlcideSystem/Webcam
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

init: 64 bits version of the webcam

Browse Source
This commit is contained in:
Nicolas RACOT
2026-07-13 16:31:52 +02:00
parent c0f3eeb00d
commit 07e526544d
381 changed files with 43996 additions and 9097 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
include/opencv2/core/cuda/block.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_DEVICE_BLOCK_HPP__
#define __OPENCV_CUDA_DEVICE_BLOCK_HPP__
#ifndef OPENCV_CUDA_DEVICE_BLOCK_HPP
#define OPENCV_CUDA_DEVICE_BLOCK_HPP
/** @file
* @deprecated Use @ref cudev instead.
@ -106,7 +106,7 @@ namespace cv { namespace cuda { namespace device
}
template<typename InIt, typename OutIt, class UnOp>
static __device__ __forceinline__ void transfrom(InIt beg, InIt end, OutIt out, UnOp op)
static __device__ __forceinline__ void transform(InIt beg, InIt end, OutIt out, UnOp op)
{
int STRIDE = stride();
InIt t = beg + flattenedThreadId();
@ -117,7 +117,7 @@ namespace cv { namespace cuda { namespace device
}
template<typename InIt1, typename InIt2, typename OutIt, class BinOp>
static __device__ __forceinline__ void transfrom(InIt1 beg1, InIt1 end1, InIt2 beg2, OutIt out, BinOp op)
static __device__ __forceinline__ void transform(InIt1 beg1, InIt1 end1, InIt2 beg2, OutIt out, BinOp op)
{
int STRIDE = stride();
InIt1 t1 = beg1 + flattenedThreadId();
@ -208,4 +208,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif /* __OPENCV_CUDA_DEVICE_BLOCK_HPP__ */
#endif /* OPENCV_CUDA_DEVICE_BLOCK_HPP */

14
include/opencv2/core/cuda/border_interpolate.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__
#define __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__
#ifndef OPENCV_CUDA_BORDER_INTERPOLATE_HPP
#define OPENCV_CUDA_BORDER_INTERPOLATE_HPP
#include "saturate_cast.hpp"
#include "vec_traits.hpp"
@ -632,12 +632,12 @@ namespace cv { namespace cuda { namespace device
__device__ __forceinline__ int idx_row_low(int y) const
{
return (y >= 0) * y + (y < 0) * (y - ((y - height + 1) / height) * height);
return (y >= 0) ? y : (y - ((y - height + 1) / height) * height);
}
__device__ __forceinline__ int idx_row_high(int y) const
{
return (y < height) * y + (y >= height) * (y % height);
return (y < height) ? y : (y % height);
}
__device__ __forceinline__ int idx_row(int y) const
@ -647,12 +647,12 @@ namespace cv { namespace cuda { namespace device
__device__ __forceinline__ int idx_col_low(int x) const
{
return (x >= 0) * x + (x < 0) * (x - ((x - width + 1) / width) * width);
return (x >= 0) ? x : (x - ((x - width + 1) / width) * width);
}
__device__ __forceinline__ int idx_col_high(int x) const
{
return (x < width) * x + (x >= width) * (x % width);
return (x < width) ? x : (x % width);
}
__device__ __forceinline__ int idx_col(int x) const
@ -719,4 +719,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__
#endif // OPENCV_CUDA_BORDER_INTERPOLATE_HPP

6
include/opencv2/core/cuda/color.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_COLOR_HPP__
#define __OPENCV_CUDA_COLOR_HPP__
#ifndef OPENCV_CUDA_COLOR_HPP
#define OPENCV_CUDA_COLOR_HPP
#include "detail/color_detail.hpp"
@ -306,4 +306,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__
#endif // OPENCV_CUDA_COLOR_HPP

6
include/opencv2/core/cuda/common.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_COMMON_HPP__
#define __OPENCV_CUDA_COMMON_HPP__
#ifndef OPENCV_CUDA_COMMON_HPP
#define OPENCV_CUDA_COMMON_HPP
#include <cuda_runtime.h>
#include "opencv2/core/cuda_types.hpp"
@ -106,4 +106,4 @@ namespace cv { namespace cuda
//! @endcond
#endif // __OPENCV_CUDA_COMMON_HPP__
#endif // OPENCV_CUDA_COMMON_HPP

6
include/opencv2/core/cuda/datamov_utils.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_DATAMOV_UTILS_HPP__
#define __OPENCV_CUDA_DATAMOV_UTILS_HPP__
#ifndef OPENCV_CUDA_DATAMOV_UTILS_HPP
#define OPENCV_CUDA_DATAMOV_UTILS_HPP
#include "common.hpp"
@ -110,4 +110,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_DATAMOV_UTILS_HPP__
#endif // OPENCV_CUDA_DATAMOV_UTILS_HPP

6
include/opencv2/core/cuda/detail/color_detail.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_COLOR_DETAIL_HPP__
#define __OPENCV_CUDA_COLOR_DETAIL_HPP__
#ifndef OPENCV_CUDA_COLOR_DETAIL_HPP
#define OPENCV_CUDA_COLOR_DETAIL_HPP
#include "../common.hpp"
#include "../vec_traits.hpp"
@ -1977,4 +1977,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_COLOR_DETAIL_HPP__
#endif // OPENCV_CUDA_COLOR_DETAIL_HPP

16
include/opencv2/core/cuda/detail/reduce.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_REDUCE_DETAIL_HPP__
#define __OPENCV_CUDA_REDUCE_DETAIL_HPP__
#ifndef OPENCV_CUDA_REDUCE_DETAIL_HPP
#define OPENCV_CUDA_REDUCE_DETAIL_HPP
#include <thrust/tuple.h>
#include "../warp.hpp"
@ -275,9 +275,9 @@ namespace cv { namespace cuda { namespace device
template <typename Pointer, typename Reference, class Op>
static __device__ void reduce(Pointer smem, Reference val, unsigned int tid, Op op)
{
#if __CUDA_ARCH__ >= 300
(void) smem;
(void) tid;
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
CV_UNUSED(smem);
CV_UNUSED(tid);
Unroll<N / 2, Pointer, Reference, Op>::loopShfl(val, op, N);
#else
@ -298,7 +298,7 @@ namespace cv { namespace cuda { namespace device
{
const unsigned int laneId = Warp::laneId();
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
Unroll<16, Pointer, Reference, Op>::loopShfl(val, op, warpSize);
if (laneId == 0)
@ -321,7 +321,7 @@ namespace cv { namespace cuda { namespace device
if (tid < 32)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
Unroll<M / 2, Pointer, Reference, Op>::loopShfl(val, op, M);
#else
Unroll<M / 2, Pointer, Reference, Op>::loop(smem, val, tid, op);
@ -362,4 +362,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_REDUCE_DETAIL_HPP__
#endif // OPENCV_CUDA_REDUCE_DETAIL_HPP

12
include/opencv2/core/cuda/detail/reduce_key_val.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP__
#define __OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP__
#ifndef OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP
#define OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP
#include <thrust/tuple.h>
#include "../warp.hpp"
@ -402,9 +402,9 @@ namespace cv { namespace cuda { namespace device
static __device__ void reduce(KP skeys, KR key, VP svals, VR val, unsigned int tid, Cmp cmp)
{
#if 0 // __CUDA_ARCH__ >= 300
(void) skeys;
(void) svals;
(void) tid;
CV_UNUSED(skeys);
CV_UNUSED(svals);
CV_UNUSED(tid);
Unroll<N / 2, KP, KR, VP, VR, Cmp>::loopShfl(key, val, cmp, N);
#else
@ -499,4 +499,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP__
#endif // OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP

17
include/opencv2/core/cuda/detail/transform_detail.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_TRANSFORM_DETAIL_HPP__
#define __OPENCV_CUDA_TRANSFORM_DETAIL_HPP__
#ifndef OPENCV_CUDA_TRANSFORM_DETAIL_HPP
#define OPENCV_CUDA_TRANSFORM_DETAIL_HPP
#include "../common.hpp"
#include "../vec_traits.hpp"
@ -223,11 +223,7 @@ namespace cv { namespace cuda { namespace device
if (x_shifted + ft::smart_shift - 1 < src_.cols)
{
const read_type src_n_el = ((const read_type*)src)[x];
write_type dst_n_el = ((const write_type*)dst)[x];
OpUnroller<ft::smart_shift>::unroll(src_n_el, dst_n_el, mask, op, x_shifted, y);
((write_type*)dst)[x] = dst_n_el;
OpUnroller<ft::smart_shift>::unroll(src_n_el, ((write_type*)dst)[x], mask, op, x_shifted, y);
}
else
{
@ -275,11 +271,8 @@ namespace cv { namespace cuda { namespace device
{
const read_type1 src1_n_el = ((const read_type1*)src1)[x];
const read_type2 src2_n_el = ((const read_type2*)src2)[x];
write_type dst_n_el = ((const write_type*)dst)[x];
OpUnroller<ft::smart_shift>::unroll(src1_n_el, src2_n_el, dst_n_el, mask, op, x_shifted, y);
((write_type*)dst)[x] = dst_n_el;
OpUnroller<ft::smart_shift>::unroll(src1_n_el, src2_n_el, ((write_type*)dst)[x], mask, op, x_shifted, y);
}
else
{
@ -396,4 +389,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_TRANSFORM_DETAIL_HPP__
#endif // OPENCV_CUDA_TRANSFORM_DETAIL_HPP

6
include/opencv2/core/cuda/detail/type_traits_detail.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP__
#define __OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP__
#ifndef OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP
#define OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP
#include "../common.hpp"
#include "../vec_traits.hpp"
@ -188,4 +188,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP__
#endif // OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP

6
include/opencv2/core/cuda/detail/vec_distance_detail.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__
#define __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__
#ifndef OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP
#define OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP
#include "../datamov_utils.hpp"
@ -118,4 +118,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__
#endif // OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP

6
include/opencv2/core/cuda/dynamic_smem.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_DYNAMIC_SMEM_HPP__
#define __OPENCV_CUDA_DYNAMIC_SMEM_HPP__
#ifndef OPENCV_CUDA_DYNAMIC_SMEM_HPP
#define OPENCV_CUDA_DYNAMIC_SMEM_HPP
/** @file
* @deprecated Use @ref cudev instead.
@ -85,4 +85,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_DYNAMIC_SMEM_HPP__
#endif // OPENCV_CUDA_DYNAMIC_SMEM_HPP

20
include/opencv2/core/cuda/emulation.hpp
View File

@ -177,8 +177,8 @@ namespace cv { namespace cuda { namespace device
} while (assumed != old);
return __longlong_as_double(old);
#else
(void) address;
(void) val;
CV_UNUSED(address);
CV_UNUSED(val);
return 0.0;
#endif
}
@ -199,8 +199,8 @@ namespace cv { namespace cuda { namespace device
} while (assumed != old);
return __int_as_float(old);
#else
(void) address;
(void) val;
CV_UNUSED(address);
CV_UNUSED(val);
return 0.0f;
#endif
}
@ -216,8 +216,8 @@ namespace cv { namespace cuda { namespace device
} while (assumed != old);
return __longlong_as_double(old);
#else
(void) address;
(void) val;
CV_UNUSED(address);
CV_UNUSED(val);
return 0.0;
#endif
}
@ -238,8 +238,8 @@ namespace cv { namespace cuda { namespace device
} while (assumed != old);
return __int_as_float(old);
#else
(void) address;
(void) val;
CV_UNUSED(address);
CV_UNUSED(val);
return 0.0f;
#endif
}
@ -255,8 +255,8 @@ namespace cv { namespace cuda { namespace device
} while (assumed != old);
return __longlong_as_double(old);
#else
(void) address;
(void) val;
CV_UNUSED(address);
CV_UNUSED(val);
return 0.0;
#endif
}

18
include/opencv2/core/cuda/filters.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_FILTERS_HPP__
#define __OPENCV_CUDA_FILTERS_HPP__
#ifndef OPENCV_CUDA_FILTERS_HPP
#define OPENCV_CUDA_FILTERS_HPP
#include "saturate_cast.hpp"
#include "vec_traits.hpp"
@ -64,8 +64,8 @@ namespace cv { namespace cuda { namespace device
explicit __host__ __device__ __forceinline__ PointFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
: src(src_)
{
(void)fx;
(void)fy;
CV_UNUSED(fx);
CV_UNUSED(fy);
}
__device__ __forceinline__ elem_type operator ()(float y, float x) const
@ -84,8 +84,8 @@ namespace cv { namespace cuda { namespace device
explicit __host__ __device__ __forceinline__ LinearFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
: src(src_)
{
(void)fx;
(void)fy;
CV_UNUSED(fx);
CV_UNUSED(fy);
}
__device__ __forceinline__ elem_type operator ()(float y, float x) const
{
@ -125,8 +125,8 @@ namespace cv { namespace cuda { namespace device
explicit __host__ __device__ __forceinline__ CubicFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
: src(src_)
{
(void)fx;
(void)fy;
CV_UNUSED(fx);
CV_UNUSED(fy);
}
static __device__ __forceinline__ float bicubicCoeff(float x_)
@ -283,4 +283,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_FILTERS_HPP__
#endif // OPENCV_CUDA_FILTERS_HPP

6
include/opencv2/core/cuda/funcattrib.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP_
#define __OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP_
#ifndef OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP
#define OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP
#include <cstdio>
@ -76,4 +76,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif /* __OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP_ */
#endif /* OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP */

26
include/opencv2/core/cuda/functional.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_FUNCTIONAL_HPP__
#define __OPENCV_CUDA_FUNCTIONAL_HPP__
#ifndef OPENCV_CUDA_FUNCTIONAL_HPP
#define OPENCV_CUDA_FUNCTIONAL_HPP
#include <functional>
#include "saturate_cast.hpp"
@ -58,8 +58,22 @@
namespace cv { namespace cuda { namespace device
{
// Function Objects
#ifdef CV_CXX11
template<typename Argument, typename Result> struct unary_function
{
typedef Argument argument_type;
typedef Result result_type;
};
template<typename Argument1, typename Argument2, typename Result> struct binary_function
{
typedef Argument1 first_argument_type;
typedef Argument2 second_argument_type;
typedef Result result_type;
};
#else
template<typename Argument, typename Result> struct unary_function : public std::unary_function<Argument, Result> {};
template<typename Argument1, typename Argument2, typename Result> struct binary_function : public std::binary_function<Argument1, Argument2, Result> {};
#endif
// Arithmetic Operations
template <typename T> struct plus : binary_function<T, T, T>
@ -583,7 +597,7 @@ namespace cv { namespace cuda { namespace device
template <typename T> struct thresh_trunc_func : unary_function<T, T>
{
explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}
explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {CV_UNUSED(maxVal_);}
__device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const
{
@ -599,7 +613,7 @@ namespace cv { namespace cuda { namespace device
template <typename T> struct thresh_to_zero_func : unary_function<T, T>
{
explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}
explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {CV_UNUSED(maxVal_);}
__device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const
{
@ -615,7 +629,7 @@ namespace cv { namespace cuda { namespace device
template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T>
{
explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}
explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {CV_UNUSED(maxVal_);}
__device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const
{
@ -794,4 +808,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_FUNCTIONAL_HPP__
#endif // OPENCV_CUDA_FUNCTIONAL_HPP

6
include/opencv2/core/cuda/limits.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_LIMITS_HPP__
#define __OPENCV_CUDA_LIMITS_HPP__
#ifndef OPENCV_CUDA_LIMITS_HPP
#define OPENCV_CUDA_LIMITS_HPP
#include <limits.h>
#include <float.h>
@ -125,4 +125,4 @@ template <> struct numeric_limits<double>
//! @endcond
#endif // __OPENCV_CUDA_LIMITS_HPP__
#endif // OPENCV_CUDA_LIMITS_HPP

10
include/opencv2/core/cuda/reduce.hpp
View File

@ -40,8 +40,12 @@
//
//M*/
#ifndef __OPENCV_CUDA_REDUCE_HPP__
#define __OPENCV_CUDA_REDUCE_HPP__
#ifndef OPENCV_CUDA_REDUCE_HPP
#define OPENCV_CUDA_REDUCE_HPP
#ifndef THRUST_DEBUG // eliminate -Wundef warning
#define THRUST_DEBUG 0
#endif
#include <thrust/tuple.h>
#include "detail/reduce.hpp"
@ -202,4 +206,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_UTILITY_HPP__
#endif // OPENCV_CUDA_REDUCE_HPP

14
include/opencv2/core/cuda/saturate_cast.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_SATURATE_CAST_HPP__
#define __OPENCV_CUDA_SATURATE_CAST_HPP__
#ifndef OPENCV_CUDA_SATURATE_CAST_HPP
#define OPENCV_CUDA_SATURATE_CAST_HPP
#include "common.hpp"
@ -101,7 +101,7 @@ namespace cv { namespace cuda { namespace device
}
template<> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
uint res = 0;
asm("cvt.rni.sat.u8.f64 %0, %1;" : "=r"(res) : "d"(v));
return res;
@ -149,7 +149,7 @@ namespace cv { namespace cuda { namespace device
}
template<> __device__ __forceinline__ schar saturate_cast<schar>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
uint res = 0;
asm("cvt.rni.sat.s8.f64 %0, %1;" : "=r"(res) : "d"(v));
return res;
@ -191,7 +191,7 @@ namespace cv { namespace cuda { namespace device
}
template<> __device__ __forceinline__ ushort saturate_cast<ushort>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
ushort res = 0;
asm("cvt.rni.sat.u16.f64 %0, %1;" : "=h"(res) : "d"(v));
return res;
@ -226,7 +226,7 @@ namespace cv { namespace cuda { namespace device
}
template<> __device__ __forceinline__ short saturate_cast<short>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
short res = 0;
asm("cvt.rni.sat.s16.f64 %0, %1;" : "=h"(res) : "d"(v));
return res;
@ -289,4 +289,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif /* __OPENCV_CUDA_SATURATE_CAST_HPP__ */
#endif /* OPENCV_CUDA_SATURATE_CAST_HPP */

12
include/opencv2/core/cuda/scan.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_SCAN_HPP__
#define __OPENCV_CUDA_SCAN_HPP__
#ifndef OPENCV_CUDA_SCAN_HPP
#define OPENCV_CUDA_SCAN_HPP
#include "opencv2/core/cuda/common.hpp"
#include "opencv2/core/cuda/utility.hpp"
@ -61,7 +61,7 @@ namespace cv { namespace cuda { namespace device
template <ScanKind Kind, typename T, typename F> struct WarpScan
{
__device__ __forceinline__ WarpScan() {}
__device__ __forceinline__ WarpScan(const WarpScan& other) { (void)other; }
__device__ __forceinline__ WarpScan(const WarpScan& other) { CV_UNUSED(other); }
__device__ __forceinline__ T operator()( volatile T *ptr , const unsigned int idx)
{
@ -95,7 +95,7 @@ namespace cv { namespace cuda { namespace device
template <ScanKind Kind , typename T, typename F> struct WarpScanNoComp
{
__device__ __forceinline__ WarpScanNoComp() {}
__device__ __forceinline__ WarpScanNoComp(const WarpScanNoComp& other) { (void)other; }
__device__ __forceinline__ WarpScanNoComp(const WarpScanNoComp& other) { CV_UNUSED(other); }
__device__ __forceinline__ T operator()( volatile T *ptr , const unsigned int idx)
{
@ -135,7 +135,7 @@ namespace cv { namespace cuda { namespace device
template <ScanKind Kind , typename T, typename Sc, typename F> struct BlockScan
{
__device__ __forceinline__ BlockScan() {}
__device__ __forceinline__ BlockScan(const BlockScan& other) { (void)other; }
__device__ __forceinline__ BlockScan(const BlockScan& other) { CV_UNUSED(other); }
__device__ __forceinline__ T operator()(volatile T *ptr)
{
@ -255,4 +255,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_SCAN_HPP__
#endif // OPENCV_CUDA_SCAN_HPP

6
include/opencv2/core/cuda/simd_functions.hpp
View File

@ -70,8 +70,8 @@
* POSSIBILITY OF SUCH DAMAGE.
*/
#ifndef __OPENCV_CUDA_SIMD_FUNCTIONS_HPP__
#define __OPENCV_CUDA_SIMD_FUNCTIONS_HPP__
#ifndef OPENCV_CUDA_SIMD_FUNCTIONS_HPP
#define OPENCV_CUDA_SIMD_FUNCTIONS_HPP
#include "common.hpp"
@ -866,4 +866,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_SIMD_FUNCTIONS_HPP__
#endif // OPENCV_CUDA_SIMD_FUNCTIONS_HPP

6
include/opencv2/core/cuda/transform.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_TRANSFORM_HPP__
#define __OPENCV_CUDA_TRANSFORM_HPP__
#ifndef OPENCV_CUDA_TRANSFORM_HPP
#define OPENCV_CUDA_TRANSFORM_HPP
#include "common.hpp"
#include "utility.hpp"
@ -72,4 +72,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_TRANSFORM_HPP__
#endif // OPENCV_CUDA_TRANSFORM_HPP

6
include/opencv2/core/cuda/type_traits.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_TYPE_TRAITS_HPP__
#define __OPENCV_CUDA_TYPE_TRAITS_HPP__
#ifndef OPENCV_CUDA_TYPE_TRAITS_HPP
#define OPENCV_CUDA_TYPE_TRAITS_HPP
#include "detail/type_traits_detail.hpp"
@ -87,4 +87,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_TYPE_TRAITS_HPP__
#endif // OPENCV_CUDA_TYPE_TRAITS_HPP

15
include/opencv2/core/cuda/utility.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_UTILITY_HPP__
#define __OPENCV_CUDA_UTILITY_HPP__
#ifndef OPENCV_CUDA_UTILITY_HPP
#define OPENCV_CUDA_UTILITY_HPP
#include "saturate_cast.hpp"
#include "datamov_utils.hpp"
@ -54,6 +54,15 @@
namespace cv { namespace cuda { namespace device
{
struct CV_EXPORTS ThrustAllocator
{
typedef uchar value_type;
virtual ~ThrustAllocator();
virtual __device__ __host__ uchar* allocate(size_t numBytes) = 0;
virtual __device__ __host__ void deallocate(uchar* ptr, size_t numBytes) = 0;
static ThrustAllocator& getAllocator();
static void setAllocator(ThrustAllocator* allocator);
};
#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
@ -218,4 +227,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_UTILITY_HPP__
#endif // OPENCV_CUDA_UTILITY_HPP

6
include/opencv2/core/cuda/vec_distance.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_VEC_DISTANCE_HPP__
#define __OPENCV_CUDA_VEC_DISTANCE_HPP__
#ifndef OPENCV_CUDA_VEC_DISTANCE_HPP
#define OPENCV_CUDA_VEC_DISTANCE_HPP
#include "reduce.hpp"
#include "functional.hpp"
@ -229,4 +229,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_VEC_DISTANCE_HPP__
#endif // OPENCV_CUDA_VEC_DISTANCE_HPP

6
include/opencv2/core/cuda/vec_math.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_VECMATH_HPP__
#define __OPENCV_CUDA_VECMATH_HPP__
#ifndef OPENCV_CUDA_VECMATH_HPP
#define OPENCV_CUDA_VECMATH_HPP
#include "vec_traits.hpp"
#include "saturate_cast.hpp"
@ -927,4 +927,4 @@ CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double)
//! @endcond
#endif // __OPENCV_CUDA_VECMATH_HPP__
#endif // OPENCV_CUDA_VECMATH_HPP

6
include/opencv2/core/cuda/vec_traits.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_VEC_TRAITS_HPP__
#define __OPENCV_CUDA_VEC_TRAITS_HPP__
#ifndef OPENCV_CUDA_VEC_TRAITS_HPP
#define OPENCV_CUDA_VEC_TRAITS_HPP
#include "common.hpp"
@ -285,4 +285,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif // __OPENCV_CUDA_VEC_TRAITS_HPP__
#endif // OPENCV_CUDA_VEC_TRAITS_HPP

8
include/opencv2/core/cuda/warp.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_DEVICE_WARP_HPP__
#define __OPENCV_CUDA_DEVICE_WARP_HPP__
#ifndef OPENCV_CUDA_DEVICE_WARP_HPP
#define OPENCV_CUDA_DEVICE_WARP_HPP
/** @file
* @deprecated Use @ref cudev instead.
@ -64,7 +64,7 @@ namespace cv { namespace cuda { namespace device
static __device__ __forceinline__ unsigned int laneId()
{
unsigned int ret;
asm("mov.u32 %0, %laneid;" : "=r"(ret) );
asm("mov.u32 %0, %%laneid;" : "=r"(ret) );
return ret;
}
@ -136,4 +136,4 @@ namespace cv { namespace cuda { namespace device
//! @endcond
#endif /* __OPENCV_CUDA_DEVICE_WARP_HPP__ */
#endif /* OPENCV_CUDA_DEVICE_WARP_HPP */

33
include/opencv2/core/cuda/warp_shuffle.hpp
View File

@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_CUDA_WARP_SHUFFLE_HPP__
#define __OPENCV_CUDA_WARP_SHUFFLE_HPP__
#ifndef OPENCV_CUDA_WARP_SHUFFLE_HPP
#define OPENCV_CUDA_WARP_SHUFFLE_HPP
/** @file
* @deprecated Use @ref cudev instead.
@ -51,10 +51,15 @@
namespace cv { namespace cuda { namespace device
{
#if __CUDACC_VER_MAJOR__ >= 9
# define __shfl(x, y, z) __shfl_sync(0xFFFFFFFFU, x, y, z)
# define __shfl_up(x, y, z) __shfl_up_sync(0xFFFFFFFFU, x, y, z)
# define __shfl_down(x, y, z) __shfl_down_sync(0xFFFFFFFFU, x, y, z)
#endif
template <typename T>
__device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
return __shfl(val, srcLane, width);
#else
return T();
@ -62,7 +67,7 @@ namespace cv { namespace cuda { namespace device
}
__device__ __forceinline__ unsigned int shfl(unsigned int val, int srcLane, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
return (unsigned int) __shfl((int) val, srcLane, width);
#else
return 0;
@ -70,7 +75,7 @@ namespace cv { namespace cuda { namespace device
}
__device__ __forceinline__ double shfl(double val, int srcLane, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
int lo = __double2loint(val);
int hi = __double2hiint(val);
@ -86,7 +91,7 @@ namespace cv { namespace cuda { namespace device
template <typename T>
__device__ __forceinline__ T shfl_down(T val, unsigned int delta, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
return __shfl_down(val, delta, width);
#else
return T();
@ -94,7 +99,7 @@ namespace cv { namespace cuda { namespace device
}
__device__ __forceinline__ unsigned int shfl_down(unsigned int val, unsigned int delta, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
return (unsigned int) __shfl_down((int) val, delta, width);
#else
return 0;
@ -102,7 +107,7 @@ namespace cv { namespace cuda { namespace device
}
__device__ __forceinline__ double shfl_down(double val, unsigned int delta, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
int lo = __double2loint(val);
int hi = __double2hiint(val);
@ -118,7 +123,7 @@ namespace cv { namespace cuda { namespace device
template <typename T>
__device__ __forceinline__ T shfl_up(T val, unsigned int delta, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
return __shfl_up(val, delta, width);
#else
return T();
@ -126,7 +131,7 @@ namespace cv { namespace cuda { namespace device
}
__device__ __forceinline__ unsigned int shfl_up(unsigned int val, unsigned int delta, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
return (unsigned int) __shfl_up((int) val, delta, width);
#else
return 0;
@ -134,7 +139,7 @@ namespace cv { namespace cuda { namespace device
}
__device__ __forceinline__ double shfl_up(double val, unsigned int delta, int width = warpSize)
{
#if __CUDA_ARCH__ >= 300
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
int lo = __double2loint(val);
int hi = __double2hiint(val);
@ -148,6 +153,10 @@ namespace cv { namespace cuda { namespace device
}
}}}
# undef __shfl
# undef __shfl_up
# undef __shfl_down
//! @endcond
#endif // __OPENCV_CUDA_WARP_SHUFFLE_HPP__
#endif // OPENCV_CUDA_WARP_SHUFFLE_HPP