mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-01-18 06:03:15 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

fixed hundreds of "anonymous" warnings for gpu module.

Browse Source
This commit is contained in:
Marina Kolpakova 2012-06-08 17:09:38 +00:00
parent ffa44fb114
commit e86f0aaea1
13 changed files with 220 additions and 49 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
modules/core/CMakeLists.txt
View File

@ -5,8 +5,11 @@ ocv_module_include_directories(${ZLIB_INCLUDE_DIR})
if(HAVE_CUDA)
file(GLOB lib_cuda "src/cuda/*.cu")
source_group("Cuda" FILES "${lib_cuda}")
ocv_include_directories(${CUDA_INCLUDE_DIRS} "${OpenCV_SOURCE_DIR}/modules/gpu/src" "${OpenCV_SOURCE_DIR}/modules/gpu/src/cuda")
include_directories(AFTER SYSTEM ${CUDA_INCLUDE_DIRS})
ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpu/src" "${OpenCV_SOURCE_DIR}/modules/gpu/src/cuda")
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
OCV_CUDA_COMPILE(cuda_objs ${lib_cuda})
set(cuda_link_libs ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})

4
modules/gpu/CMakeLists.txt
View File

@ -30,7 +30,9 @@ if (HAVE_CUDA)
set(ncv_files ${ncv_srcs} ${ncv_hdrs} ${ncv_cuda})
source_group("Src\\NVidia" FILES ${ncv_files})
ocv_include_directories("src/nvidia" "src/nvidia/core" "src/nvidia/NPP_staging" ${CUDA_INCLUDE_DIRS})
include_directories(AFTER SYSTEM ${CUDA_INCLUDE_DIRS})
ocv_include_directories("src/nvidia" "src/nvidia/core" "src/nvidia/NPP_staging")
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-keep")
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-Xcompiler;/EHsc-;")

16
modules/gpu/src/cuda/matrix_reductions.cu
View File

