diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp index 45681ba64cb649cac5a061e7d52d4d1cc426ca47..35653a0ffc0474b99ef0786156041544be8b2c84 100644 --- a/modules/gpu/src/cascadeclassifier.cpp +++ b/modules/gpu/src/cascadeclassifier.cpp @@ -336,24 +336,24 @@ NCVStatus loadFromXML(const std::string &filename, haar.NumClassifierTotalNodes = 0; haar.NumFeatures = 0; haar.ClassifierSize.width = 0; - haar.ClassifierSize.height = 0; + haar.ClassifierSize.height = 0; haar.bHasStumpsOnly = true; haar.bNeedsTiltedII = false; Ncv32u curMaxTreeDepth; - std::vector xmlFileCont; + std::vector xmlFileCont; std::vector h_TmpClassifierNotRootNodes; haarStages.resize(0); haarClassifierNodes.resize(0); - haarFeatures.resize(0); + haarFeatures.resize(0); Ptr oldCascade = (CvHaarClassifierCascade*)cvLoad(filename.c_str(), 0, 0, 0); if (oldCascade.empty()) return NCV_HAAR_XML_LOADING_EXCEPTION; - - /////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - + +/////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// + haar.ClassifierSize.width = oldCascade->orig_window_size.width; haar.ClassifierSize.height = oldCascade->orig_window_size.height; @@ -366,53 +366,58 @@ NCVStatus loadFromXML(const std::string &filename, curStage.setStageThreshold(oldCascade->stage_classifier[s].threshold); int treesCount = oldCascade->stage_classifier[s].count; - for(int t = 0; t < treesCount; ++t) // bytrees - { + for(int t = 0; t < treesCount; ++t) // by trees + { Ncv32u nodeId = 0; CvHaarClassifier* tree = &oldCascade->stage_classifier[s].classifier[t]; int nodesCount = tree->count; - for(int n = 0; n < nodesCount; ++n) //by features - { + for(int n = 0; n < nodesCount; ++n) //by features + { CvHaarFeature* feature = &tree->haar_feature[n]; - HaarClassifierNode128 curNode; + HaarClassifierNode128 curNode; curNode.setThreshold(tree->threshold[n]); - + + NcvBool bIsLeftNodeLeaf = false; + NcvBool bIsRightNodeLeaf = false; + HaarClassifierNodeDescriptor32 nodeLeft; if ( tree->left[n] <= 0 ) { Ncv32f leftVal = tree->alpha[-tree->left[n]]; ncvStat = nodeLeft.create(leftVal); - ncvAssertReturn(ncvStat == NCV_SUCCESS, ncvStat); + ncvAssertReturn(ncvStat == NCV_SUCCESS, ncvStat); + bIsLeftNodeLeaf = true; } else { - Ncv32u leftNodeOffset = tree->left[n]; + Ncv32u leftNodeOffset = tree->left[n]; nodeLeft.create((Ncv32u)(h_TmpClassifierNotRootNodes.size() + leftNodeOffset - 1)); haar.bHasStumpsOnly = false; } curNode.setLeftNodeDesc(nodeLeft); - + HaarClassifierNodeDescriptor32 nodeRight; if ( tree->right[n] <= 0 ) - { - Ncv32f rightVal = tree->alpha[-tree->right[n]]; + { + Ncv32f rightVal = tree->alpha[-tree->right[n]]; ncvStat = nodeRight.create(rightVal); ncvAssertReturn(ncvStat == NCV_SUCCESS, ncvStat); + bIsRightNodeLeaf = true; } else - { - Ncv32u rightNodeOffset = tree->right[n]; + { + Ncv32u rightNodeOffset = tree->right[n]; nodeRight.create((Ncv32u)(h_TmpClassifierNotRootNodes.size() + rightNodeOffset - 1)); haar.bHasStumpsOnly = false; } - curNode.setRightNodeDesc(nodeRight); + curNode.setRightNodeDesc(nodeRight); Ncv32u tiltedVal = feature->tilted; - haar.bNeedsTiltedII = (tiltedVal != 0); + haar.bNeedsTiltedII = (tiltedVal != 0); - Ncv32u featureId = 0; + Ncv32u featureId = 0; for(int l = 0; l < CV_HAAR_FEATURE_MAX; ++l) //by rects { Ncv32u rectX = feature->rect[l].r.x; @@ -435,7 +440,8 @@ NCVStatus loadFromXML(const std::string &filename, } HaarFeatureDescriptor32 tmpFeatureDesc; - ncvStat = tmpFeatureDesc.create(haar.bNeedsTiltedII, featureId, haarFeatures.size() - featureId); + ncvStat = tmpFeatureDesc.create(haar.bNeedsTiltedII, bIsLeftNodeLeaf, bIsRightNodeLeaf, + featureId, haarFeatures.size() - featureId); ncvAssertReturn(NCV_SUCCESS == ncvStat, ncvStat); curNode.setFeatureDesc(tmpFeatureDesc); @@ -453,14 +459,14 @@ NCVStatus loadFromXML(const std::string &filename, } nodeId++; - } + } } curStage.setNumClassifierRootNodes(treesCount); - haarStages.push_back(curStage); + haarStages.push_back(curStage); } -/////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// +/////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// //fill in cascade stats haar.NumStages = haarStages.size(); @@ -472,8 +478,10 @@ NCVStatus loadFromXML(const std::string &filename, Ncv32u offsetRoot = haarClassifierNodes.size(); for (Ncv32u i=0; i &h_integralImage, while (bMoreNodesToTraverse) { HaarClassifierNode128 curNode = h_HaarNodes.ptr()[curNodeOffset]; - Ncv32u curNodeFeaturesNum = curNode.getFeatureDesc().getNumFeatures(); - Ncv32u curNodeFeaturesOffs = curNode.getFeatureDesc().getFeaturesOffset(); + HaarFeatureDescriptor32 curFeatDesc = curNode.getFeatureDesc(); + Ncv32u curNodeFeaturesNum = curFeatDesc.getNumFeatures(); + Ncv32u curNodeFeaturesOffs = curFeatDesc.getFeaturesOffset(); Ncv32f curNodeVal = 0.f; for (Ncv32u iRect=0; iRect &h_integralImage, HaarClassifierNodeDescriptor32 nodeLeft = curNode.getLeftNodeDesc(); HaarClassifierNodeDescriptor32 nodeRight = curNode.getRightNodeDesc(); Ncv32f nodeThreshold = curNode.getThreshold(); + HaarClassifierNodeDescriptor32 nextNodeDescriptor; + NcvBool nextNodeIsLeaf; if (curNodeVal < scaleAreaPixels * h_weights.ptr()[i * h_weights.stride() + j] * nodeThreshold) { nextNodeDescriptor = nodeLeft; + nextNodeIsLeaf = curFeatDesc.isLeftNodeLeaf(); } else { nextNodeDescriptor = nodeRight; + nextNodeIsLeaf = curFeatDesc.isRightNodeLeaf(); } - NcvBool tmpIsLeaf = nextNodeDescriptor.isLeaf(); - if (tmpIsLeaf) + if (nextNodeIsLeaf) { Ncv32f tmpLeafValue = nextNodeDescriptor.getLeafValueHost(); curStageSum += tmpLeafValue; diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp b/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp index bb463d8f4c7c37ca52f6112c8bdb5bc640ae7142..5abe2548584a0abfdf9885498ceb06e43b809853 100644 --- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp +++ b/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp @@ -112,7 +112,9 @@ struct HaarFeatureDescriptor32 private: #define HaarFeatureDescriptor32_Interpret_MaskFlagTilted 0x80000000 -#define HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures 0x7F +#define HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf 0x40000000 +#define HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf 0x20000000 +#define HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures 0x1F #define HaarFeatureDescriptor32_NumFeatures_Shift 24 #define HaarFeatureDescriptor32_CreateCheck_MaxFeatureOffset 0x00FFFFFF @@ -120,7 +122,8 @@ private: public: - __host__ NCVStatus create(NcvBool bTilted, Ncv32u numFeatures, Ncv32u offsetFeatures) + __host__ NCVStatus create(NcvBool bTilted, NcvBool bLeftLeaf, NcvBool bRightLeaf, + Ncv32u numFeatures, Ncv32u offsetFeatures) { if (numFeatures > HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures) { @@ -132,6 +135,8 @@ public: } this->desc = 0; this->desc |= (bTilted ? HaarFeatureDescriptor32_Interpret_MaskFlagTilted : 0); + this->desc |= (bLeftLeaf ? HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf : 0); + this->desc |= (bRightLeaf ? HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf : 0); this->desc |= (numFeatures << HaarFeatureDescriptor32_NumFeatures_Shift); this->desc |= offsetFeatures; return NCV_SUCCESS; @@ -142,9 +147,19 @@ public: return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagTilted) != 0; } + __device__ __host__ NcvBool isLeftNodeLeaf(void) + { + return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf) != 0; + } + + __device__ __host__ NcvBool isRightNodeLeaf(void) + { + return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf) != 0; + } + __device__ __host__ Ncv32u getNumFeatures(void) { - return (this->desc & ~HaarFeatureDescriptor32_Interpret_MaskFlagTilted) >> HaarFeatureDescriptor32_NumFeatures_Shift; + return (this->desc >> HaarFeatureDescriptor32_NumFeatures_Shift) & HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures; } __device__ __host__ Ncv32u getFeaturesOffset(void) @@ -158,34 +173,18 @@ struct HaarClassifierNodeDescriptor32 { uint1 _ui1; -#define HaarClassifierNodeDescriptor32_Interpret_MaskSwitch (1 << 30) - __host__ NCVStatus create(Ncv32f leafValue) { - if ((*(Ncv32u *)&leafValue) & HaarClassifierNodeDescriptor32_Interpret_MaskSwitch) - { - return NCV_HAAR_XML_LOADING_EXCEPTION; - } *(Ncv32f *)&this->_ui1 = leafValue; return NCV_SUCCESS; } __host__ NCVStatus create(Ncv32u offsetHaarClassifierNode) { - if (offsetHaarClassifierNode >= HaarClassifierNodeDescriptor32_Interpret_MaskSwitch) - { - return NCV_HAAR_XML_LOADING_EXCEPTION; - } this->_ui1.x = offsetHaarClassifierNode; - this->_ui1.x |= HaarClassifierNodeDescriptor32_Interpret_MaskSwitch; return NCV_SUCCESS; } - __device__ __host__ NcvBool isLeaf(void) - { - return !(this->_ui1.x & HaarClassifierNodeDescriptor32_Interpret_MaskSwitch); - } - __host__ Ncv32f getLeafValueHost(void) { return *(Ncv32f *)&this->_ui1.x; @@ -200,7 +199,7 @@ struct HaarClassifierNodeDescriptor32 __device__ __host__ Ncv32u getNextNodeOffset(void) { - return (this->_ui1.x & ~HaarClassifierNodeDescriptor32_Interpret_MaskSwitch); + return this->_ui1.x; } };