reverted r8003 (CascadeClassifier_GPU)

diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp b/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp
index a877c9d..79872cb 100644 (file)
--- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp
+++ b/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp
@@ -59,7 +59,6 @@
 #define _ncvhaarobjectdetection_hpp_\r
 \r
 #include <string>\r
-#include <vector_types.h>\r
 #include "NCV.hpp"\r
 \r
 \r
@@ -69,43 +68,41 @@
 //\r
 //==============================================================================\r
 \r
+\r
 struct HaarFeature64\r
 {\r
-    union\r
-    {\r
-        uint2 _ui2;\r
-        struct {NcvRect8u__ _rect; Ncv32f _f;};\r
-    };\r
-\r
+    uint2 _ui2;\r
 \r
 #define HaarFeature64_CreateCheck_MaxRectField                  0xFF\r
 \r
     __host__ NCVStatus setRect(Ncv32u rectX, Ncv32u rectY, Ncv32u rectWidth, Ncv32u rectHeight, Ncv32u /*clsWidth*/, Ncv32u /*clsHeight*/)\r
     {\r
         ncvAssertReturn(rectWidth <= HaarFeature64_CreateCheck_MaxRectField && rectHeight <= HaarFeature64_CreateCheck_MaxRectField, NCV_HAAR_TOO_LARGE_FEATURES);\r
-        _rect = NcvRect8u(static_cast<Ncv8u>(rectX),static_cast<Ncv8u>(rectY),static_cast<Ncv8u>(rectWidth),static_cast<Ncv8u>(rectHeight));\r
-\r
+        ((NcvRect8u*)&(this->_ui2.x))->x = (Ncv8u)rectX;\r
+        ((NcvRect8u*)&(this->_ui2.x))->y = (Ncv8u)rectY;\r
+        ((NcvRect8u*)&(this->_ui2.x))->width = (Ncv8u)rectWidth;\r
+        ((NcvRect8u*)&(this->_ui2.x))->height = (Ncv8u)rectHeight;\r
         return NCV_SUCCESS;\r
     }\r
 \r
     __host__ NCVStatus setWeight(Ncv32f weight)\r
     {\r
-        _f = weight;\r
-\r
+        ((Ncv32f*)&(this->_ui2.y))[0] = weight;\r
         return NCV_SUCCESS;\r
     }\r
 \r
     __device__ __host__ void getRect(Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight)\r
     {\r
-        *rectX = _rect.x;\r
-        *rectY = _rect.y;\r
-        *rectWidth = _rect.width;\r
-        *rectHeight = _rect.height;\r
+        NcvRect8u tmpRect = *(NcvRect8u*)(&this->_ui2.x);\r
+        *rectX = tmpRect.x;\r
+        *rectY = tmpRect.y;\r
+        *rectWidth = tmpRect.width;\r
+        *rectHeight = tmpRect.height;\r
     }\r
 \r
     __device__ __host__ Ncv32f getWeight(void)\r
     {\r
-        return _f;\r
+        return *(Ncv32f*)(&this->_ui2.y);\r
     }\r
 };\r
 \r
@@ -174,27 +171,23 @@ public:
 \r
 struct HaarClassifierNodeDescriptor32\r
 {\r
-union\r
-{\r
     uint1 _ui1;\r
-    Ncv32f _f;\r
-};\r
 \r
     __host__ NCVStatus create(Ncv32f leafValue)\r
     {\r
-        _f = leafValue;\r
+        *(Ncv32f *)&this->_ui1 = leafValue;\r
         return NCV_SUCCESS;\r
     }\r
 \r
     __host__ NCVStatus create(Ncv32u offsetHaarClassifierNode)\r
     {\r
-        _ui1.x = offsetHaarClassifierNode;\r
+        this->_ui1.x = offsetHaarClassifierNode;\r
         return NCV_SUCCESS;\r
     }\r
 \r
     __host__ Ncv32f getLeafValueHost(void)\r
     {\r
-        return _f;\r
+        return *(Ncv32f *)&this->_ui1.x;\r
     }\r
 \r
 #ifdef __CUDACC__\r
@@ -206,67 +199,57 @@ union
 \r
     __device__ __host__ Ncv32u getNextNodeOffset(void)\r
     {\r
-        return _ui1.x;\r
+        return this->_ui1.x;\r
     }\r
 };\r
 \r
 \r
 struct HaarClassifierNode128\r
 {\r
-union\r
-{\r
     uint4 _ui4;\r
-    struct\r
-    {\r
-        HaarFeatureDescriptor32 _f;\r
-        Ncv32f _t;\r
-        HaarClassifierNodeDescriptor32 _nl;\r
-        HaarClassifierNodeDescriptor32 _nr;\r
-    };\r
-};\r
 \r
     __host__ NCVStatus setFeatureDesc(HaarFeatureDescriptor32 f)\r
     {\r
-        _f = f;\r
+        this->_ui4.x = *(Ncv32u *)&f;\r
         return NCV_SUCCESS;\r
     }\r
 \r
     __host__ NCVStatus setThreshold(Ncv32f t)\r
     {\r
-        _t = t;\r
+        this->_ui4.y = *(Ncv32u *)&t;\r
         return NCV_SUCCESS;\r
     }\r
 \r
     __host__ NCVStatus setLeftNodeDesc(HaarClassifierNodeDescriptor32 nl)\r
     {\r
-        _nl = nl;\r
+        this->_ui4.z = *(Ncv32u *)&nl;\r
         return NCV_SUCCESS;\r
     }\r
 \r
     __host__ NCVStatus setRightNodeDesc(HaarClassifierNodeDescriptor32 nr)\r
     {\r
-        _nr = nr;\r
+        this->_ui4.w = *(Ncv32u *)&nr;\r
         return NCV_SUCCESS;\r
     }\r
 \r
     __host__ __device__ HaarFeatureDescriptor32 getFeatureDesc(void)\r
     {\r
-        return _f;\r
+        return *(HaarFeatureDescriptor32 *)&this->_ui4.x;\r
     }\r
 \r
     __host__ __device__ Ncv32f getThreshold(void)\r
     {\r
-        return _t;\r
+        return *(Ncv32f*)&this->_ui4.y;\r
     }\r
 \r
     __host__ __device__ HaarClassifierNodeDescriptor32 getLeftNodeDesc(void)\r
     {\r
-        return _nl;\r
+        return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.z;\r
     }\r
 \r
     __host__ __device__ HaarClassifierNodeDescriptor32 getRightNodeDesc(void)\r
     {\r
-        return _nr;\r
+        return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.w;\r
     }\r
 };\r
 \r
