fixed hundreds of "anonymous" warnings for gpu module.
authorMarina Kolpakova <no@email>
Fri, 8 Jun 2012 17:09:38 +0000 (17:09 +0000)
committerMarina Kolpakova <no@email>
Fri, 8 Jun 2012 17:09:38 +0000 (17:09 +0000)
13 files changed:
modules/core/CMakeLists.txt
modules/gpu/CMakeLists.txt
modules/gpu/src/cuda/matrix_reductions.cu
modules/gpu/src/cuda/surf.cu
modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp
modules/gpu/src/opencv2/gpu/device/datamov_utils.hpp
modules/gpu/src/opencv2/gpu/device/detail/color_detail.hpp
modules/gpu/src/opencv2/gpu/device/functional.hpp
modules/gpu/src/opencv2/gpu/device/saturate_cast.hpp
modules/gpu/src/opencv2/gpu/device/transform.hpp
modules/gpu/src/opencv2/gpu/device/utility.hpp
modules/gpu/test/nvidia/NCVTest.hpp

index de9c846..37e6e95 100644 (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})
index 5e7b185..b5f19f8 100644 (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-;")
index fbf5ce5..d6b6d94 100644 (file)
@@ -87,7 +87,9 @@ namespace cv { namespace gpu { namespace device
             __device__ __forceinline__ bool operator()(int y, int x) const \r
             { \r
                 return true; \r
-            } \r
+            }\r
+            __device__ __forceinline__ MaskTrue(){}\r
+            __device__ __forceinline__ MaskTrue(const MaskTrue& mask_){}\r
         };\r
 \r
         //////////////////////////////////////////////////////////////////////////////\r
@@ -1795,6 +1797,9 @@ namespace cv { namespace gpu { namespace device
                 return 0;\r
             }\r
 \r
+            __device__ __forceinline__ SumReductor(const SumReductor& other){}\r
+            __device__ __forceinline__ SumReductor(){}\r
+\r
             __device__ __forceinline__ S operator ()(volatile S a, volatile S b) const\r
             {\r
                 return a + b;\r
@@ -1813,6 +1818,9 @@ namespace cv { namespace gpu { namespace device
                 return 0;\r
             }\r
 \r
+            __device__ __forceinline__ AvgReductor(const AvgReductor& other){}\r
+            __device__ __forceinline__ AvgReductor(){}\r
+\r
             __device__ __forceinline__ S operator ()(volatile S a, volatile S b) const\r
             {\r
                 return a + b;\r
@@ -1831,6 +1839,9 @@ namespace cv { namespace gpu { namespace device
                 return numeric_limits<S>::max();\r
             }\r
 \r
+            __device__ __forceinline__ MinReductor(const MinReductor& other){}\r
+            __device__ __forceinline__ MinReductor(){}\r
+\r
             template <typename T> __device__ __forceinline__ T operator ()(volatile T a, volatile T b) const\r
             {\r
                 return saturate_cast<T>(::min(a, b));\r
@@ -1853,6 +1864,9 @@ namespace cv { namespace gpu { namespace device
                 return numeric_limits<S>::min();\r
             }\r
 \r
+            __device__ __forceinline__ MaxReductor(const MaxReductor& other){}\r
+            __device__ __forceinline__ MaxReductor(){}\r
+\r
             template <typename T> __device__ __forceinline__ int operator ()(volatile T a, volatile T b) const\r
             {\r
                 return ::max(a, b);\r
index ac7b6c2..e65f4aa 100644 (file)
@@ -116,7 +116,7 @@ namespace cv { namespace gpu { namespace device
 \r
         template <int N> __device__ float icvCalcHaarPatternSum(const float src[][5], int oldSize, int newSize, int y, int x)\r
         {\r
-        #if __CUDA_ARCH__ >= 200\r
+        #if __CUDA_ARCH__ && __CUDA_ARCH__ >= 200\r
             typedef double real_t;\r
         #else\r
             typedef float  real_t;\r
@@ -248,7 +248,7 @@ namespace cv { namespace gpu { namespace device
         template <typename Mask>\r
         __global__ void icvFindMaximaInLayer(const PtrStepf det, const PtrStepf trace, int4* maxPosBuffer, unsigned int* maxCounter)\r
         {\r
-            #if __CUDA_ARCH__ >= 110\r
+            #if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110\r
 \r
             extern __shared__ float N9[];\r
 \r
@@ -371,7 +371,7 @@ namespace cv { namespace gpu { namespace device
             float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian,\r
             unsigned int* featureCounter)\r
         {\r
-            #if __CUDA_ARCH__ >= 110\r
+            #if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110\r
 \r
             const int4 maxPos = maxPosBuffer[blockIdx.x];\r
 \r
index fded861..781d411 100644 (file)
@@ -231,7 +231,7 @@ __device__ Ncv32u d_outMaskPosition;
 \r
 __device__ void compactBlockWriteOutAnchorParallel(Ncv32u threadPassFlag, Ncv32u threadElem, Ncv32u *vectorOut)\r
 {\r
-#if __CUDA_ARCH__ >= 110\r
+#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110\r
     \r
     __shared__ Ncv32u shmem[NUM_THREADS_ANCHORSPARALLEL * 2];\r
     __shared__ Ncv32u numPassed;\r
@@ -587,7 +587,7 @@ __global__ void applyHaarClassifierClassifierParallel(Ncv32u *d_IImg, Ncv32u IIm
     }\r
     else\r
     {\r
-#if __CUDA_ARCH__ >= 110\r
+#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110\r
         if (bPass && !threadIdx.x)\r
         {\r
             Ncv32u outMaskOffset = atomicAdd(&d_outMaskPosition, 1);\r
index 2fcc2c5..a13d344 100644 (file)
@@ -41,7 +41,7 @@
 \r
 #ifndef _ncvruntimetemplates_hpp_\r
 #define _ncvruntimetemplates_hpp_\r
-#if _MSC_VER >= 1200\r
+#if defined _MSC_VER &&_MSC_VER >= 1200\r
 #pragma warning( disable: 4800 )\r
 #endif\r
 \r
index 50b9c7e..bd5c49f 100644 (file)
@@ -47,7 +47,7 @@
 \r
 namespace cv { namespace gpu { namespace device \r
 {\r
-    #if __CUDA_ARCH__ >= 200\r
+    #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200\r
 \r
         // for Fermi memory space is detected automatically\r
         template <typename T> struct ForceGlob\r
index 2f1a65a..2b2115e 100644 (file)
@@ -114,6 +114,11 @@ namespace cv { namespace gpu { namespace device
 \r
                 return dst;\r
             }\r
+\r
+            __device__ __forceinline__ RGB2RGB()\r
+                : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
+            __device__ __forceinline__ RGB2RGB(const RGB2RGB& other_)\r
+                :unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
         };\r
 \r
         template <> struct RGB2RGB<uchar, 4, 4, 2> : unary_function<uint, uint>\r
@@ -129,6 +134,9 @@ namespace cv { namespace gpu { namespace device
 \r
                 return dst;\r
             }\r
+\r
+            __device__ __forceinline__ RGB2RGB():unary_function<uint, uint>(){}\r
+            __device__ __forceinline__ RGB2RGB(const RGB2RGB& other_):unary_function<uint, uint>(){}\r
         };\r
     }\r
 \r
@@ -184,6 +192,8 @@ namespace cv { namespace gpu { namespace device
             {\r
                 return RGB2RGB5x5Converter<green_bits, bidx>::cvt(src);\r
             }\r
+            __device__ __forceinline__ RGB2RGB5x5():unary_function<uchar3, ushort>(){}\r
+            __device__ __forceinline__ RGB2RGB5x5(const RGB2RGB5x5& other_):unary_function<uchar3, ushort>(){}\r
         };\r
         template<int bidx, int green_bits> struct RGB2RGB5x5<4, bidx,green_bits> : unary_function<uint, ushort>\r
         {\r
@@ -191,6 +201,9 @@ namespace cv { namespace gpu { namespace device
             {\r
                 return RGB2RGB5x5Converter<green_bits, bidx>::cvt(src);\r
             }\r
+            __device__ __forceinline__ RGB2RGB5x5():unary_function<uint, ushort>(){}\r
+            __device__ __forceinline__ RGB2RGB5x5(const RGB2RGB5x5& other_):unary_function<uint, ushort>(){}\r
+\r
         };\r
     }\r
 \r
@@ -252,7 +265,11 @@ namespace cv { namespace gpu { namespace device
                 RGB5x52RGBConverter<green_bits, bidx>::cvt(src, dst);\r
                 return dst;\r
             }\r
+            __device__ __forceinline__ RGB5x52RGB():unary_function<ushort, uchar3>(){}\r
+            __device__ __forceinline__ RGB5x52RGB(const RGB5x52RGB& other_):unary_function<ushort, uchar3>(){}\r
+\r
         };\r
+\r
         template <int bidx, int green_bits> struct RGB5x52RGB<4, bidx, green_bits> : unary_function<ushort, uint>\r
         {\r
             __device__ __forceinline__ uint operator()(ushort src) const\r
@@ -261,6 +278,8 @@ namespace cv { namespace gpu { namespace device
                 RGB5x52RGBConverter<green_bits, bidx>::cvt(src, dst);\r
                 return dst;\r
             }\r
+            __device__ __forceinline__ RGB5x52RGB():unary_function<ushort, uint>(){}\r
+            __device__ __forceinline__ RGB5x52RGB(const RGB5x52RGB& other_):unary_function<ushort, uint>(){}\r
         };\r
     }\r
 \r
@@ -289,6 +308,8 @@ namespace cv { namespace gpu { namespace device
 \r
                 return dst;\r
             }\r
+            __device__ __forceinline__ Gray2RGB():unary_function<T, typename TypeVec<T, dcn>::vec_type>(){}\r
+            __device__ __forceinline__ Gray2RGB(const Gray2RGB& other_):unary_function<T, typename TypeVec<T, dcn>::vec_type>(){}\r
         };\r
         template <> struct Gray2RGB<uchar, 4> : unary_function<uchar, uint>\r
         {\r
@@ -302,6 +323,8 @@ namespace cv { namespace gpu { namespace device
 \r
                 return dst;\r
             }\r
+            __device__ __forceinline__ Gray2RGB():unary_function<uchar, uint>(){}\r
+            __device__ __forceinline__ Gray2RGB(const Gray2RGB& other_):unary_function<uchar, uint>(){}\r
         };\r
     }\r
 \r
@@ -340,6 +363,8 @@ namespace cv { namespace gpu { namespace device
             {\r
                 return Gray2RGB5x5Converter<green_bits>::cvt(src);\r
             }\r
+            __device__ __forceinline__ Gray2RGB5x5():unary_function<uchar, ushort>(){}\r
+            __device__ __forceinline__ Gray2RGB5x5(const Gray2RGB5x5& other_):unary_function<uchar, ushort>(){}\r
         };\r
     }\r
 \r
@@ -471,6 +496,8 @@ namespace cv { namespace gpu { namespace device
                 RGB2YUVConvert<bidx>(&src.x, dst);\r
                 return dst;\r
             }\r
+            __device__ __forceinline__ RGB2YUV():unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
+            __device__ __forceinline__ RGB2YUV(const RGB2YUV& other_):unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
         };\r
     }\r
 \r
@@ -535,7 +562,10 @@ namespace cv { namespace gpu { namespace device
 \r
                 return dst;\r
             }\r
+            __device__ __forceinline__ YUV2RGB():unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
+            __device__ __forceinline__ YUV2RGB(const YUV2RGB& other_):unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
         };\r
+\r
         template <int bidx> struct YUV2RGB<uchar, 4, 4, bidx> : unary_function<uint, uint>\r
         {\r
             __device__ __forceinline__ uint operator ()(uint src) const\r
@@ -605,7 +635,10 @@ namespace cv { namespace gpu { namespace device
                 RGB2YCrCbConvert<bidx>(&src.x, dst);\r
                 return dst;\r
             }\r
+            __device__ __forceinline__ RGB2YCrCb():unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
+            __device__ __forceinline__ RGB2YCrCb(const RGB2YCrCb& other_):unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
         };\r
+\r
         template <int bidx> struct RGB2YCrCb<uchar, 4, 4, bidx> : unary_function<uint, uint>\r
         {\r
             __device__ __forceinline__ uint operator ()(uint src) const\r
@@ -676,7 +709,10 @@ namespace cv { namespace gpu { namespace device
 \r
                 return dst;\r
             }\r
+            __device__ __forceinline__ YCrCb2RGB():unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
+            __device__ __forceinline__ YCrCb2RGB(const YCrCb2RGB& other_):unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}\r
         };\r
+\r
         template <int bidx> struct YCrCb2RGB<uchar, 4, 4, bidx> : unary_function<uint, uint>\r
         {\r
             __device__ __forceinline__ uint operator ()(uint src) const\r
@@ -1331,6 +1367,7 @@ namespace cv { namespace gpu { namespace device
             {\r
                 return HLS2RGBConvert<bidx, hr>(src);\r
             }\r
+\r
         };\r
     }\r
 \r
index d21f728..435fe65 100644 (file)
@@ -56,158 +56,224 @@ namespace cv { namespace gpu { namespace device
     using thrust::binary_function;\r
 \r
     // Arithmetic Operations\r
-\r
     template <typename T> struct plus : binary_function<T, T, T>\r
     {\r
-        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                 typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a + b;\r
         }\r
+        __device__ __forceinline__ plus(const plus& other):binary_function<T,T,T>(){}\r
+        __device__ __forceinline__ plus():binary_function<T,T,T>(){}\r
     };\r
+\r
     template <typename T> struct minus : binary_function<T, T, T>\r
     {\r
-        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                 typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a - b;\r
         }\r
+        __device__ __forceinline__ minus(const minus& other):binary_function<T,T,T>(){}\r
+        __device__ __forceinline__ minus():binary_function<T,T,T>(){}\r
     };\r
+\r
     template <typename T> struct multiplies : binary_function<T, T, T>\r
     {\r
-        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                 typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a * b;\r
         }\r
+        __device__ __forceinline__ multiplies(const multiplies& other):binary_function<T,T,T>(){}\r
+        __device__ __forceinline__ multiplies():binary_function<T,T,T>(){}\r
     };\r
+\r
     template <typename T> struct divides : binary_function<T, T, T>\r
     {\r
-        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                 typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a / b;\r
         }\r
+        __device__ __forceinline__ divides(const divides& other):binary_function<T,T,T>(){}\r
+        __device__ __forceinline__ divides():binary_function<T,T,T>(){}\r
     };\r
+\r
     template <typename T> struct modulus : binary_function<T, T, T>\r
     {\r
-        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                 typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a % b;\r
         }\r
+        __device__ __forceinline__ modulus(const modulus& other):binary_function<T,T,T>(){}\r
+        __device__ __forceinline__ modulus():binary_function<T,T,T>(){}\r
     };\r
+\r
     template <typename T> struct negate : unary_function<T, T>\r
     {\r
         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a) const\r
         {\r
             return -a;\r
         }\r
+        __device__ __forceinline__ negate(const negate& other):unary_function<T,T>(){}\r
+        __device__ __forceinline__ negate():unary_function<T,T>(){}\r
     };\r
 \r
     // Comparison Operations\r
-\r
     template <typename T> struct equal_to : binary_function<T, T, bool>\r
     {\r
-        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                    typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a == b;\r
         }\r
+        __device__ __forceinline__ equal_to(const equal_to& other):binary_function<T,T,bool>(){}\r
+        __device__ __forceinline__ equal_to():binary_function<T,T,bool>(){}\r
     };\r
+\r
     template <typename T> struct not_equal_to : binary_function<T, T, bool>\r
     {\r
-        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                    typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a != b;\r
         }\r
+        __device__ __forceinline__ not_equal_to(const not_equal_to& other):binary_function<T,T,bool>(){}\r
+        __device__ __forceinline__ not_equal_to():binary_function<T,T,bool>(){}\r
     };\r
+\r
     template <typename T> struct greater : binary_function<T, T, bool>\r
     {\r
-        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                    typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a > b;\r
         }\r
+        __device__ __forceinline__ greater(const greater& other):binary_function<T,T,bool>(){}\r
+        __device__ __forceinline__ greater():binary_function<T,T,bool>(){}\r
     };\r
+\r
     template <typename T> struct less : binary_function<T, T, bool>\r
     {\r
-        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                    typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a < b;\r
         }\r
+        __device__ __forceinline__ less(const less& other):binary_function<T,T,bool>(){}\r
+        __device__ __forceinline__ less():binary_function<T,T,bool>(){}\r
     };\r
+\r
     template <typename T> struct greater_equal : binary_function<T, T, bool>\r
     {\r
-        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                    typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a >= b;\r
         }\r
+        __device__ __forceinline__ greater_equal(const greater_equal& other):binary_function<T,T,bool>(){}\r
+        __device__ __forceinline__ greater_equal():binary_function<T,T,bool>(){}\r
     };\r
+\r
     template <typename T> struct less_equal : binary_function<T, T, bool>\r
     {\r
-        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                    typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a <= b;\r
         }\r
+        __device__ __forceinline__ less_equal(const less_equal& other):binary_function<T,T,bool>(){}\r
+        __device__ __forceinline__ less_equal():binary_function<T,T,bool>(){}\r
     };\r
 \r
     // Logical Operations\r
-\r
     template <typename T> struct logical_and : binary_function<T, T, bool>\r
     {\r
-        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                    typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a && b;\r
         }\r
+        __device__ __forceinline__ logical_and(const logical_and& other):binary_function<T,T,bool>(){}\r
+        __device__ __forceinline__ logical_and():binary_function<T,T,bool>(){}\r
     };\r
+\r
     template <typename T> struct logical_or : binary_function<T, T, bool>\r
     {\r
-        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                    typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a || b;\r
         }\r
+        __device__ __forceinline__ logical_or(const logical_or& other):binary_function<T,T,bool>(){}\r
+        __device__ __forceinline__ logical_or():binary_function<T,T,bool>(){}\r
     };\r
+\r
     template <typename T> struct logical_not : unary_function<T, bool>\r
     {\r
         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a) const\r
         {\r
             return !a;\r
         }\r
+        __device__ __forceinline__ logical_not(const logical_not& other):unary_function<T,bool>(){}\r
+        __device__ __forceinline__ logical_not():unary_function<T,bool>(){}\r
     };\r
 \r
     // Bitwise Operations\r
-\r
     template <typename T> struct bit_and : binary_function<T, T, T>\r
     {\r
-        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                 typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a & b;\r
         }\r
+        __device__ __forceinline__ bit_and(const bit_and& other):binary_function<T,T,T>(){}\r
+        __device__ __forceinline__ bit_and():binary_function<T,T,T>(){}\r
     };\r
+\r
     template <typename T> struct bit_or : binary_function<T, T, T>\r
     {\r
-        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                 typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a | b;\r
         }\r
+        __device__ __forceinline__ bit_or(const bit_or& other):binary_function<T,T,T>(){}\r
+        __device__ __forceinline__ bit_or():binary_function<T,T,T>(){}\r
     };\r
+\r
     template <typename T> struct bit_xor : binary_function<T, T, T>\r
     {\r
-        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const\r
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
+                                                 typename TypeTraits<T>::ParameterType b) const\r
         {\r
             return a ^ b;\r
         }\r
+        __device__ __forceinline__ bit_xor(const bit_xor& other):binary_function<T,T,T>(){}\r
+        __device__ __forceinline__ bit_xor():binary_function<T,T,T>(){}\r
     };\r
+\r
     template <typename T> struct bit_not : unary_function<T, T>\r
     {\r
         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType v) const \r
         {\r
             return ~v;\r
         }\r
+        __device__ __forceinline__ bit_not(const bit_not& other):unary_function<T,T>(){}\r
+        __device__ __forceinline__ bit_not():unary_function<T,T>(){}\r
     };\r
 \r
     // Generalized Identity Operations\r
-\r
     template <typename T> struct identity : unary_function<T, T>\r
     {\r
         __device__ __forceinline__ typename TypeTraits<T>::ParameterType operator()(typename TypeTraits<T>::ParameterType x) const \r
         {\r
             return x;\r
         }\r
+        __device__ __forceinline__ identity(const identity& other):unary_function<T,T>(){}\r
+        __device__ __forceinline__ identity():unary_function<T,T>(){}\r
     };\r
 \r
     template <typename T1, typename T2> struct project1st : binary_function<T1, T2, T1>\r
@@ -216,13 +282,18 @@ namespace cv { namespace gpu { namespace device
         {\r
             return lhs;\r
         }\r
+        __device__ __forceinline__ project1st(const project1st& other):binary_function<T1,T2,T1>(){}\r
+        __device__ __forceinline__ project1st():binary_function<T1,T2,T1>(){}\r
     };\r
+\r
     template <typename T1, typename T2> struct project2nd : binary_function<T1, T2, T2>\r
     {\r
         __device__ __forceinline__ typename TypeTraits<T2>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const \r
         {\r
             return rhs;\r
         }\r
+        __device__ __forceinline__ project2nd(const project2nd& other):binary_function<T1,T2,T2>(){}\r
+        __device__ __forceinline__ project2nd():binary_function<T1,T2,T2>(){}\r
     };\r
 \r
     // Min/Max Operations\r
@@ -231,6 +302,8 @@ namespace cv { namespace gpu { namespace device
     template <> struct name<type> : binary_function<type, type, type> \\r
     { \\r
         __device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \\r
+        __device__ __forceinline__ name(const name& other):binary_function<type, type, type>(){}\\r
+        __device__ __forceinline__ name():binary_function<type, type, type>(){}\\r
     };\r
 \r
     template <typename T> struct maximum : binary_function<T, T, T>\r
@@ -239,6 +312,8 @@ namespace cv { namespace gpu { namespace device
         {\r
             return lhs < rhs ? rhs : lhs;\r
         }\r
+        __device__ __forceinline__ maximum(const maximum& other):binary_function<T, T, T>(){}\r
+        __device__ __forceinline__ maximum():binary_function<T, T, T>(){}\r
     };\r
 \r
     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, uchar, ::max)\r
@@ -257,6 +332,8 @@ namespace cv { namespace gpu { namespace device
         {\r
             return lhs < rhs ? lhs : rhs;\r
         }\r
+        __device__ __forceinline__ minimum(const minimum& other):binary_function<T, T, T>(){}\r
+        __device__ __forceinline__ minimum():binary_function<T, T, T>(){}\r
     };\r
 \r
     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, uchar, ::min)\r
@@ -272,7 +349,7 @@ namespace cv { namespace gpu { namespace device
 #undef OPENCV_GPU_IMPLEMENT_MINMAX\r
 \r
     // Math functions\r
-\r
+///bound=========================================\r
 #define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \\r
     template <typename T> struct name ## _func : unary_function<T, float> \\r
     { \\r
@@ -342,17 +419,17 @@ namespace cv { namespace gpu { namespace device
     };\r
 \r
     // Saturate Cast Functor\r
-\r
     template <typename T, typename D> struct saturate_cast_func : unary_function<T, D>\r
     {\r
         __device__ __forceinline__ D operator ()(typename TypeTraits<T>::ParameterType v) const\r
         {\r
             return saturate_cast<D>(v);\r
         }\r
+        __device__ __forceinline__ saturate_cast_func(const saturate_cast_func& other):unary_function<T, D>(){}\r
+        __device__ __forceinline__ saturate_cast_func():unary_function<T, D>(){}\r
     };\r
 \r
     // Threshold Functors\r
-\r
     template <typename T> struct thresh_binary_func : unary_function<T, T>\r
     {\r
         __host__ __device__ __forceinline__ thresh_binary_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {}\r
@@ -361,10 +438,15 @@ namespace cv { namespace gpu { namespace device
         {\r
             return (src > thresh) * maxVal;\r
         }\r
+        __device__ __forceinline__ thresh_binary_func(const thresh_binary_func& other)\r
+            : unary_function<T, T>(), thresh(other.thresh), maxVal(other.maxVal){}\r
+\r
+        __device__ __forceinline__ thresh_binary_func():unary_function<T, T>(){}\r
 \r
         const T thresh;\r
         const T maxVal;\r
     };\r
+\r
     template <typename T> struct thresh_binary_inv_func : unary_function<T, T>\r
     {\r
         __host__ __device__ __forceinline__ thresh_binary_inv_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {}\r
@@ -373,10 +455,15 @@ namespace cv { namespace gpu { namespace device
         {\r
             return (src <= thresh) * maxVal;\r
         }\r
+        __device__ __forceinline__ thresh_binary_inv_func(const thresh_binary_inv_func& other)\r
+            : unary_function<T, T>(), thresh(other.thresh), maxVal(other.maxVal){}\r
+\r
+        __device__ __forceinline__ thresh_binary_inv_func():unary_function<T, T>(){}\r
 \r
         const T thresh;\r
         const T maxVal;\r
     };\r
+\r
     template <typename T> struct thresh_trunc_func : unary_function<T, T>\r
     {\r
         explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {}\r
@@ -386,8 +473,14 @@ namespace cv { namespace gpu { namespace device
             return minimum<T>()(src, thresh);\r
         }\r
 \r
+        __device__ __forceinline__ thresh_trunc_func(const thresh_trunc_func& other)\r
+            : unary_function<T, T>(), thresh(other.thresh){}\r
+\r
+        __device__ __forceinline__ thresh_trunc_func():unary_function<T, T>(){}\r
+\r
         const T thresh;\r
     };\r
+\r
     template <typename T> struct thresh_to_zero_func : unary_function<T, T>\r
     {\r
         explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {}\r
@@ -396,9 +489,14 @@ namespace cv { namespace gpu { namespace device
         {\r
             return (src > thresh) * src;\r
         }\r
+        __device__ __forceinline__ thresh_to_zero_func(const thresh_to_zero_func& other)\r
+            : unary_function<T, T>(), thresh(other.thresh){}\r
+\r
+        __device__ __forceinline__ thresh_to_zero_func():unary_function<T, T>(){}\r
 \r
         const T thresh;\r
     };\r
+\r
     template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T>\r
     {\r
         explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {}\r
@@ -407,12 +505,15 @@ namespace cv { namespace gpu { namespace device
         {\r
             return (src <= thresh) * src;\r
         }\r
+        __device__ __forceinline__ thresh_to_zero_inv_func(const thresh_to_zero_inv_func& other)\r
+            : unary_function<T, T>(), thresh(other.thresh){}\r
+\r
+        __device__ __forceinline__ thresh_to_zero_inv_func():unary_function<T, T>(){}\r
 \r
         const T thresh;\r
     };\r
-\r
+//bound!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! ============>\r
     // Function Object Adaptors\r
-\r
     template <typename Predicate> struct unary_negate : unary_function<typename Predicate::argument_type, bool>\r
     {\r
       explicit __host__ __device__ __forceinline__ unary_negate(const Predicate& p) : pred(p) {}\r
index 35575a2..d9fa5ce 100644 (file)
@@ -84,7 +84,7 @@ namespace cv { namespace gpu { namespace device
     }\r
     template<> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)\r
     {\r
-    #if __CUDA_ARCH__ >= 130\r
+    #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
         int iv = __double2int_rn(v); \r
         return saturate_cast<uchar>(iv);\r
     #else\r
@@ -120,7 +120,7 @@ namespace cv { namespace gpu { namespace device
     }\r
     template<> __device__ __forceinline__ schar saturate_cast<schar>(double v)\r
     {             \r
-    #if __CUDA_ARCH__ >= 130\r
+    #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
         int iv = __double2int_rn(v); \r
         return saturate_cast<schar>(iv);\r
     #else\r
@@ -151,7 +151,7 @@ namespace cv { namespace gpu { namespace device
     }\r
     template<> __device__ __forceinline__ ushort saturate_cast<ushort>(double v)\r
     {             \r
-    #if __CUDA_ARCH__ >= 130\r
+    #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
         int iv = __double2int_rn(v); \r
         return saturate_cast<ushort>(iv);\r
     #else\r
@@ -178,7 +178,7 @@ namespace cv { namespace gpu { namespace device
     }\r
     template<> __device__ __forceinline__ short saturate_cast<short>(double v)\r
     {            \r
-    #if __CUDA_ARCH__ >= 130\r
+    #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
         int iv = __double2int_rn(v); \r
         return saturate_cast<short>(iv);\r
     #else\r
@@ -192,7 +192,7 @@ namespace cv { namespace gpu { namespace device
     }\r
     template<> __device__ __forceinline__ int saturate_cast<int>(double v) \r
     {\r
-    #if __CUDA_ARCH__ >= 130 \r
+    #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
         return __double2int_rn(v);\r
     #else\r
         return saturate_cast<int>((float)v);\r
@@ -205,7 +205,7 @@ namespace cv { namespace gpu { namespace device
     }\r
     template<> __device__ __forceinline__ uint saturate_cast<uint>(double v) \r
     {            \r
-    #if __CUDA_ARCH__ >= 130\r
+    #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
         return __double2uint_rn(v);\r
     #else\r
         return saturate_cast<uint>((float)v);\r
@@ -213,4 +213,4 @@ namespace cv { namespace gpu { namespace device
     }\r
 }}}\r
 \r
-#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */
\ No newline at end of file
+#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */\r
index 89eed7e..a0e79df 100644 (file)
 namespace cv { namespace gpu { namespace device \r
 {\r
     template <typename T, typename D, typename UnOp, typename Mask>\r
-    static inline void transform(DevMem2D_<T> src, DevMem2D_<D> dst, UnOp op, Mask mask, cudaStream_t stream)\r
+    static inline void transform(DevMem2D_<T> src, DevMem2D_<D> dst, UnOp op, const Mask& mask, cudaStream_t stream)\r
     {\r
         typedef TransformFunctorTraits<UnOp> ft;\r
         transform_detail::TransformDispatcher<VecTraits<T>::cn == 1 && VecTraits<D>::cn == 1 && ft::smart_shift != 1>::call(src, dst, op, mask, stream);\r
     }\r
 \r
     template <typename T1, typename T2, typename D, typename BinOp, typename Mask>\r
-    static inline void transform(DevMem2D_<T1> src1, DevMem2D_<T2> src2, DevMem2D_<D> dst, BinOp op, Mask mask, cudaStream_t stream)\r
+    static inline void transform(DevMem2D_<T1> src1, DevMem2D_<T2> src2, DevMem2D_<D> dst, BinOp op, const Mask& mask, cudaStream_t stream)\r
     {\r
         typedef TransformFunctorTraits<BinOp> ft;\r
         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);\r
index cb702b7..cb4db80 100644 (file)
@@ -70,6 +70,7 @@ namespace cv { namespace gpu { namespace device
     struct SingleMask\r
     {\r
         explicit __host__ __device__ __forceinline__ SingleMask(PtrStepb mask_) : mask(mask_) {}\r
+        __host__ __device__ __forceinline__ SingleMask(const SingleMask& mask_): mask(mask_.mask){}\r
         \r
         __device__ __forceinline__ bool operator()(int y, int x) const\r
         {            \r
@@ -81,7 +82,10 @@ namespace cv { namespace gpu { namespace device
 \r
     struct SingleMaskChannels\r
     {\r
-        __host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_) : mask(mask_), channels(channels_) {}\r
+        __host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_) 
+        : mask(mask_), channels(channels_) {}\r
+        __host__ __device__ __forceinline__ SingleMaskChannels(const SingleMaskChannels& mask_)\r
+            :mask(mask_.mask), channels(mask_.channels){}\r
         \r
         __device__ __forceinline__ bool operator()(int y, int x) const\r
         {            \r
@@ -94,7 +98,11 @@ namespace cv { namespace gpu { namespace device
 \r
     struct MaskCollection\r
     {\r
-        explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_) : maskCollection(maskCollection_) {}\r
+        explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_)\r
+            : maskCollection(maskCollection_) {}\r
+\r
+        __device__ __forceinline__ MaskCollection(const MaskCollection& masks_)\r
+            : maskCollection(masks_.maskCollection), curMask(masks_.curMask){}\r
 \r
         __device__ __forceinline__ void next()\r
         {\r
@@ -117,6 +125,9 @@ namespace cv { namespace gpu { namespace device
 \r
     struct WithOutMask\r
     {\r
+        __device__ __forceinline__ WithOutMask(){}\r
+        __device__ __forceinline__ WithOutMask(const WithOutMask& mask){}\r
+\r
         __device__ __forceinline__ void next() const\r
         {\r
         }\r
index 00be0f0..94ec46c 100644 (file)
@@ -11,7 +11,9 @@
 #ifndef _ncvtest_hpp_\r
 #define _ncvtest_hpp_\r
 \r
-#pragma warning( disable : 4201 4408 4127 4100)\r
+#if defined _MSC_VER\r
+# pragma warning( disable : 4201 4408 4127 4100)\r
+#endif\r
 \r
 #include <string>\r
 #include <vector>\r
@@ -36,6 +38,7 @@ class INCVTest
 public:\r
     virtual bool executeTest(NCVTestReport &report) = 0;\r
     virtual std::string getName() const = 0;\r
+    virtual ~INCVTest(){}\r
 };\r
 \r
 \r