From e86f0aaea1d356f1bb3eb899082610e885b550d4 Mon Sep 17 00:00:00 2001
From: Marina Kolpakova <no@email>
Date: Fri, 8 Jun 2012 17:09:38 +0000
Subject: [PATCH] fixed hundreds of "anonymous" warnings for gpu module.

---
 modules/core/CMakeLists.txt                   |   7 +-
 modules/gpu/CMakeLists.txt                    |   4 +-
 modules/gpu/src/cuda/matrix_reductions.cu     |  16 +-
 modules/gpu/src/cuda/surf.cu                  |   6 +-
 .../gpu/src/nvidia/NCVHaarObjectDetection.cu  |   4 +-
 .../src/nvidia/core/NCVRuntimeTemplates.hpp   |   2 +-
 .../src/opencv2/gpu/device/datamov_utils.hpp  |   2 +-
 .../gpu/device/detail/color_detail.hpp        |  37 +++++
 .../gpu/src/opencv2/gpu/device/functional.hpp | 153 +++++++++++++++---
 .../src/opencv2/gpu/device/saturate_cast.hpp  |  14 +-
 .../gpu/src/opencv2/gpu/device/transform.hpp  |   4 +-
 .../gpu/src/opencv2/gpu/device/utility.hpp    |  15 +-
 modules/gpu/test/nvidia/NCVTest.hpp           |   5 +-
 13 files changed, 220 insertions(+), 49 deletions(-)

diff --git a/modules/core/CMakeLists.txt b/modules/core/CMakeLists.txt
index de9c846894..37e6e95388 100644
--- a/modules/core/CMakeLists.txt
+++ b/modules/core/CMakeLists.txt
@@ -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})
diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt
index 5e7b18582a..b5f19f83ea 100644
--- a/modules/gpu/CMakeLists.txt
+++ b/modules/gpu/CMakeLists.txt
@@ -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-;")
diff --git a/modules/gpu/src/cuda/matrix_reductions.cu b/modules/gpu/src/cuda/matrix_reductions.cu
index fbf5ce5240..d6b6d94193 100644
--- a/modules/gpu/src/cuda/matrix_reductions.cu
+++ b/modules/gpu/src/cuda/matrix_reductions.cu
@@ -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);
diff --git a/modules/gpu/src/cuda/surf.cu b/modules/gpu/src/cuda/surf.cu
index ac7b6c27cb..e65f4aafc9 100644
--- a/modules/gpu/src/cuda/surf.cu
+++ b/modules/gpu/src/cuda/surf.cu
@@ -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];
 
diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu b/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
index fded861893..781d411c24 100644
--- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
+++ b/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
@@ -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);
diff --git a/modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp b/modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp
index 2fcc2c5fdf..a13d34489d 100644
--- a/modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp
+++ b/modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp
@@ -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
 
diff --git a/modules/gpu/src/opencv2/gpu/device/datamov_utils.hpp b/modules/gpu/src/opencv2/gpu/device/datamov_utils.hpp
index 50b9c7e49c..bd5c49fbec 100644
--- a/modules/gpu/src/opencv2/gpu/device/datamov_utils.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/datamov_utils.hpp
@@ -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
diff --git a/modules/gpu/src/opencv2/gpu/device/detail/color_detail.hpp b/modules/gpu/src/opencv2/gpu/device/detail/color_detail.hpp
index 2f1a65ad37..2b2115e1c0 100644
--- a/modules/gpu/src/opencv2/gpu/device/detail/color_detail.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/detail/color_detail.hpp
@@ -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);
             }
+
         };
     }
 
diff --git a/modules/gpu/src/opencv2/gpu/device/functional.hpp b/modules/gpu/src/opencv2/gpu/device/functional.hpp
index d21f728e31..435fe65f36 100644
--- a/modules/gpu/src/opencv2/gpu/device/functional.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/functional.hpp
@@ -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) {}
diff --git a/modules/gpu/src/opencv2/gpu/device/saturate_cast.hpp b/modules/gpu/src/opencv2/gpu/device/saturate_cast.hpp
index 35575a2b6d..d9fa5ce0c0 100644
--- a/modules/gpu/src/opencv2/gpu/device/saturate_cast.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/saturate_cast.hpp
@@ -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__ */
\ No newline at end of file
+#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */
diff --git a/modules/gpu/src/opencv2/gpu/device/transform.hpp b/modules/gpu/src/opencv2/gpu/device/transform.hpp
index 89eed7e678..a0e79dfa15 100644
--- a/modules/gpu/src/opencv2/gpu/device/transform.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/transform.hpp
@@ -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);
diff --git a/modules/gpu/src/opencv2/gpu/device/utility.hpp b/modules/gpu/src/opencv2/gpu/device/utility.hpp
index cb702b7392..cb4db80a04 100644
--- a/modules/gpu/src/opencv2/gpu/device/utility.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/utility.hpp
@@ -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
         {
         }
diff --git a/modules/gpu/test/nvidia/NCVTest.hpp b/modules/gpu/test/nvidia/NCVTest.hpp
index 00be0f0d19..94ec46ce6f 100644
--- a/modules/gpu/test/nvidia/NCVTest.hpp
+++ b/modules/gpu/test/nvidia/NCVTest.hpp
@@ -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(){}
 };