@@ -277,15 +260,11 @@ struct HaarStage64
 #define HaarStage64_Interpret_MaskRootNodeOffset    0xFFFF0000\r
 #define HaarStage64_Interpret_ShiftRootNodeOffset   16\r
 \r
-union\r
-{\r
     uint2 _ui2;\r
-    struct {Ncv32f _t; Ncv32u _root;};\r
-};\r
 \r
     __host__ NCVStatus setStageThreshold(Ncv32f t)\r
     {\r
-        _t = t;\r
+        this->_ui2.x = *(Ncv32u *)&t;\r
         return NCV_SUCCESS;\r
     }\r
 \r
@@ -311,7 +290,7 @@ union
 \r
     __host__ __device__ Ncv32f getStageThreshold(void)\r
     {\r
-        return _t;\r
+        return *(Ncv32f*)&this->_ui2.x;\r
     }\r
 \r
     __host__ __device__ Ncv32u getStartClassifierRootNodeOffset(void)\r
@@ -325,12 +304,14 @@ union
     }\r
 };\r
 \r
+\r
 NCV_CT_ASSERT(sizeof(HaarFeature64) == 8);\r
 NCV_CT_ASSERT(sizeof(HaarFeatureDescriptor32) == 4);\r
 NCV_CT_ASSERT(sizeof(HaarClassifierNodeDescriptor32) == 4);\r
 NCV_CT_ASSERT(sizeof(HaarClassifierNode128) == 16);\r
 NCV_CT_ASSERT(sizeof(HaarStage64) == 8);\r
 \r
+\r
 //==============================================================================\r
 //\r
 // Classifier cascade descriptor\r
@@ -469,4 +450,4 @@ NCV_EXPORTS NCVStatus ncvHaarStoreNVBIN_host(const std::string &filename,
 \r
 \r
 \r
-#endif // _ncvhaarobjectdetection_hpp_\r
+#endif // _ncvhaarobjectdetection_hpp_
\ No newline at end of file

diff --git a/modules/gpu/src/nvidia/core/NCV.hpp b/modules/gpu/src/nvidia/core/NCV.hpp
index ae2a144..bab90f8 100644 (file)
--- a/modules/gpu/src/nvidia/core/NCV.hpp
+++ b/modules/gpu/src/nvidia/core/NCV.hpp
@@ -134,24 +134,15 @@ typedef      unsigned char Ncv8u;
 typedef              float Ncv32f;\r
 typedef             double Ncv64f;\r
 \r
-struct NcvRect8u__\r
+\r
+struct NcvRect8u\r
 {\r
     Ncv8u x;\r
     Ncv8u y;\r
     Ncv8u width;\r
     Ncv8u height;\r
-};\r
-\r
-struct NcvRect8u : NcvRect8u__\r
-{\r
-    __host__ __device__ NcvRect8u() {}\r
-    __host__ __device__ NcvRect8u(Ncv8u x, Ncv8u y, Ncv8u width, Ncv8u height)\r
-    {\r
-        x = x;\r
-        y = y;\r
-        width = width;\r
-        height = height;\r
-    }\r
+    __host__ __device__ NcvRect8u() : x(0), y(0), width(0), height(0) {};\r
+    __host__ __device__ NcvRect8u(Ncv8u x, Ncv8u y, Ncv8u width, Ncv8u height) : x(x), y(y), width(width), height(height) {}\r
 };\r
 \r
 \r
@@ -1029,4 +1020,4 @@ NCV_EXPORTS NCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst, Ncv32u dstStride, N
 \r
 \r
 \r
-#endif // _ncv_hpp_\r
+#endif // _ncv_hpp_
\ No newline at end of file