generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

make oclHaarDetectObjects running on more ocl platforms

Browse Source
This commit is contained in:
yao 2013-01-16 17:13:32 +08:00 committed by Andrey Kamaev
parent b5bd2cde9e
commit 56c1a7fab6
2 changed files with 344 additions and 399 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

283
modules/ocl/src/haar.cpp
View File

@ -63,13 +63,13 @@ using namespace std;
namespace cv namespace cv
{ {
namespace ocl namespace ocl
{ {
///////////////////////////OpenCL kernel strings/////////////////////////// ///////////////////////////OpenCL kernel strings///////////////////////////
extern const char *haarobjectdetect; extern const char *haarobjectdetect;
extern const char *haarobjectdetectbackup; extern const char *haarobjectdetectbackup;
extern const char *haarobjectdetect_scaled2; extern const char *haarobjectdetect_scaled2;
} }
} }
/* these settings affect the quality of detection: change with care */ /* these settings affect the quality of detection: change with care */
@ -883,13 +883,6 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
bool findBiggestObject = (flags & CV_HAAR_FIND_BIGGEST_OBJECT) != 0; bool findBiggestObject = (flags & CV_HAAR_FIND_BIGGEST_OBJECT) != 0;
// bool roughSearch = (flags & CV_HAAR_DO_ROUGH_SEARCH) != 0; // bool roughSearch = (flags & CV_HAAR_DO_ROUGH_SEARCH) != 0;
//the Intel HD Graphics is unsupported
if (gimg.clCxt->impl->devName.find("Intel(R) HD Graphics") != string::npos)
{
cout << " Intel HD GPU device unsupported " << endl;
return NULL;
}
//double t = 0; //double t = 0;
if( maxSize.height == 0 || maxSize.width == 0 ) if( maxSize.height == 0 || maxSize.width == 0 )
{ {
@ -937,7 +930,7 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
if( gimg.cols < minSize.width || gimg.rows < minSize.height ) if( gimg.cols < minSize.width || gimg.rows < minSize.height )
CV_Error(CV_StsError, "Image too small"); CV_Error(CV_StsError, "Image too small");
if( flags & CV_HAAR_SCALE_IMAGE ) if( (flags & CV_HAAR_SCALE_IMAGE) && gimg.clCxt->impl->devName.find("Intel(R) HD Graphics") == string::npos )
{ {
CvSize winSize0 = cascade->orig_window_size; CvSize winSize0 = cascade->orig_window_size;
//float scalefactor = 1.1f; //float scalefactor = 1.1f;
@ -2170,41 +2163,41 @@ CvType haar_type( CV_TYPE_NAME_HAAR, gpuIsHaarClassifier,
namespace cv namespace cv
{ {
HaarClassifierCascade::HaarClassifierCascade() {} HaarClassifierCascade::HaarClassifierCascade() {}
HaarClassifierCascade::HaarClassifierCascade(const String &filename) HaarClassifierCascade::HaarClassifierCascade(const String &filename)
{ {
load(filename); load(filename);
} }
bool HaarClassifierCascade::load(const String &filename) bool HaarClassifierCascade::load(const String &filename)
{ {
cascade = Ptr<CvHaarClassifierCascade>((CvHaarClassifierCascade *)cvLoad(filename.c_str(), 0, 0, 0)); cascade = Ptr<CvHaarClassifierCascade>((CvHaarClassifierCascade *)cvLoad(filename.c_str(), 0, 0, 0));
return (CvHaarClassifierCascade *)cascade != 0; return (CvHaarClassifierCascade *)cascade != 0;
} }
void HaarClassifierCascade::detectMultiScale( const Mat &image, void HaarClassifierCascade::detectMultiScale( const Mat &image,
Vector<Rect> &objects, double scaleFactor, Vector<Rect> &objects, double scaleFactor,
int minNeighbors, int flags, int minNeighbors, int flags,
Size minSize ) Size minSize )
{ {
MemStorage storage(cvCreateMemStorage(0)); MemStorage storage(cvCreateMemStorage(0));
CvMat _image = image; CvMat _image = image;
CvSeq *_objects = gpuHaarDetectObjects( &_image, cascade, storage, scaleFactor, CvSeq *_objects = gpuHaarDetectObjects( &_image, cascade, storage, scaleFactor,
minNeighbors, flags, minSize ); minNeighbors, flags, minSize );
Seq<Rect>(_objects).copyTo(objects); Seq<Rect>(_objects).copyTo(objects);
} }
int HaarClassifierCascade::runAt(Point pt, int startStage, int) const int HaarClassifierCascade::runAt(Point pt, int startStage, int) const
{ {
return gpuRunHaarClassifierCascade(cascade, pt, startStage); return gpuRunHaarClassifierCascade(cascade, pt, startStage);
} }
void HaarClassifierCascade::setImages( const Mat &sum, const Mat &sqsum, void HaarClassifierCascade::setImages( const Mat &sum, const Mat &sqsum,
const Mat &tilted, double scale ) const Mat &tilted, double scale )
{ {
CvMat _sum = sum, _sqsum = sqsum, _tilted = tilted; CvMat _sum = sum, _sqsum = sqsum, _tilted = tilted;
gpuSetImagesForHaarClassifierCascade( cascade, &_sum, &_sqsum, &_tilted, scale ); gpuSetImagesForHaarClassifierCascade( cascade, &_sum, &_sqsum, &_tilted, scale );
} }
} }
#endif #endif
@ -2579,116 +2572,116 @@ CvPoint pt, int start_stage */)
namespace cv namespace cv
{ {
namespace ocl namespace ocl
{
struct gpuHaarDetectObjects_ScaleImage_Invoker
{
gpuHaarDetectObjects_ScaleImage_Invoker( const CvHaarClassifierCascade *_cascade,
int _stripSize, double _factor,
const Mat &_sum1, const Mat &_sqsum1, Mat *_norm1,
Mat *_mask1, Rect _equRect, ConcurrentRectVector &_vec )
{ {
cascade = _cascade;
stripSize = _stripSize;
factor = _factor;
sum1 = _sum1;
sqsum1 = _sqsum1;
norm1 = _norm1;
mask1 = _mask1;
equRect = _equRect;
vec = &_vec;
}
struct gpuHaarDetectObjects_ScaleImage_Invoker void operator()( const BlockedRange &range ) const
{
Size winSize0 = cascade->orig_window_size;
Size winSize(cvRound(winSize0.width * factor), cvRound(winSize0.height * factor));
int y1 = range.begin() * stripSize, y2 = min(range.end() * stripSize, sum1.rows - 1 - winSize0.height);
Size ssz(sum1.cols - 1 - winSize0.width, y2 - y1);
int x, y, ystep = factor > 2 ? 1 : 2;
for( y = y1; y < y2; y += ystep )
for( x = 0; x < ssz.width; x += ystep )
{
if( gpuRunHaarClassifierCascade( /*cascade, cvPoint(x, y), 0*/ ) > 0 )
vec->push_back(Rect(cvRound(x * factor), cvRound(y * factor),
winSize.width, winSize.height));
}
}
const CvHaarClassifierCascade *cascade;
int stripSize;
double factor;
Mat sum1, sqsum1, *norm1, *mask1;
Rect equRect;
ConcurrentRectVector *vec;
};
struct gpuHaarDetectObjects_ScaleCascade_Invoker
{
gpuHaarDetectObjects_ScaleCascade_Invoker( const CvHaarClassifierCascade *_cascade,
Size _winsize, const Range &_xrange, double _ystep,
size_t _sumstep, const int **_p, const int **_pq,
ConcurrentRectVector &_vec )
{
cascade = _cascade;
winsize = _winsize;
xrange = _xrange;
ystep = _ystep;
sumstep = _sumstep;
p = _p;
pq = _pq;
vec = &_vec;
}
void operator()( const BlockedRange &range ) const
{
int iy, startY = range.begin(), endY = range.end();
const int *p0 = p[0], *p1 = p[1], *p2 = p[2], *p3 = p[3];
const int *pq0 = pq[0], *pq1 = pq[1], *pq2 = pq[2], *pq3 = pq[3];
bool doCannyPruning = p0 != 0;
int sstep = (int)(sumstep / sizeof(p0[0]));
for( iy = startY; iy < endY; iy++ )
{ {
gpuHaarDetectObjects_ScaleImage_Invoker( const CvHaarClassifierCascade *_cascade, int ix, y = cvRound(iy * ystep), ixstep = 1;
int _stripSize, double _factor, for( ix = xrange.start; ix < xrange.end; ix += ixstep )
const Mat &_sum1, const Mat &_sqsum1, Mat *_norm1,
Mat *_mask1, Rect _equRect, ConcurrentRectVector &_vec )
{ {
cascade = _cascade; int x = cvRound(ix * ystep); // it should really be ystep, not ixstep
stripSize = _stripSize;
factor = _factor;
sum1 = _sum1;
sqsum1 = _sqsum1;
norm1 = _norm1;
mask1 = _mask1;
equRect = _equRect;
vec = &_vec;
}
void operator()( const BlockedRange &range ) const if( doCannyPruning )
{
Size winSize0 = cascade->orig_window_size;
Size winSize(cvRound(winSize0.width * factor), cvRound(winSize0.height * factor));
int y1 = range.begin() * stripSize, y2 = min(range.end() * stripSize, sum1.rows - 1 - winSize0.height);
Size ssz(sum1.cols - 1 - winSize0.width, y2 - y1);
int x, y, ystep = factor > 2 ? 1 : 2;
for( y = y1; y < y2; y += ystep )
for( x = 0; x < ssz.width; x += ystep )
{
if( gpuRunHaarClassifierCascade( /*cascade, cvPoint(x, y), 0*/ ) > 0 )
vec->push_back(Rect(cvRound(x * factor), cvRound(y * factor),
winSize.width, winSize.height));
}
}
const CvHaarClassifierCascade *cascade;
int stripSize;
double factor;
Mat sum1, sqsum1, *norm1, *mask1;
Rect equRect;
ConcurrentRectVector *vec;
};
struct gpuHaarDetectObjects_ScaleCascade_Invoker
{
gpuHaarDetectObjects_ScaleCascade_Invoker( const CvHaarClassifierCascade *_cascade,
Size _winsize, const Range &_xrange, double _ystep,
size_t _sumstep, const int **_p, const int **_pq,
ConcurrentRectVector &_vec )
{
cascade = _cascade;
winsize = _winsize;
xrange = _xrange;
ystep = _ystep;
sumstep = _sumstep;
p = _p;
pq = _pq;
vec = &_vec;
}
void operator()( const BlockedRange &range ) const
{
int iy, startY = range.begin(), endY = range.end();
const int *p0 = p[0], *p1 = p[1], *p2 = p[2], *p3 = p[3];
const int *pq0 = pq[0], *pq1 = pq[1], *pq2 = pq[2], *pq3 = pq[3];
bool doCannyPruning = p0 != 0;
int sstep = (int)(sumstep / sizeof(p0[0]));
for( iy = startY; iy < endY; iy++ )
{ {
int ix, y = cvRound(iy * ystep), ixstep = 1; int offset = y * sstep + x;
for( ix = xrange.start; ix < xrange.end; ix += ixstep ) int s = p0[offset] - p1[offset] - p2[offset] + p3[offset];
int sq = pq0[offset] - pq1[offset] - pq2[offset] + pq3[offset];
if( s < 100 || sq < 20 )
{ {
int x = cvRound(ix * ystep); // it should really be ystep, not ixstep ixstep = 2;
continue;
if( doCannyPruning )
{
int offset = y * sstep + x;
int s = p0[offset] - p1[offset] - p2[offset] + p3[offset];
int sq = pq0[offset] - pq1[offset] - pq2[offset] + pq3[offset];
if( s < 100 || sq < 20 )
{
ixstep = 2;
continue;
}
}
int result = gpuRunHaarClassifierCascade(/* cascade, cvPoint(x, y), 0 */);
if( result > 0 )
vec->push_back(Rect(x, y, winsize.width, winsize.height));
ixstep = result != 0 ? 1 : 2;
} }
} }
int result = gpuRunHaarClassifierCascade(/* cascade, cvPoint(x, y), 0 */);
if( result > 0 )
vec->push_back(Rect(x, y, winsize.width, winsize.height));
ixstep = result != 0 ? 1 : 2;
} }
}
const CvHaarClassifierCascade *cascade;
double ystep;
size_t sumstep;
Size winsize;
Range xrange;
const int **p;
const int **pq;
ConcurrentRectVector *vec;
};
} }
const CvHaarClassifierCascade *cascade;
double ystep;
size_t sumstep;
Size winsize;
Range xrange;
const int **p;
const int **pq;
ConcurrentRectVector *vec;
};
}
} }
/* /*

460
modules/ocl/src/kernels/haarobjectdetect_scaled2.cl
View File

@ -44,75 +44,75 @@
//M*/ //M*/
// Enter your kernel in this window // Enter your kernel in this window
#pragma OPENCL EXTENSION cl_amd_printf:enable //#pragma OPENCL EXTENSION cl_amd_printf:enable
#define CV_HAAR_FEATURE_MAX 3 #define CV_HAAR_FEATURE_MAX 3
typedef int sumtype; typedef int sumtype;
typedef float sqsumtype; typedef float sqsumtype;
typedef struct __attribute__((aligned (128))) GpuHidHaarFeature typedef struct __attribute__((aligned(128))) GpuHidHaarFeature
{ {
struct __attribute__((aligned (32))) struct __attribute__((aligned(32)))
{ {
int p0 __attribute__((aligned (4))); int p0 __attribute__((aligned(4)));
int p1 __attribute__((aligned (4))); int p1 __attribute__((aligned(4)));
int p2 __attribute__((aligned (4))); int p2 __attribute__((aligned(4)));
int p3 __attribute__((aligned (4))); int p3 __attribute__((aligned(4)));
float weight __attribute__((aligned (4))); float weight __attribute__((aligned(4)));
} }
rect[CV_HAAR_FEATURE_MAX] __attribute__((aligned (32))); rect[CV_HAAR_FEATURE_MAX] __attribute__((aligned(32)));
} }
GpuHidHaarFeature; GpuHidHaarFeature;
typedef struct __attribute__((aligned (128) )) GpuHidHaarTreeNode typedef struct __attribute__((aligned(128))) GpuHidHaarTreeNode
{ {
int p[CV_HAAR_FEATURE_MAX][4] __attribute__((aligned (64))); int p[CV_HAAR_FEATURE_MAX][4] __attribute__((aligned(64)));
float weight[CV_HAAR_FEATURE_MAX] /*__attribute__((aligned (16)))*/; float weight[CV_HAAR_FEATURE_MAX] /*__attribute__((aligned (16)))*/;
float threshold /*__attribute__((aligned (4)))*/; float threshold /*__attribute__((aligned (4)))*/;
float alpha[2] __attribute__((aligned (8))); float alpha[2] __attribute__((aligned(8)));
int left __attribute__((aligned (4))); int left __attribute__((aligned(4)));
int right __attribute__((aligned (4))); int right __attribute__((aligned(4)));
} }
GpuHidHaarTreeNode; GpuHidHaarTreeNode;
typedef struct __attribute__((aligned (32))) GpuHidHaarClassifier typedef struct __attribute__((aligned(32))) GpuHidHaarClassifier
{ {
int count __attribute__((aligned (4))); int count __attribute__((aligned(4)));
GpuHidHaarTreeNode* node __attribute__((aligned (8))); GpuHidHaarTreeNode *node __attribute__((aligned(8)));
float* alpha __attribute__((aligned (8))); float *alpha __attribute__((aligned(8)));
} }
GpuHidHaarClassifier; GpuHidHaarClassifier;
typedef struct __attribute__((aligned (64))) GpuHidHaarStageClassifier typedef struct __attribute__((aligned(64))) GpuHidHaarStageClassifier
{ {
int count __attribute__((aligned (4))); int count __attribute__((aligned(4)));
float threshold __attribute__((aligned (4))); float threshold __attribute__((aligned(4)));
int two_rects __attribute__((aligned (4))); int two_rects __attribute__((aligned(4)));
int reserved0 __attribute__((aligned (8))); int reserved0 __attribute__((aligned(8)));
int reserved1 __attribute__((aligned (8))); int reserved1 __attribute__((aligned(8)));
int reserved2 __attribute__((aligned (8))); int reserved2 __attribute__((aligned(8)));
int reserved3 __attribute__((aligned (8))); int reserved3 __attribute__((aligned(8)));
} }
GpuHidHaarStageClassifier; GpuHidHaarStageClassifier;
typedef struct __attribute__((aligned (64))) GpuHidHaarClassifierCascade typedef struct __attribute__((aligned(64))) GpuHidHaarClassifierCascade
{ {
int count __attribute__((aligned (4))); int count __attribute__((aligned(4)));
int is_stump_based __attribute__((aligned (4))); int is_stump_based __attribute__((aligned(4)));
int has_tilted_features __attribute__((aligned (4))); int has_tilted_features __attribute__((aligned(4)));
int is_tree __attribute__((aligned (4))); int is_tree __attribute__((aligned(4)));
int pq0 __attribute__((aligned (4))); int pq0 __attribute__((aligned(4)));
int pq1 __attribute__((aligned (4))); int pq1 __attribute__((aligned(4)));
int pq2 __attribute__((aligned (4))); int pq2 __attribute__((aligned(4)));
int pq3 __attribute__((aligned (4))); int pq3 __attribute__((aligned(4)));
int p0 __attribute__((aligned (4))); int p0 __attribute__((aligned(4)));
int p1 __attribute__((aligned (4))); int p1 __attribute__((aligned(4)));
int p2 __attribute__((aligned (4))); int p2 __attribute__((aligned(4)));
int p3 __attribute__((aligned (4))); int p3 __attribute__((aligned(4)));
float inv_window_area __attribute__((aligned (4))); float inv_window_area __attribute__((aligned(4)));
}GpuHidHaarClassifierCascade; } GpuHidHaarClassifierCascade;
__kernel void gpuRunHaarClassifierCascade_scaled2( __kernel void gpuRunHaarClassifierCascade_scaled2(
global GpuHidHaarStageClassifier * stagecascadeptr, global GpuHidHaarStageClassifier *stagecascadeptr,
global int4 * info, global int4 *info,
global GpuHidHaarTreeNode * nodeptr, global GpuHidHaarTreeNode *nodeptr,
global const int * restrict sum, global const int *restrict sum,
global const float * restrict sqsum, global const float *restrict sqsum,
global int4 * candidate, global int4 *candidate,
const int step, const int step,
const int loopcount, const int loopcount,
const int start_stage, const int start_stage,
@ -120,215 +120,167 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
const int end_stage, const int end_stage,
const int startnode, const int startnode,
const int splitnode, const int splitnode,
global int4 * p, global int4 *p,
//const int4 * pq, //const int4 * pq,
global float * correction, global float *correction,
const int nodecount) const int nodecount)
{ {
int grpszx = get_local_size(0); int grpszx = get_local_size(0);
int grpszy = get_local_size(1); int grpszy = get_local_size(1);
int grpnumx = get_num_groups(0); int grpnumx = get_num_groups(0);
int grpidx=get_group_id(0); int grpidx = get_group_id(0);
int lclidx = get_local_id(0); int lclidx = get_local_id(0);
int lclidy = get_local_id(1); int lclidy = get_local_id(1);
int lcl_sz = mul24(grpszx,grpszy); int lcl_sz = mul24(grpszx, grpszy);
int lcl_id = mad24(lclidy,grpszx,lclidx); int lcl_id = mad24(lclidy, grpszx, lclidx);
__local int lclshare[1024]; __local int lclshare[1024];
__local int* glboutindex=lclshare+0; __local int *glboutindex = lclshare + 0;
__local int* lclcount=glboutindex+1; __local int *lclcount = glboutindex + 1;
__local int* lcloutindex=lclcount+1; __local int *lcloutindex = lclcount + 1;
__local float* partialsum=(__local float*)(lcloutindex+(lcl_sz<<1)); __local float *partialsum = (__local float *)(lcloutindex + (lcl_sz << 1));
glboutindex[0]=0; glboutindex[0] = 0;
int outputoff = mul24(grpidx,256); int outputoff = mul24(grpidx, 256);
candidate[outputoff+(lcl_id<<2)] = (int4)0; candidate[outputoff + (lcl_id << 2)] = (int4)0;
candidate[outputoff+(lcl_id<<2)+1] = (int4)0; candidate[outputoff + (lcl_id << 2) + 1] = (int4)0;
candidate[outputoff+(lcl_id<<2)+2] = (int4)0; candidate[outputoff + (lcl_id << 2) + 2] = (int4)0;
candidate[outputoff+(lcl_id<<2)+3] = (int4)0; candidate[outputoff + (lcl_id << 2) + 3] = (int4)0;
for(int scalei = 0; scalei <loopcount; scalei++)
{
int4 scaleinfo1;
scaleinfo1 = info[scalei];
int width = (scaleinfo1.x & 0xffff0000) >> 16;
int height = scaleinfo1.x & 0xffff;
int grpnumperline =(scaleinfo1.y & 0xffff0000) >> 16;
int totalgrp = scaleinfo1.y & 0xffff;
float factor = as_float(scaleinfo1.w);
float correction_t=correction[scalei];
int ystep=(int)(max(2.0f,factor)+0.5f);
for(int grploop=get_group_id(0);grploop<totalgrp;grploop+=grpnumx){
int4 cascadeinfo=p[scalei];
int grpidy = grploop / grpnumperline;
int grpidx = grploop - mul24(grpidy, grpnumperline);
int ix = mad24(grpidx,grpszx,lclidx);
int iy = mad24(grpidy,grpszy,lclidy);
int x=ix*ystep;
int y=iy*ystep;
lcloutindex[lcl_id]=0;
lclcount[0]=0;
int result=1,nodecounter;
float mean,variance_norm_factor;
//if((ix < width) && (iy < height))
{
const int p_offset = mad24(y, step, x);
cascadeinfo.x +=p_offset;
cascadeinfo.z +=p_offset;
mean = (sum[mad24(cascadeinfo.y,step,cascadeinfo.x)] - sum[mad24(cascadeinfo.y,step,cascadeinfo.z)] -
sum[mad24(cascadeinfo.w,step,cascadeinfo.x)] + sum[mad24(cascadeinfo.w,step,cascadeinfo.z)])
*correction_t;
variance_norm_factor =sqsum[mad24(cascadeinfo.y,step, cascadeinfo.x)] - sqsum[mad24(cascadeinfo.y, step, cascadeinfo.z)] -
sqsum[mad24(cascadeinfo.w, step, cascadeinfo.x)] + sqsum[mad24(cascadeinfo.w, step, cascadeinfo.z)];
variance_norm_factor = variance_norm_factor * correction_t - mean * mean;
variance_norm_factor = variance_norm_factor >=0.f ? sqrt(variance_norm_factor) : 1.f;
result = 1;
nodecounter = startnode+nodecount*scalei;
for(int stageloop = start_stage; stageloop < split_stage&&result; stageloop++ )
{
float stage_sum = 0.f;
int4 stageinfo = *(global int4*)(stagecascadeptr+stageloop);
float stagethreshold = as_float(stageinfo.y);
for(int nodeloop = 0; nodeloop < stageinfo.x; nodeloop++ )
{
__global GpuHidHaarTreeNode* currentnodeptr = (nodeptr + nodecounter);
int4 info1 = *(__global int4*)(&(currentnodeptr->p[0][0]));
int4 info2 = *(__global int4*)(&(currentnodeptr->p[1][0]));
int4 info3 = *(__global int4*)(&(currentnodeptr->p[2][0]));
float4 w = *(__global float4*)(&(currentnodeptr->weight[0]));
float2 alpha2 = *(__global float2*)(&(currentnodeptr->alpha[0]));
float nodethreshold = w.w * variance_norm_factor;
info1.x +=p_offset;
info1.z +=p_offset;
info2.x +=p_offset;
info2.z +=p_offset;
float classsum = (sum[mad24(info1.y,step,info1.x)] - sum[mad24(info1.y,step,info1.z)] -
sum[mad24(info1.w,step,info1.x)] + sum[mad24(info1.w,step,info1.z)]) * w.x;
classsum += (sum[mad24(info2.y,step,info2.x)] - sum[mad24(info2.y,step,info2.z)] -
sum[mad24(info2.w,step,info2.x)] + sum[mad24(info2.w,step,info2.z)]) * w.y;
info3.x +=p_offset;
info3.z +=p_offset;
classsum += (sum[mad24(info3.y,step,info3.x)] - sum[mad24(info3.y,step,info3.z)] -
sum[mad24(info3.w,step,info3.x)] + sum[mad24(info3.w,step,info3.z)]) * w.z;
stage_sum += classsum >= nodethreshold ? alpha2.y : alpha2.x;
nodecounter++;
}
result=(stage_sum>=stagethreshold);
}
if(result&&(ix<width)&&(iy<height))
{
int queueindex=atomic_inc(lclcount);
lcloutindex[queueindex<<1]=(y<<16)|x;
lcloutindex[(queueindex<<1)+1]=as_int(variance_norm_factor);
}
barrier(CLK_LOCAL_MEM_FENCE);
int queuecount=lclcount[0];
nodecounter=splitnode+nodecount*scalei;
for(int stageloop=split_stage;stageloop<end_stage&&queuecount>0;stageloop++)
{
lclcount[0]=0;
barrier(CLK_LOCAL_MEM_FENCE);
int2 stageinfo=*(global int2*)(stagecascadeptr+stageloop);
float stagethreshold=as_float(stageinfo.y);
int perfscale=queuecount>4?3:2;
int queuecount_loop=(queuecount+(1<<perfscale)-1)>>perfscale;
int lcl_compute_win=lcl_sz>>perfscale;
int lcl_compute_win_id=(lcl_id>>(6-perfscale));
int lcl_loops=(stageinfo.x+lcl_compute_win-1)>>(6-perfscale);
int lcl_compute_id=lcl_id-(lcl_compute_win_id<<(6-perfscale));
for(int queueloop=0;queueloop<queuecount_loop&&lcl_compute_win_id<queuecount;queueloop++)
{
float stage_sum=0.f;
int temp_coord=lcloutindex[lcl_compute_win_id<<1];
float variance_norm_factor=as_float(lcloutindex[(lcl_compute_win_id<<1)+1]);
int queue_offset=mad24(((temp_coord&(int)0xffff0000)>>16),step,temp_coord&0xffff);
int tempnodecounter=lcl_compute_id;
float part_sum=0.f;
for(int lcl_loop=0;lcl_loop<lcl_loops&&tempnodecounter<stageinfo.x;lcl_loop++)
{
__global GpuHidHaarTreeNode* currentnodeptr = (nodeptr + nodecounter + tempnodecounter);
int4 info1 = *(__global int4*)(&(currentnodeptr->p[0][0]));
int4 info2 = *(__global int4*)(&(currentnodeptr->p[1][0]));
int4 info3 = *(__global int4*)(&(currentnodeptr->p[2][0]));
float4 w = *(__global float4*)(&(currentnodeptr->weight[0]));
float2 alpha2 = *(__global float2*)(&(currentnodeptr->alpha[0]));
float nodethreshold = w.w * variance_norm_factor;
info1.x +=queue_offset;
info1.z +=queue_offset;
info2.x +=queue_offset;
info2.z +=queue_offset;
float classsum = (sum[mad24(info1.y,step,info1.x)] - sum[mad24(info1.y,step,info1.z)] -
sum[mad24(info1.w,step,info1.x)] + sum[mad24(info1.w,step,info1.z)]) * w.x;
classsum += (sum[mad24(info2.y,step,info2.x)] - sum[mad24(info2.y,step,info2.z)] -
sum[mad24(info2.w,step,info2.x)] + sum[mad24(info2.w,step,info2.z)]) * w.y;
info3.x +=queue_offset; for (int scalei = 0; scalei < loopcount; scalei++)
info3.z +=queue_offset; {
classsum += (sum[mad24(info3.y,step,info3.x)] - sum[mad24(info3.y,step,info3.z)] - int4 scaleinfo1;
sum[mad24(info3.w,step,info3.x)] + sum[mad24(info3.w,step,info3.z)]) * w.z; scaleinfo1 = info[scalei];
part_sum += classsum >= nodethreshold ? alpha2.y : alpha2.x; int width = (scaleinfo1.x & 0xffff0000) >> 16;
tempnodecounter+=lcl_compute_win; int height = scaleinfo1.x & 0xffff;
} int grpnumperline = (scaleinfo1.y & 0xffff0000) >> 16;
partialsum[lcl_id]=part_sum; int totalgrp = scaleinfo1.y & 0xffff;
barrier(CLK_LOCAL_MEM_FENCE); float factor = as_float(scaleinfo1.w);
for(int i=0;i<lcl_compute_win&&(lcl_compute_id==0);i++) float correction_t = correction[scalei];
{ int ystep = (int)(max(2.0f, factor) + 0.5f);
stage_sum+=partialsum[lcl_id+i];
} for (int grploop = get_group_id(0); grploop < totalgrp; grploop += grpnumx)
if(stage_sum>=stagethreshold&&(lcl_compute_id==0)) {
{ int4 cascadeinfo = p[scalei];
int queueindex=atomic_inc(lclcount); int grpidy = grploop / grpnumperline;
lcloutindex[queueindex<<1]=temp_coord; int grpidx = grploop - mul24(grpidy, grpnumperline);
lcloutindex[(queueindex<<1)+1]=as_int(variance_norm_factor); int ix = mad24(grpidx, grpszx, lclidx);
} int iy = mad24(grpidy, grpszy, lclidy);
lcl_compute_win_id+=(1<<perfscale); int x = ix * ystep;
barrier(CLK_LOCAL_MEM_FENCE); int y = iy * ystep;
} lcloutindex[lcl_id] = 0;
queuecount=lclcount[0]; lclcount[0] = 0;
nodecounter+=stageinfo.x; int result = 1, nodecounter;
} float mean, variance_norm_factor;
if(lcl_id<queuecount) //if((ix < width) && (iy < height))
{ {
int temp=lcloutindex[lcl_id<<1]; const int p_offset = mad24(y, step, x);
int x=temp&0xffff; cascadeinfo.x += p_offset;
int y=(temp&(int)0xffff0000)>>16; cascadeinfo.z += p_offset;
temp=glboutindex[0]; mean = (sum[mad24(cascadeinfo.y, step, cascadeinfo.x)] - sum[mad24(cascadeinfo.y, step, cascadeinfo.z)] -
int4 candidate_result; sum[mad24(cascadeinfo.w, step, cascadeinfo.x)] + sum[mad24(cascadeinfo.w, step, cascadeinfo.z)])
candidate_result.zw=(int2)convert_int_rtn(factor*20.f); * correction_t;
candidate_result.x=x; variance_norm_factor = sqsum[mad24(cascadeinfo.y, step, cascadeinfo.x)] - sqsum[mad24(cascadeinfo.y, step, cascadeinfo.z)] -
candidate_result.y=y; sqsum[mad24(cascadeinfo.w, step, cascadeinfo.x)] + sqsum[mad24(cascadeinfo.w, step, cascadeinfo.z)];
atomic_inc(glboutindex); variance_norm_factor = variance_norm_factor * correction_t - mean * mean;
candidate[outputoff+temp+lcl_id]=candidate_result; variance_norm_factor = variance_norm_factor >= 0.f ? sqrt(variance_norm_factor) : 1.f;
} result = 1;
barrier(CLK_LOCAL_MEM_FENCE); nodecounter = startnode + nodecount * scalei;
}
for (int stageloop = start_stage; stageloop < end_stage && result; stageloop++)
{
float stage_sum = 0.f;
int4 stageinfo = *(global int4 *)(stagecascadeptr + stageloop);
float stagethreshold = as_float(stageinfo.y);
for (int nodeloop = 0; nodeloop < stageinfo.x; nodeloop++)
{
__global GpuHidHaarTreeNode *currentnodeptr = (nodeptr + nodecounter);
int4 info1 = *(__global int4 *)(&(currentnodeptr->p[0][0]));
int4 info2 = *(__global int4 *)(&(currentnodeptr->p[1][0]));
int4 info3 = *(__global int4 *)(&(currentnodeptr->p[2][0]));
float4 w = *(__global float4 *)(&(currentnodeptr->weight[0]));
float2 alpha2 = *(__global float2 *)(&(currentnodeptr->alpha[0]));
float nodethreshold = w.w * variance_norm_factor;
info1.x += p_offset;
info1.z += p_offset;
info2.x += p_offset;
info2.z += p_offset;
float classsum = (sum[mad24(info1.y, step, info1.x)] - sum[mad24(info1.y, step, info1.z)] -
sum[mad24(info1.w, step, info1.x)] + sum[mad24(info1.w, step, info1.z)]) * w.x;
classsum += (sum[mad24(info2.y, step, info2.x)] - sum[mad24(info2.y, step, info2.z)] -
sum[mad24(info2.w, step, info2.x)] + sum[mad24(info2.w, step, info2.z)]) * w.y;
info3.x += p_offset;
info3.z += p_offset;
classsum += (sum[mad24(info3.y, step, info3.x)] - sum[mad24(info3.y, step, info3.z)] -
sum[mad24(info3.w, step, info3.x)] + sum[mad24(info3.w, step, info3.z)]) * w.z;
stage_sum += classsum >= nodethreshold ? alpha2.y : alpha2.x;
nodecounter++;
}
result = (stage_sum >= stagethreshold);
}
if (result && (ix < width) && (iy < height))
{
int queueindex = atomic_inc(lclcount);
lcloutindex[queueindex << 1] = (y << 16) | x;
lcloutindex[(queueindex << 1) + 1] = as_int(variance_norm_factor);
}
barrier(CLK_LOCAL_MEM_FENCE);
int queuecount = lclcount[0];
nodecounter = splitnode + nodecount * scalei;
if (lcl_id < queuecount)
{
int temp = lcloutindex[lcl_id << 1];
int x = temp & 0xffff;
int y = (temp & (int)0xffff0000) >> 16;
temp = glboutindex[0];
int4 candidate_result;
candidate_result.zw = (int2)convert_int_rtn(factor * 20.f);
candidate_result.x = x;
candidate_result.y = y;
atomic_inc(glboutindex);
candidate[outputoff + temp + lcl_id] = candidate_result;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
}
} }
}
} }
__kernel void gpuscaleclassifier(global GpuHidHaarTreeNode * orinode, global GpuHidHaarTreeNode * newnode,float scale,float weight_scale,int nodenum) __kernel void gpuscaleclassifier(global GpuHidHaarTreeNode *orinode, global GpuHidHaarTreeNode *newnode, float scale, float weight_scale, int nodenum)
{ {
int counter=get_global_id(0); int counter = get_global_id(0);
int tr_x[3],tr_y[3],tr_h[3],tr_w[3],i=0; int tr_x[3], tr_y[3], tr_h[3], tr_w[3], i = 0;
GpuHidHaarTreeNode t1 = *(orinode + counter); GpuHidHaarTreeNode t1 = *(orinode + counter);
#pragma unroll #pragma unroll
for(i=0;i<3;i++){
tr_x[i]=(int)(t1.p[i][0]*scale+0.5f); for (i = 0; i < 3; i++)
tr_y[i]=(int)(t1.p[i][1]*scale+0.5f); {
tr_w[i]=(int)(t1.p[i][2]*scale+0.5f); tr_x[i] = (int)(t1.p[i][0] * scale + 0.5f);
tr_h[i]=(int)(t1.p[i][3]*scale+0.5f); tr_y[i] = (int)(t1.p[i][1] * scale + 0.5f);
} tr_w[i] = (int)(t1.p[i][2] * scale + 0.5f);
t1.weight[0]=t1.p[2][0]?-(t1.weight[1]*tr_h[1]*tr_w[1]+t1.weight[2]*tr_h[2]*tr_w[2])/(tr_h[0]*tr_w[0]):-t1.weight[1]*tr_h[1]*tr_w[1]/(tr_h[0]*tr_w[0]); tr_h[i] = (int)(t1.p[i][3] * scale + 0.5f);
counter+=nodenum; }
#pragma unroll
for(i=0;i<3;i++) t1.weight[0] = t1.p[2][0] ? -(t1.weight[1] * tr_h[1] * tr_w[1] + t1.weight[2] * tr_h[2] * tr_w[2]) / (tr_h[0] * tr_w[0]) : -t1.weight[1] * tr_h[1] * tr_w[1] / (tr_h[0] * tr_w[0]);
{ counter += nodenum;
newnode[counter].p[i][0]=tr_x[i]; #pragma unroll
newnode[counter].p[i][1]=tr_y[i];
newnode[counter].p[i][2]=tr_x[i]+tr_w[i]; for (i = 0; i < 3; i++)
newnode[counter].p[i][3]=tr_y[i]+tr_h[i]; {
newnode[counter].weight[i]=t1.weight[i]*weight_scale; newnode[counter].p[i][0] = tr_x[i];
} newnode[counter].p[i][1] = tr_y[i];
newnode[counter].left=t1.left; newnode[counter].p[i][2] = tr_x[i] + tr_w[i];
newnode[counter].right=t1.right; newnode[counter].p[i][3] = tr_y[i] + tr_h[i];
newnode[counter].threshold=t1.threshold; newnode[counter].weight[i] = t1.weight[i] * weight_scale;
newnode[counter].alpha[0]=t1.alpha[0]; }
newnode[counter].alpha[1]=t1.alpha[1];
newnode[counter].left = t1.left;
newnode[counter].right = t1.right;
newnode[counter].threshold = t1.threshold;
newnode[counter].alpha[0] = t1.alpha[0];
newnode[counter].alpha[1] = t1.alpha[1];
} }