mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-25 11:40:44 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

reverted r8003 (CascadeClassifier_GPU)

Browse Source
This commit is contained in:
Vladislav Vinogradov 2012-04-28 08:41:10 +00:00
parent 76dc827648
commit 917063b7f7
2 changed files with 35 additions and 63 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

79
modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp
View File

@ -59,7 +59,6 @@
#define _ncvhaarobjectdetection_hpp_
#include <string>
#include <vector_types.h>
#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(static_cast<Ncv8u>(rectX),static_cast<Ncv8u>(rectY),static_cast<Ncv8u>(rectWidth),static_cast<Ncv8u>(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);
}
};
@ -173,28 +170,24 @@ 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_

19
modules/gpu/src/nvidia/core/NCV.hpp
View File

@ -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_