From 9e868272e917eeb79fd1eceac85e552a0b9272cc Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Sat, 28 Apr 2012 08:56:17 +0000 Subject: [PATCH] reverted r8003 (CascadeClassifier_GPU didn't work) --- modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp | 79 +++++++++-------------- modules/gpu/src/nvidia/core/NCV.hpp | 19 ++---- 2 files changed, 35 insertions(+), 63 deletions(-) diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp b/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp index eddc390..79872cb 100644 --- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp +++ b/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp @@ -59,7 +59,6 @@ #define _ncvhaarobjectdetection_hpp_ #include -#include #include "NCV.hpp" @@ -69,43 +68,41 @@ // //============================================================================== + struct HaarFeature64 { - union - { - uint2 _ui2; - struct {NcvRect8u__ _rect; Ncv32f _f;}; - }; - + uint2 _ui2; #define HaarFeature64_CreateCheck_MaxRectField 0xFF __host__ NCVStatus setRect(Ncv32u rectX, Ncv32u rectY, Ncv32u rectWidth, Ncv32u rectHeight, Ncv32u /*clsWidth*/, Ncv32u /*clsHeight*/) { ncvAssertReturn(rectWidth <= HaarFeature64_CreateCheck_MaxRectField && rectHeight <= HaarFeature64_CreateCheck_MaxRectField, NCV_HAAR_TOO_LARGE_FEATURES); - _rect = NcvRect8u(rectX,rectY,rectWidth,rectHeight); - + ((NcvRect8u*)&(this->_ui2.x))->x = (Ncv8u)rectX; + ((NcvRect8u*)&(this->_ui2.x))->y = (Ncv8u)rectY; + ((NcvRect8u*)&(this->_ui2.x))->width = (Ncv8u)rectWidth; + ((NcvRect8u*)&(this->_ui2.x))->height = (Ncv8u)rectHeight; return NCV_SUCCESS; } __host__ NCVStatus setWeight(Ncv32f weight) { - _f = weight; - + ((Ncv32f*)&(this->_ui2.y))[0] = weight; return NCV_SUCCESS; } __device__ __host__ void getRect(Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight) { - *rectX = _rect.x; - *rectY = _rect.y; - *rectWidth = _rect.width; - *rectHeight = _rect.height; + NcvRect8u tmpRect = *(NcvRect8u*)(&this->_ui2.x); + *rectX = tmpRect.x; + *rectY = tmpRect.y; + *rectWidth = tmpRect.width; + *rectHeight = tmpRect.height; } __device__ __host__ Ncv32f getWeight(void) { - return _f; + return *(Ncv32f*)(&this->_ui2.y); } }; @@ -174,27 +171,23 @@ public: struct HaarClassifierNodeDescriptor32 { -union -{ uint1 _ui1; - Ncv32f _f; -}; __host__ NCVStatus create(Ncv32f leafValue) { - _f = leafValue; + *(Ncv32f *)&this->_ui1 = leafValue; return NCV_SUCCESS; } __host__ NCVStatus create(Ncv32u offsetHaarClassifierNode) { - _ui1.x = offsetHaarClassifierNode; + this->_ui1.x = offsetHaarClassifierNode; return NCV_SUCCESS; } __host__ Ncv32f getLeafValueHost(void) { - return _f; + return *(Ncv32f *)&this->_ui1.x; } #ifdef __CUDACC__ @@ -206,67 +199,57 @@ union __device__ __host__ Ncv32u getNextNodeOffset(void) { - return _ui1.x; + return this->_ui1.x; } }; struct HaarClassifierNode128 { -union -{ uint4 _ui4; - struct - { - HaarFeatureDescriptor32 _f; - Ncv32f _t; - HaarClassifierNodeDescriptor32 _nl; - HaarClassifierNodeDescriptor32 _nr; - }; -}; __host__ NCVStatus setFeatureDesc(HaarFeatureDescriptor32 f) { - _f = f; + this->_ui4.x = *(Ncv32u *)&f; return NCV_SUCCESS; } __host__ NCVStatus setThreshold(Ncv32f t) { - _t = t; + this->_ui4.y = *(Ncv32u *)&t; return NCV_SUCCESS; } __host__ NCVStatus setLeftNodeDesc(HaarClassifierNodeDescriptor32 nl) { - _nl = nl; + this->_ui4.z = *(Ncv32u *)&nl; return NCV_SUCCESS; } __host__ NCVStatus setRightNodeDesc(HaarClassifierNodeDescriptor32 nr) { - _nr = nr; + this->_ui4.w = *(Ncv32u *)&nr; return NCV_SUCCESS; } __host__ __device__ HaarFeatureDescriptor32 getFeatureDesc(void) { - return _f; + return *(HaarFeatureDescriptor32 *)&this->_ui4.x; } __host__ __device__ Ncv32f getThreshold(void) { - return _t; + return *(Ncv32f*)&this->_ui4.y; } __host__ __device__ HaarClassifierNodeDescriptor32 getLeftNodeDesc(void) { - return _nl; + return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.z; } __host__ __device__ HaarClassifierNodeDescriptor32 getRightNodeDesc(void) { - return _nr; + return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.w; } }; @@ -277,15 +260,11 @@ struct HaarStage64 #define HaarStage64_Interpret_MaskRootNodeOffset 0xFFFF0000 #define HaarStage64_Interpret_ShiftRootNodeOffset 16 -union -{ uint2 _ui2; - struct {Ncv32f _t; Ncv32u _root;}; -}; __host__ NCVStatus setStageThreshold(Ncv32f t) { - _t = t; + this->_ui2.x = *(Ncv32u *)&t; return NCV_SUCCESS; } @@ -311,7 +290,7 @@ union __host__ __device__ Ncv32f getStageThreshold(void) { - return _t; + return *(Ncv32f*)&this->_ui2.x; } __host__ __device__ Ncv32u getStartClassifierRootNodeOffset(void) @@ -325,12 +304,14 @@ union } }; + NCV_CT_ASSERT(sizeof(HaarFeature64) == 8); NCV_CT_ASSERT(sizeof(HaarFeatureDescriptor32) == 4); NCV_CT_ASSERT(sizeof(HaarClassifierNodeDescriptor32) == 4); NCV_CT_ASSERT(sizeof(HaarClassifierNode128) == 16); NCV_CT_ASSERT(sizeof(HaarStage64) == 8); + //============================================================================== // // Classifier cascade descriptor @@ -469,4 +450,4 @@ NCV_EXPORTS NCVStatus ncvHaarStoreNVBIN_host(const std::string &filename, -#endif // _ncvhaarobjectdetection_hpp_ +#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 --- 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; typedef double Ncv64f; -struct NcvRect8u__ + +struct NcvRect8u { Ncv8u x; Ncv8u y; Ncv8u width; Ncv8u height; -}; - -struct NcvRect8u : NcvRect8u__ -{ - __host__ __device__ NcvRect8u() {} - __host__ __device__ NcvRect8u(Ncv8u x, Ncv8u y, Ncv8u width, Ncv8u height) - { - x = x; - y = y; - width = width; - height = height; - } + __host__ __device__ NcvRect8u() : x(0), y(0), width(0), height(0) {}; + __host__ __device__ NcvRect8u(Ncv8u x, Ncv8u y, Ncv8u width, Ncv8u height) : x(x), y(y), width(width), height(height) {} }; @@ -1029,4 +1020,4 @@ NCV_EXPORTS NCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst, Ncv32u dstStride, N -#endif // _ncv_hpp_ +#endif // _ncv_hpp_ \ No newline at end of file -- 2.7.4