fixed -Wstrict_alliasing warning for GCC

modules/gpu/src/imgproc.cpp

@@ -261,6 +261,12 @@ namespace
     }
 }
 
+#if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__  > 4
+typedef Npp32s __attribute__((__may_alias__)) Npp32s_a;
+#else
+typedef Npp32s Npp32s_a;
+#endif
+
 void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value, Stream& s)
 {
     CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);
@@ -308,7 +314,7 @@ void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom
         case CV_32FC1:
             {
                 Npp32f val = saturate_cast<Npp32f>(value[0]);
-                Npp32s nVal = *(reinterpret_cast<Npp32s*>(&val));
+                Npp32s nVal = *(reinterpret_cast<Npp32s_a*>(&val));
                 nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), srcsz,
                     dst.ptr<Npp32s>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );
                 break;

modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp

@@ -67,7 +67,11 @@
 // Guaranteed size cross-platform classifier structures
 //
 //==============================================================================
-
+#if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__  > 4
+typedef Ncv32f __attribute__((__may_alias__)) Ncv32f_a;
+#else
+typedef Ncv32f Ncv32f_a;
+#endif
 
 struct HaarFeature64
 {
@@ -87,7 +91,7 @@ struct HaarFeature64
 
     __host__ NCVStatus setWeight(Ncv32f weight)
     {
-        ((Ncv32f*)&(this->_ui2.y))[0] = weight;
+        ((Ncv32f_a*)&(this->_ui2.y))[0] = weight;
         return NCV_SUCCESS;
     }
 
@@ -102,7 +106,7 @@ struct HaarFeature64
 
     __device__ __host__ Ncv32f getWeight(void)
     {
-        return *(Ncv32f*)(&this->_ui2.y);
+        return *(Ncv32f_a*)(&this->_ui2.y);
     }
 };
 
@@ -168,14 +172,13 @@ public:
     }
 };
 
-
 struct HaarClassifierNodeDescriptor32
 {
     uint1 _ui1;
 
     __host__ NCVStatus create(Ncv32f leafValue)
     {
-        *(Ncv32f *)&this->_ui1 = leafValue;
+        *(Ncv32f_a *)&this->_ui1 = leafValue;
         return NCV_SUCCESS;
     }
 
@@ -187,7 +190,7 @@ struct HaarClassifierNodeDescriptor32
 
     __host__ Ncv32f getLeafValueHost(void)
     {
-        return *(Ncv32f *)&this->_ui1.x;
+        return *(Ncv32f_a *)&this->_ui1.x;
     }
 
 #ifdef __CUDACC__
@@ -203,6 +206,11 @@ struct HaarClassifierNodeDescriptor32
     }
 };
 
+#if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__  > 4
+typedef Ncv32u __attribute__((__may_alias__)) Ncv32u_a;
+#else
+typedef Ncv32u Ncv32u_a;
+#endif
 
 struct HaarClassifierNode128
 {
@@ -216,19 +224,19 @@ struct HaarClassifierNode128
 
     __host__ NCVStatus setThreshold(Ncv32f t)
     {
-        this->_ui4.y = *(Ncv32u *)&t;
+        this->_ui4.y = *(Ncv32u_a *)&t;
         return NCV_SUCCESS;
     }
 
     __host__ NCVStatus setLeftNodeDesc(HaarClassifierNodeDescriptor32 nl)
     {
-        this->_ui4.z = *(Ncv32u *)&nl;
+        this->_ui4.z = *(Ncv32u_a *)&nl;
         return NCV_SUCCESS;
     }
 
     __host__ NCVStatus setRightNodeDesc(HaarClassifierNodeDescriptor32 nr)
     {
-        this->_ui4.w = *(Ncv32u *)&nr;
+        this->_ui4.w = *(Ncv32u_a *)&nr;
         return NCV_SUCCESS;
     }
 
@@ -239,7 +247,7 @@ struct HaarClassifierNode128
 
     __host__ __device__ Ncv32f getThreshold(void)
     {
-        return *(Ncv32f*)&this->_ui4.y;
+        return *(Ncv32f_a*)&this->_ui4.y;
     }
 
     __host__ __device__ HaarClassifierNodeDescriptor32 getLeftNodeDesc(void)
@@ -264,7 +272,7 @@ struct HaarStage64
 
     __host__ NCVStatus setStageThreshold(Ncv32f t)
     {
-        this->_ui2.x = *(Ncv32u *)&t;
+        this->_ui2.x = *(Ncv32u_a *)&t;
         return NCV_SUCCESS;
     }
 
@@ -290,7 +298,7 @@ struct HaarStage64
 
     __host__ __device__ Ncv32f getStageThreshold(void)
     {
-        return *(Ncv32f*)&this->_ui2.x;
+        return *(Ncv32f_a*)&this->_ui2.x;
     }
 
     __host__ __device__ Ncv32u getStartClassifierRootNodeOffset(void)

modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu

@@ -1423,7 +1423,7 @@ NCVStatus compactVector_32u_device(Ncv32u *d_src, Ncv32u srcLen,
                     (d_hierSums.ptr() + partSumOffsets[i],
                      partSumNums[i], NULL,
                      d_hierSums.ptr() + partSumOffsets[i+1],
-                     NULL);
+                     0);
             }
             else
             {
@@ -1433,7 +1433,7 @@ NCVStatus compactVector_32u_device(Ncv32u *d_src, Ncv32u srcLen,
                     (d_hierSums.ptr() + partSumOffsets[i],
                      partSumNums[i], NULL,
                      NULL,
-                     NULL);
+                     0);
             }
 
             ncvAssertCUDALastErrorReturn(NPPST_CUDA_KERNEL_EXECUTION_ERROR);
