From ede3781e3efdcf7e5333638c0808ea80190aa4dd Mon Sep 17 00:00:00 2001 From: "marina.kolpakova" Date: Mon, 20 Aug 2012 02:27:54 +0400 Subject: [PATCH] fixed -Wstrict_alliasing warning for GCC --- modules/gpu/src/imgproc.cpp | 8 +++++- modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp | 32 ++++++++++++++-------- modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu | 18 +++++++----- .../gpu/test/nvidia/TestHaarCascadeApplication.cpp | 6 ++-- modules/gpu/test/nvidia/main_nvidia.cpp | 4 ++- 5 files changed, 44 insertions(+), 24 deletions(-) diff --git a/modules/gpu/src/imgproc.cpp b/modules/gpu/src/imgproc.cpp index 31ab44c..fe1ad7b 100644 --- a/modules/gpu/src/imgproc.cpp +++ b/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(value[0]); - Npp32s nVal = *(reinterpret_cast(&val)); + Npp32s nVal = *(reinterpret_cast(&val)); nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), static_cast(src.step), srcsz, dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); break; diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp b/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp index 79872cb..fd0e53e 100644 --- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp +++ b/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) diff --git a/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu b/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu index 6e0f5eb..f754539 100644 --- a/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu +++ b/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 diff --git a/modules/gpu/test/nvidia/TestHaarCascadeApplication.cpp b/modules/gpu/test/nvidia/TestHaarCascadeApplication.cpp index e53f7f3..fe4f5d6 100644 --- a/modules/gpu/test/nvidia/TestHaarCascadeApplication.cpp +++ b/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 h_pixelMask_d(*this->allocatorCPU.get(), this->width, this->height); diff --git a/modules/gpu/test/nvidia/main_nvidia.cpp b/modules/gpu/test/nvidia/main_nvidia.cpp index bb61608..15759e1 100644 --- a/modules/gpu/test/nvidia/main_nvidia.cpp +++ b/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 -- 2.7.4