@ -87,7 +87,9 @@ namespace cv { namespace gpu { namespace device
__device__ __forceinline__ bool operator()(int y, int x) const
{
return true;
}
}
__device__ __forceinline__ MaskTrue(){}
__device__ __forceinline__ MaskTrue(const MaskTrue& mask_){}
};
//////////////////////////////////////////////////////////////////////////////
@ -1795,6 +1797,9 @@ namespace cv { namespace gpu { namespace device
return 0;
}
__device__ __forceinline__ SumReductor(const SumReductor& other){}
__device__ __forceinline__ SumReductor(){}
__device__ __forceinline__ S operator ()(volatile S a, volatile S b) const
{
return a + b;
@ -1813,6 +1818,9 @@ namespace cv { namespace gpu { namespace device
return 0;
}
__device__ __forceinline__ AvgReductor(const AvgReductor& other){}
__device__ __forceinline__ AvgReductor(){}
__device__ __forceinline__ S operator ()(volatile S a, volatile S b) const
{
return a + b;
@ -1831,6 +1839,9 @@ namespace cv { namespace gpu { namespace device
return numeric_limits<S>::max();
}
__device__ __forceinline__ MinReductor(const MinReductor& other){}
__device__ __forceinline__ MinReductor(){}
template <typename T> __device__ __forceinline__ T operator ()(volatile T a, volatile T b) const
{
return saturate_cast<T>(::min(a, b));
@ -1853,6 +1864,9 @@ namespace cv { namespace gpu { namespace device
return numeric_limits<S>::min();
}
__device__ __forceinline__ MaxReductor(const MaxReductor& other){}
__device__ __forceinline__ MaxReductor(){}
template <typename T> __device__ __forceinline__ int operator ()(volatile T a, volatile T b) const
{
return ::max(a, b);

6
modules/gpu/src/cuda/surf.cu
View File

@ -116,7 +116,7 @@ namespace cv { namespace gpu { namespace device
template <int N> __device__ float icvCalcHaarPatternSum(const float src[][5], int oldSize, int newSize, int y, int x)
{
#if __CUDA_ARCH__ >= 200
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
typedef double real_t;
#else
typedef float real_t;
@ -248,7 +248,7 @@ namespace cv { namespace gpu { namespace device
template <typename Mask>
__global__ void icvFindMaximaInLayer(const PtrStepf det, const PtrStepf trace, int4* maxPosBuffer, unsigned int* maxCounter)
{
#if __CUDA_ARCH__ >= 110
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110
extern __shared__ float N9[];
@ -371,7 +371,7 @@ namespace cv { namespace gpu { namespace device
float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian,
unsigned int* featureCounter)
{
#if __CUDA_ARCH__ >= 110
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110
const int4 maxPos = maxPosBuffer[blockIdx.x];

4
modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
View File

@ -231,7 +231,7 @@ __device__ Ncv32u d_outMaskPosition;
__device__ void compactBlockWriteOutAnchorParallel(Ncv32u threadPassFlag, Ncv32u threadElem, Ncv32u *vectorOut)
{
#if __CUDA_ARCH__ >= 110
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110
__shared__ Ncv32u shmem[NUM_THREADS_ANCHORSPARALLEL * 2];
__shared__ Ncv32u numPassed;
@ -587,7 +587,7 @@ __global__ void applyHaarClassifierClassifierParallel(Ncv32u *d_IImg, Ncv32u IIm
}
else
{
#if __CUDA_ARCH__ >= 110
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110
if (bPass && !threadIdx.x)
{
Ncv32u outMaskOffset = atomicAdd(&d_outMaskPosition, 1);

2
modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp
View File

@ -41,7 +41,7 @@
#ifndef _ncvruntimetemplates_hpp_
#define _ncvruntimetemplates_hpp_
#if _MSC_VER >= 1200
#if defined _MSC_VER &&_MSC_VER >= 1200
#pragma warning( disable: 4800 )
#endif

2
modules/gpu/src/opencv2/gpu/device/datamov_utils.hpp
View File

@ -47,7 +47,7 @@
namespace cv { namespace gpu { namespace device
{
#if __CUDA_ARCH__ >= 200
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
// for Fermi memory space is detected automatically
template <typename T> struct ForceGlob

37
modules/gpu/src/opencv2/gpu/device/detail/color_detail.hpp
View File

@ -114,6 +114,11 @@ namespace cv { namespace gpu { namespace device
return dst;
}
__device__ __forceinline__ RGB2RGB()
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
__device__ __forceinline__ RGB2RGB(const RGB2RGB& other_)
:unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
};
template <> struct RGB2RGB<uchar, 4, 4, 2> : unary_function<uint, uint>
@ -129,6 +134,9 @@ namespace cv { namespace gpu { namespace device
return dst;
}
__device__ __forceinline__ RGB2RGB():unary_function<uint, uint>(){}
__device__ __forceinline__ RGB2RGB(const RGB2RGB& other_):unary_function<uint, uint>(){}
};
}
@ -184,6 +192,8 @@ namespace cv { namespace gpu { namespace device
{
return RGB2RGB5x5Converter<green_bits, bidx>::cvt(src);
}
__device__ __forceinline__ RGB2RGB5x5():unary_function<uchar3, ushort>(){}
__device__ __forceinline__ RGB2RGB5x5(const RGB2RGB5x5& other_):unary_function<uchar3, ushort>(){}
};
template<int bidx, int green_bits> struct RGB2RGB5x5<4, bidx,green_bits> : unary_function<uint, ushort>
{
@ -191,6 +201,9 @@ namespace cv { namespace gpu { namespace device
{
return RGB2RGB5x5Converter<green_bits, bidx>::cvt(src);
}
__device__ __forceinline__ RGB2RGB5x5():unary_function<uint, ushort>(){}
__device__ __forceinline__ RGB2RGB5x5(const RGB2RGB5x5& other_):unary_function<uint, ushort>(){}
};
}
@ -252,7 +265,11 @@ namespace cv { namespace gpu { namespace device
RGB5x52RGBConverter<green_bits, bidx>::cvt(src, dst);
return dst;
}
__device__ __forceinline__ RGB5x52RGB():unary_function<ushort, uchar3>(){}
__device__ __forceinline__ RGB5x52RGB(const RGB5x52RGB& other_):unary_function<ushort, uchar3>(){}
};
template <int bidx, int green_bits> struct RGB5x52RGB<4, bidx, green_bits> : unary_function<ushort, uint>
{
__device__ __forceinline__ uint operator()(ushort src) const
@ -261,6 +278,8 @@ namespace cv { namespace gpu { namespace device
RGB5x52RGBConverter<green_bits, bidx>::cvt(src, dst);
return dst;
}
__device__ __forceinline__ RGB5x52RGB():unary_function<ushort, uint>(){}
__device__ __forceinline__ RGB5x52RGB(const RGB5x52RGB& other_):unary_function<ushort, uint>(){}
};
}
@ -289,6 +308,8 @@ namespace cv { namespace gpu { namespace device
return dst;
}
__device__ __forceinline__ Gray2RGB():unary_function<T, typename TypeVec<T, dcn>::vec_type>(){}
__device__ __forceinline__ Gray2RGB(const Gray2RGB& other_):unary_function<T, typename TypeVec<T, dcn>::vec_type>(){}
};
template <> struct Gray2RGB<uchar, 4> : unary_function<uchar, uint>
{
@ -302,6 +323,8 @@ namespace cv { namespace gpu { namespace device
return dst;
}
__device__ __forceinline__ Gray2RGB():unary_function<uchar, uint>(){}
__device__ __forceinline__ Gray2RGB(const Gray2RGB& other_):unary_function<uchar, uint>(){}
};
}
@ -340,6 +363,8 @@ namespace cv { namespace gpu { namespace device
{
return Gray2RGB5x5Converter<green_bits>::cvt(src);
}
__device__ __forceinline__ Gray2RGB5x5():unary_function<uchar, ushort>(){}
__device__ __forceinline__ Gray2RGB5x5(const Gray2RGB5x5& other_):unary_function<uchar, ushort>(){}
};
}
@ -471,6 +496,8 @@ namespace cv { namespace gpu { namespace device
RGB2YUVConvert<bidx>(&src.x, dst);
return dst;
}
__device__ __forceinline__ RGB2YUV():unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
__device__ __forceinline__ RGB2YUV(const RGB2YUV& other_):unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
};
}
@ -535,7 +562,10 @@ namespace cv { namespace gpu { namespace device
return dst;
}
__device__ __forceinline__ YUV2RGB():unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
__device__ __forceinline__ YUV2RGB(const YUV2RGB& other_):unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
};
template <int bidx> struct YUV2RGB<uchar, 4, 4, bidx> : unary_function<uint, uint>
{
__device__ __forceinline__ uint operator ()(uint src) const
@ -605,7 +635,10 @@ namespace cv { namespace gpu { namespace device
RGB2YCrCbConvert<bidx>(&src.x, dst);
return dst;
}
__device__ __forceinline__ RGB2YCrCb():unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
__device__ __forceinline__ RGB2YCrCb(const RGB2YCrCb& other_):unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
};
template <int bidx> struct RGB2YCrCb<uchar, 4, 4, bidx> : unary_function<uint, uint>
{
__device__ __forceinline__ uint operator ()(uint src) const
@ -676,7 +709,10 @@ namespace cv { namespace gpu { namespace device
return dst;
}
__device__ __forceinline__ YCrCb2RGB():unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
__device__ __forceinline__ YCrCb2RGB(const YCrCb2RGB& other_):unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
};
template <int bidx> struct YCrCb2RGB<uchar, 4, 4, bidx> : unary_function<uint, uint>
{
__device__ __forceinline__ uint operator ()(uint src) const
@ -1331,6 +1367,7 @@ namespace cv { namespace gpu { namespace device
{
return HLS2RGBConvert<bidx, hr>(src);
}
};
}

153
modules/gpu/src/opencv2/gpu/device/functional.hpp
View File

@ -56,158 +56,224 @@ namespace cv { namespace gpu { namespace device
using thrust::binary_function;
// Arithmetic Operations
template <typename T> struct plus : binary_function<T, T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a + b;
}
__device__ __forceinline__ plus(const plus& other):binary_function<T,T,T>(){}
__device__ __forceinline__ plus():binary_function<T,T,T>(){}
};
template <typename T> struct minus : binary_function<T, T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a - b;
}
__device__ __forceinline__ minus(const minus& other):binary_function<T,T,T>(){}
__device__ __forceinline__ minus():binary_function<T,T,T>(){}
};
template <typename T> struct multiplies : binary_function<T, T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a * b;
}
__device__ __forceinline__ multiplies(const multiplies& other):binary_function<T,T,T>(){}
__device__ __forceinline__ multiplies():binary_function<T,T,T>(){}
};
template <typename T> struct divides : binary_function<T, T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a / b;
}
__device__ __forceinline__ divides(const divides& other):binary_function<T,T,T>(){}
__device__ __forceinline__ divides():binary_function<T,T,T>(){}
};
template <typename T> struct modulus : binary_function<T, T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a % b;
}
__device__ __forceinline__ modulus(const modulus& other):binary_function<T,T,T>(){}
__device__ __forceinline__ modulus():binary_function<T,T,T>(){}
};
template <typename T> struct negate : unary_function<T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a) const
{
return -a;
}
__device__ __forceinline__ negate(const negate& other):unary_function<T,T>(){}
__device__ __forceinline__ negate():unary_function<T,T>(){}
};
// Comparison Operations
template <typename T> struct equal_to : binary_function<T, T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a == b;
}
__device__ __forceinline__ equal_to(const equal_to& other):binary_function<T,T,bool>(){}
__device__ __forceinline__ equal_to():binary_function<T,T,bool>(){}
};
template <typename T> struct not_equal_to : binary_function<T, T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a != b;
}
__device__ __forceinline__ not_equal_to(const not_equal_to& other):binary_function<T,T,bool>(){}
__device__ __forceinline__ not_equal_to():binary_function<T,T,bool>(){}
};
template <typename T> struct greater : binary_function<T, T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a > b;
}
__device__ __forceinline__ greater(const greater& other):binary_function<T,T,bool>(){}
__device__ __forceinline__ greater():binary_function<T,T,bool>(){}
};
template <typename T> struct less : binary_function<T, T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a < b;
}
__device__ __forceinline__ less(const less& other):binary_function<T,T,bool>(){}
__device__ __forceinline__ less():binary_function<T,T,bool>(){}
};
template <typename T> struct greater_equal : binary_function<T, T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a >= b;
}
__device__ __forceinline__ greater_equal(const greater_equal& other):binary_function<T,T,bool>(){}
__device__ __forceinline__ greater_equal():binary_function<T,T,bool>(){}
};
template <typename T> struct less_equal : binary_function<T, T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a <= b;
}
__device__ __forceinline__ less_equal(const less_equal& other):binary_function<T,T,bool>(){}
__device__ __forceinline__ less_equal():binary_function<T,T,bool>(){}
};
// Logical Operations
template <typename T> struct logical_and : binary_function<T, T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a && b;
}
__device__ __forceinline__ logical_and(const logical_and& other):binary_function<T,T,bool>(){}
__device__ __forceinline__ logical_and():binary_function<T,T,bool>(){}
};
template <typename T> struct logical_or : binary_function<T, T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a || b;
}
__device__ __forceinline__ logical_or(const logical_or& other):binary_function<T,T,bool>(){}
__device__ __forceinline__ logical_or():binary_function<T,T,bool>(){}
};
template <typename T> struct logical_not : unary_function<T, bool>
{
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a) const
{
return !a;
}
__device__ __forceinline__ logical_not(const logical_not& other):unary_function<T,bool>(){}
__device__ __forceinline__ logical_not():unary_function<T,bool>(){}
};
// Bitwise Operations
template <typename T> struct bit_and : binary_function<T, T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a & b;
}
__device__ __forceinline__ bit_and(const bit_and& other):binary_function<T,T,T>(){}
__device__ __forceinline__ bit_and():binary_function<T,T,T>(){}
};
template <typename T> struct bit_or : binary_function<T, T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a | b;
}
__device__ __forceinline__ bit_or(const bit_or& other):binary_function<T,T,T>(){}
__device__ __forceinline__ bit_or():binary_function<T,T,T>(){}
};
template <typename T> struct bit_xor : binary_function<T, T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
typename TypeTraits<T>::ParameterType b) const
{
return a ^ b;
}
__device__ __forceinline__ bit_xor(const bit_xor& other):binary_function<T,T,T>(){}
__device__ __forceinline__ bit_xor():binary_function<T,T,T>(){}
};
template <typename T> struct bit_not : unary_function<T, T>
{
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType v) const
{
return ~v;
}
__device__ __forceinline__ bit_not(const bit_not& other):unary_function<T,T>(){}
__device__ __forceinline__ bit_not():unary_function<T,T>(){}
};
// Generalized Identity Operations
template <typename T> struct identity : unary_function<T, T>
{
__device__ __forceinline__ typename TypeTraits<T>::ParameterType operator()(typename TypeTraits<T>::ParameterType x) const
{
return x;
}
__device__ __forceinline__ identity(const identity& other):unary_function<T,T>(){}
__device__ __forceinline__ identity():unary_function<T,T>(){}
};
template <typename T1, typename T2> struct project1st : binary_function<T1, T2, T1>
@ -216,13 +282,18 @@ namespace cv { namespace gpu { namespace device
{
return lhs;
}
__device__ __forceinline__ project1st(const project1st& other):binary_function<T1,T2,T1>(){}
__device__ __forceinline__ project1st():binary_function<T1,T2,T1>(){}
};
template <typename T1, typename T2> struct project2nd : binary_function<T1, T2, T2>
{
__device__ __forceinline__ typename TypeTraits<T2>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const
{
return rhs;
}
__device__ __forceinline__ project2nd(const project2nd& other):binary_function<T1,T2,T2>(){}
__device__ __forceinline__ project2nd():binary_function<T1,T2,T2>(){}
};
// Min/Max Operations
@ -231,6 +302,8 @@ namespace cv { namespace gpu { namespace device
template <> struct name<type> : binary_function<type, type, type> \
{ \
__device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \
__device__ __forceinline__ name(const name& other):binary_function<type, type, type>(){}\
__device__ __forceinline__ name():binary_function<type, type, type>(){}\
};
template <typename T> struct maximum : binary_function<T, T, T>
@ -239,6 +312,8 @@ namespace cv { namespace gpu { namespace device
{
return lhs < rhs ? rhs : lhs;
}
__device__ __forceinline__ maximum(const maximum& other):binary_function<T, T, T>(){}
__device__ __forceinline__ maximum():binary_function<T, T, T>(){}
};
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, uchar, ::max)
@ -257,6 +332,8 @@ namespace cv { namespace gpu { namespace device
{
return lhs < rhs ? lhs : rhs;
}
__device__ __forceinline__ minimum(const minimum& other):binary_function<T, T, T>(){}
__device__ __forceinline__ minimum():binary_function<T, T, T>(){}
};
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, uchar, ::min)
@ -272,7 +349,7 @@ namespace cv { namespace gpu { namespace device
#undef OPENCV_GPU_IMPLEMENT_MINMAX
// Math functions
///bound=========================================
#define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \
template <typename T> struct name ## _func : unary_function<T, float> \
{ \
@ -342,17 +419,17 @@ namespace cv { namespace gpu { namespace device
};
// Saturate Cast Functor
template <typename T, typename D> struct saturate_cast_func : unary_function<T, D>
{
__device__ __forceinline__ D operator ()(typename TypeTraits<T>::ParameterType v) const
{
return saturate_cast<D>(v);
}
__device__ __forceinline__ saturate_cast_func(const saturate_cast_func& other):unary_function<T, D>(){}
__device__ __forceinline__ saturate_cast_func():unary_function<T, D>(){}
};
// Threshold Functors
template <typename T> struct thresh_binary_func : unary_function<T, T>
{
__host__ __device__ __forceinline__ thresh_binary_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {}
@ -361,10 +438,15 @@ namespace cv { namespace gpu { namespace device
{
return (src > thresh) * maxVal;
}
__device__ __forceinline__ thresh_binary_func(const thresh_binary_func& other)
: unary_function<T, T>(), thresh(other.thresh), maxVal(other.maxVal){}
__device__ __forceinline__ thresh_binary_func():unary_function<T, T>(){}
const T thresh;
const T maxVal;
};
template <typename T> struct thresh_binary_inv_func : unary_function<T, T>
{
__host__ __device__ __forceinline__ thresh_binary_inv_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {}
@ -373,10 +455,15 @@ namespace cv { namespace gpu { namespace device
{
return (src <= thresh) * maxVal;
}
__device__ __forceinline__ thresh_binary_inv_func(const thresh_binary_inv_func& other)
: unary_function<T, T>(), thresh(other.thresh), maxVal(other.maxVal){}
__device__ __forceinline__ thresh_binary_inv_func():unary_function<T, T>(){}
const T thresh;
const T maxVal;
};
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_) {}
@ -386,8 +473,14 @@ namespace cv { namespace gpu { namespace device
return minimum<T>()(src, thresh);
}
__device__ __forceinline__ thresh_trunc_func(const thresh_trunc_func& other)
: unary_function<T, T>(), thresh(other.thresh){}
__device__ __forceinline__ thresh_trunc_func():unary_function<T, T>(){}
const T thresh;
};
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_) {}
@ -396,9 +489,14 @@ namespace cv { namespace gpu { namespace device
{
return (src > thresh) * src;
}
__device__ __forceinline__ thresh_to_zero_func(const thresh_to_zero_func& other)
: unary_function<T, T>(), thresh(other.thresh){}
__device__ __forceinline__ thresh_to_zero_func():unary_function<T, T>(){}
const T thresh;
};
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_) {}
@ -407,12 +505,15 @@ namespace cv { namespace gpu { namespace device
{
return (src <= thresh) * src;
}
__device__ __forceinline__ thresh_to_zero_inv_func(const thresh_to_zero_inv_func& other)
: unary_function<T, T>(), thresh(other.thresh){}
__device__ __forceinline__ thresh_to_zero_inv_func():unary_function<T, T>(){}
const T thresh;
};
//bound!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! ============>
// Function Object Adaptors
template <typename Predicate> struct unary_negate : unary_function<typename Predicate::argument_type, bool>
{
explicit __host__ __device__ __forceinline__ unary_negate(const Predicate& p) : pred(p) {}

14
modules/gpu/src/opencv2/gpu/device/saturate_cast.hpp
View File

@ -84,7 +84,7 @@ namespace cv { namespace gpu { namespace device
}
template<> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
int iv = __double2int_rn(v);
return saturate_cast<uchar>(iv);
#else
@ -120,7 +120,7 @@ namespace cv { namespace gpu { namespace device
}
template<> __device__ __forceinline__ schar saturate_cast<schar>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
int iv = __double2int_rn(v);
return saturate_cast<schar>(iv);
#else
@ -151,7 +151,7 @@ namespace cv { namespace gpu { namespace device
}
template<> __device__ __forceinline__ ushort saturate_cast<ushort>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
int iv = __double2int_rn(v);
return saturate_cast<ushort>(iv);
#else
@ -178,7 +178,7 @@ namespace cv { namespace gpu { namespace device
}
template<> __device__ __forceinline__ short saturate_cast<short>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
int iv = __double2int_rn(v);
return saturate_cast<short>(iv);
#else
@ -192,7 +192,7 @@ namespace cv { namespace gpu { namespace device
}
template<> __device__ __forceinline__ int saturate_cast<int>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
return __double2int_rn(v);
#else
return saturate_cast<int>((float)v);
@ -205,7 +205,7 @@ namespace cv { namespace gpu { namespace device
}
template<> __device__ __forceinline__ uint saturate_cast<uint>(double v)
{
#if __CUDA_ARCH__ >= 130
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
return __double2uint_rn(v);
#else
return saturate_cast<uint>((float)v);
@ -213,4 +213,4 @@ namespace cv { namespace gpu { namespace device
}
}}}
#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */
#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */

4
modules/gpu/src/opencv2/gpu/device/transform.hpp
View File

@ -50,14 +50,14 @@
namespace cv { namespace gpu { namespace device
{
template <typename T, typename D, typename UnOp, typename Mask>
static inline void transform(DevMem2D_<T> src, DevMem2D_<D> dst, UnOp op, Mask mask, cudaStream_t stream)
static inline void transform(DevMem2D_<T> src, DevMem2D_<D> dst, UnOp op, const Mask& mask, cudaStream_t stream)
{
typedef TransformFunctorTraits<UnOp> ft;
transform_detail::TransformDispatcher<VecTraits<T>::cn == 1 && VecTraits<D>::cn == 1 && ft::smart_shift != 1>::call(src, dst, op, mask, stream);
}
template <typename T1, typename T2, typename D, typename BinOp, typename Mask>
static inline void transform(DevMem2D_<T1> src1, DevMem2D_<T2> src2, DevMem2D_<D> dst, BinOp op, Mask mask, cudaStream_t stream)
static inline void transform(DevMem2D_<T1> src1, DevMem2D_<T2> src2, DevMem2D_<D> dst, BinOp op, const Mask& mask, cudaStream_t stream)
{
typedef TransformFunctorTraits<BinOp> ft;
transform_detail::TransformDispatcher<VecTraits<T1>::cn == 1 && VecTraits<T2>::cn == 1 && VecTraits<D>::cn == 1 && ft::smart_shift != 1>::call(src1, src2, dst, op, mask, stream);

15
modules/gpu/src/opencv2/gpu/device/utility.hpp
View File

@ -70,6 +70,7 @@ namespace cv { namespace gpu { namespace device
struct SingleMask
{
explicit __host__ __device__ __forceinline__ SingleMask(PtrStepb mask_) : mask(mask_) {}
__host__ __device__ __forceinline__ SingleMask(const SingleMask& mask_): mask(mask_.mask){}
__device__ __forceinline__ bool operator()(int y, int x) const
{
@ -81,7 +82,10 @@ namespace cv { namespace gpu { namespace device
struct SingleMaskChannels
{
__host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_) : mask(mask_), channels(channels_) {}
__host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_)
: mask(mask_), channels(channels_) {}
__host__ __device__ __forceinline__ SingleMaskChannels(const SingleMaskChannels& mask_)
:mask(mask_.mask), channels(mask_.channels){}
__device__ __forceinline__ bool operator()(int y, int x) const
{
@ -94,7 +98,11 @@ namespace cv { namespace gpu { namespace device
struct MaskCollection
{
explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_) : maskCollection(maskCollection_) {}
explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_)
: maskCollection(maskCollection_) {}
__device__ __forceinline__ MaskCollection(const MaskCollection& masks_)
: maskCollection(masks_.maskCollection), curMask(masks_.curMask){}
__device__ __forceinline__ void next()
{
@ -117,6 +125,9 @@ namespace cv { namespace gpu { namespace device
struct WithOutMask
{
__device__ __forceinline__ WithOutMask(){}
__device__ __forceinline__ WithOutMask(const WithOutMask& mask){}
__device__ __forceinline__ void next() const
{
}

5
modules/gpu/test/nvidia/NCVTest.hpp
View File

@ -11,7 +11,9 @@
#ifndef _ncvtest_hpp_
#define _ncvtest_hpp_
#pragma warning( disable : 4201 4408 4127 4100)
#if defined _MSC_VER
# pragma warning( disable : 4201 4408 4127 4100)
#endif
#include <string>
#include <vector>
@ -36,6 +38,7 @@ class INCVTest
public:
virtual bool executeTest(NCVTestReport &report) = 0;
virtual std::string getName() const = 0;
virtual ~INCVTest(){}
};