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

commit x64 compilation from lulu cause the other branch dont seems to compile properly at home

Browse Source
This commit is contained in:
lucas doremus
2026-07-17 16:08:20 +02:00
parent c0f3eeb00d
commit 0efa4ee6f7
625 changed files with 117283 additions and 4426 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