@@ -1557,16 +1557,21 @@ NCVStatus nppsStCompact_32s(Ncv32s *d_src, Ncv32u srcLen,
 }
 
 
+#if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__  > 4
+typedef Ncv32u __attribute__((__may_alias__)) Ncv32u_a;
+#else
+typedef Ncv32u Ncv32u_a;
+#endif
+
 NCVStatus nppsStCompact_32f(Ncv32f *d_src, Ncv32u srcLen,
                             Ncv32f *d_dst, Ncv32u *p_dstLen,
                             Ncv32f elemRemove, Ncv8u *pBuffer,
                             Ncv32u bufSize, cudaDeviceProp &devProp)
 {
     return nppsStCompact_32u((Ncv32u *)d_src, srcLen, (Ncv32u *)d_dst, p_dstLen,
-                             *(Ncv32u *)&elemRemove, pBuffer, bufSize, devProp);
+                             *(Ncv32u_a *)&elemRemove, pBuffer, bufSize, devProp);
 }
 
-
 NCVStatus nppsStCompact_32u_host(Ncv32u *h_src, Ncv32u srcLen,
                                  Ncv32u *h_dst, Ncv32u *dstLen, Ncv32u elemRemove)
 {
@@ -1602,17 +1607,16 @@ NCVStatus nppsStCompact_32u_host(Ncv32u *h_src, Ncv32u srcLen,
 NCVStatus nppsStCompact_32s_host(Ncv32s *h_src, Ncv32u srcLen,
                                  Ncv32s *h_dst, Ncv32u *dstLen, Ncv32s elemRemove)
 {
-    return nppsStCompact_32u_host((Ncv32u *)h_src, srcLen, (Ncv32u *)h_dst, dstLen, *(Ncv32u *)&elemRemove);
+    return nppsStCompact_32u_host((Ncv32u *)h_src, srcLen, (Ncv32u *)h_dst, dstLen, *(Ncv32u_a *)&elemRemove);
 }
 
 
 NCVStatus nppsStCompact_32f_host(Ncv32f *h_src, Ncv32u srcLen,
                                  Ncv32f *h_dst, Ncv32u *dstLen, Ncv32f elemRemove)
 {
-    return nppsStCompact_32u_host((Ncv32u *)h_src, srcLen, (Ncv32u *)h_dst, dstLen, *(Ncv32u *)&elemRemove);
+    return nppsStCompact_32u_host((Ncv32u *)h_src, srcLen, (Ncv32u *)h_dst, dstLen, *(Ncv32u_a *)&elemRemove);
 }
 
-
 //==============================================================================
 //
 // Filter.cu

modules/gpu/test/nvidia/TestHaarCascadeApplication.cpp

@@ -245,8 +245,8 @@ bool TestHaarCascadeApplication::process()
 
     int devId;
     ncvAssertCUDAReturn(cudaGetDevice(&devId), false);
-    cudaDeviceProp devProp;
-    ncvAssertCUDAReturn(cudaGetDeviceProperties(&devProp, devId), false);
+    cudaDeviceProp _devProp;
+    ncvAssertCUDAReturn(cudaGetDeviceProperties(&_devProp, devId), false);
 
     ncvStat = ncvApplyHaarClassifierCascade_device(
         d_integralImage, d_rectStdDev, d_pixelMask,
@@ -254,7 +254,7 @@ bool TestHaarCascadeApplication::process()
         haar, h_HaarStages, d_HaarStages, d_HaarNodes, d_HaarFeatures, false,
         searchRoiU, 1, 1.0f,
         *this->allocatorGPU.get(), *this->allocatorCPU.get(),
-        devProp, 0);
+        _devProp, 0);
     ncvAssertReturn(ncvStat == NCV_SUCCESS, false);
 
     NCVMatrixAlloc<Ncv32u> h_pixelMask_d(*this->allocatorCPU.get(), this->width, this->height);

modules/gpu/test/nvidia/main_nvidia.cpp

@@ -1,4 +1,6 @@
-#pragma warning (disable : 4408 4201 4100)
+#if defined _MSC_VER && _MSC_VER >= 1200
+# pragma warning (disable : 4408 4201 4100)
+#endif
 
 #include <cstdio>