From 02fb3f0a774313a4008e9a660a05d44b81d26afd Mon Sep 17 00:00:00 2001 From: Vadim Pisarevsky Date: Tue, 17 Dec 2013 20:55:49 +0400 Subject: [PATCH] it finally works!!! --- modules/objdetect/src/cascadedetect.cpp | 81 ++- modules/objdetect/src/cascadedetect.hpp | 14 +- .../objdetect/src/opencl/haarobjectdetect.cl | 639 +++--------------- 3 files changed, 155 insertions(+), 579 deletions(-) diff --git a/modules/objdetect/src/cascadedetect.cpp b/modules/objdetect/src/cascadedetect.cpp index fc43c8c522..2b3e939be2 100644 --- a/modules/objdetect/src/cascadedetect.cpp +++ b/modules/objdetect/src/cascadedetect.cpp @@ -44,6 +44,7 @@ #include "cascadedetect.hpp" #include "opencv2/objdetect/objdetect_c.h" +#include "opencl_kernels.hpp" #if defined (LOG_CASCADE_STATISTIC) struct Logger @@ -491,7 +492,7 @@ bool HaarEvaluator::read(const FileNode& node) features->resize(n); FileNodeIterator it = node.begin(); hasTiltedFeatures = false; - std::vector ff = *features; + std::vector& ff = *features; sumSize0 = Size(); ufbuf.release(); @@ -552,30 +553,37 @@ bool HaarEvaluator::setImage( InputArray _image, Size _origWinSize, Size _sumSiz tofs = (int)((utilted.offset - usum.offset)/sizeof(int)); } else + { integral(_image, usum, noArray(), noArray(), CV_32S); + } + sqrBoxFilter(_image, usqsum, CV_32S, Size(normrect.width, normrect.height), Point(0, 0), false); + /*sqrBoxFilter(_image.getMat(), sqsum, CV_32S, + Size(normrect.width, normrect.height), + Point(0, 0), false); + sqsum.copyTo(usqsum);*/ sumStep = (int)(usum.step/usum.elemSize()); } else { sum0.create(rn*rn_scale, cn, CV_32S); - sqsum0.create(rn, cn, CV_64F); + sqsum0.create(rn, cn, CV_32S); sum = sum0(Rect(0, 0, cols+1, rows+1)); - sqsum = sqsum0(Rect(0, 0, cols+1, rows+1)); + sqsum = sqsum0(Rect(0, 0, cols, rows)); if( hasTiltedFeatures ) { Mat tilted = sum0(Rect(0, _sumSize.height, cols+1, rows+1)); - integral(_image, sum, sqsum, tilted, CV_32S); + integral(_image, sum, noArray(), tilted, CV_32S); tofs = (int)((tilted.data - sum.data)/sizeof(int)); } else - integral(_image, sum, sqsum, noArray(), CV_32S); - /*sqrBoxFilter(_image, sqsum, CV_32S, + integral(_image, sum, noArray(), noArray(), CV_32S); + sqrBoxFilter(_image, sqsum, CV_32S, Size(normrect.width, normrect.height), - Point(0, 0), false);*/ + Point(0, 0), false); sumStep = (int)(sum.step/sum.elemSize()); } @@ -592,7 +600,7 @@ bool HaarEvaluator::setImage( InputArray _image, Size _origWinSize, Size _sumSiz optfeaturesPtr[fi].setOffsets( ff[fi], sumStep, tofs ); } if( _image.isUMat() && (sumSize0 != _sumSize || ufbuf.empty()) ) - copyVectorToUMat(ff, ufbuf); + copyVectorToUMat(*optfeatures, ufbuf); sumSize0 = _sumSize; return true; @@ -608,13 +616,7 @@ bool HaarEvaluator::setWindow( Point pt ) const int* p = &sum.at(pt); int valsum = CALC_SUM_OFS(nofs, p); - - int nqofs[4]; - CV_SUM_OFS( nqofs[0], nqofs[1], nqofs[2], nqofs[3], 0, normrect, (int)(sqsum.step/sizeof(double)) ); - const double* pq = &sqsum.at(pt); - double valsqsum = CALC_SUM_OFS(nqofs, pq); - - //double valsqsum = sqsum.at(pt.y + normrect.y, pt.x + normrect.x); + double valsqsum = sqsum.at(pt.y + normrect.y, pt.x + normrect.x); double nf = (double)normrect.area() * valsqsum - (double)valsum * valsum; if( nf > 0. ) @@ -1131,8 +1133,6 @@ bool CascadeClassifierImpl::detectSingleScale( InputArray _image, Size processin bool CascadeClassifierImpl::ocl_detectSingleScale( InputArray _image, Size processingRectSize, int yStep, double factor, Size sumSize0 ) { - const int MAX_FACES = 10000; - Ptr haar = featureEvaluator.dynamicCast(); if( haar.empty() ) return false; @@ -1141,7 +1141,8 @@ bool CascadeClassifierImpl::ocl_detectSingleScale( InputArray _image, Size proce if( cascadeKernel.empty() ) { - //cascadeKernel.create(") + cascadeKernel.create("runHaarClassifierStump", ocl::objdetect::haarobjectdetect_oclsrc, + format("-D MAX_FACES=%d", MAX_FACES)); if( cascadeKernel.empty() ) return false; } @@ -1152,30 +1153,35 @@ bool CascadeClassifierImpl::ocl_detectSingleScale( InputArray _image, Size proce copyVectorToUMat(data.classifiers, uclassifiers); copyVectorToUMat(data.nodes, unodes); copyVectorToUMat(data.leaves, uleaves); - ufacepos.create(1, MAX_FACES*4 + 1, CV_32S); } std::vector bufs; haar->getUMats(bufs); CV_Assert(bufs.size() == 3); - + + Rect normrect = haar->getNormRect(); + + //processingRectSize = Size(yStep, yStep); size_t globalsize[] = { processingRectSize.width/yStep, processingRectSize.height/yStep }; - return cascadeKernel.args(ocl::KernelArg::ReadOnly(bufs[0]), // sum - ocl::KernelArg::ReadOnly(bufs[1]), // sqsum + cascadeKernel.args(ocl::KernelArg::ReadOnlyNoSize(bufs[0]), // sum + ocl::KernelArg::ReadOnlyNoSize(bufs[1]), // sqsum ocl::KernelArg::PtrReadOnly(bufs[2]), // optfeatures // cascade classifier + (int)data.stages.size(), ocl::KernelArg::PtrReadOnly(ustages), ocl::KernelArg::PtrReadOnly(uclassifiers), ocl::KernelArg::PtrReadOnly(unodes), ocl::KernelArg::PtrReadOnly(uleaves), - ocl::KernelArg::WriteOnly(ufacepos), // positions - ocl::KernelArg::PtrReadOnly(uparams), - processingRectSize.width, - processingRectSize.height, - yStep, (float)factor, MAX_FACES).run(2, globalsize, 0, false); + ocl::KernelArg::PtrWriteOnly(ufacepos), // positions + processingRectSize, + yStep, (float)factor, + normrect, data.origWinSize); + bool ok = cascadeKernel.run(2, globalsize, 0, true); + //CV_Assert(ok); + return ok; } bool CascadeClassifierImpl::isOldFormatCascade() const @@ -1234,12 +1240,13 @@ void CascadeClassifierImpl::detectMultiScaleNoGrouping( InputArray _image, std:: if( maxObjectSize.height == 0 || maxObjectSize.width == 0 ) maxObjectSize = imgsz; - bool use_ocl = false;/*ocl::useOpenCL() && + bool use_ocl = ocl::useOpenCL() && getFeatureType() == FeatureEvaluator::HAAR && !isOldFormatCascade() && + data.isStumpBased && maskGenerator.empty() && !outputRejectLevels && - tryOpenCL;*/ + tryOpenCL; if( !use_ocl ) { @@ -1268,13 +1275,20 @@ void CascadeClassifierImpl::detectMultiScaleNoGrouping( InputArray _image, std:: } Size sumSize0((imgsz.width + SUM_ALIGN) & -SUM_ALIGN, imgsz.height+1); + + if( use_ocl ) + { + ufacepos.create(1, MAX_FACES*4 + 1, CV_32S); + UMat ufacecount(ufacepos, Rect(0,0,1,1)); + ufacecount.setTo(Scalar::all(0)); + } for( double factor = 1; ; factor *= scaleFactor ) { Size originalWindowSize = getOriginalWindowSize(); Size windowSize( cvRound(originalWindowSize.width*factor), cvRound(originalWindowSize.height*factor) ); - Size scaledImageSize( cvRound( grayImage.cols/factor ), cvRound( grayImage.rows/factor ) ); + Size scaledImageSize( cvRound( imgsz.width/factor ), cvRound( imgsz.height/factor ) ); Size processingRectSize( scaledImageSize.width - originalWindowSize.width, scaledImageSize.height - originalWindowSize.height ); @@ -1331,6 +1345,7 @@ void CascadeClassifierImpl::detectMultiScaleNoGrouping( InputArray _image, std:: Mat facepos = ufacepos.getMat(ACCESS_READ); const int* fptr = facepos.ptr(); int i, nfaces = fptr[0]; + printf("nfaces = %d\n", nfaces); for( i = 0; i < nfaces; i++ ) { candidates.push_back(Rect(fptr[i*4+1], fptr[i*4+2], fptr[i*4+3], fptr[i*4+4])); @@ -1439,8 +1454,6 @@ bool CascadeClassifierImpl::Data::read(const FileNode &root) origWinSize.height = (int)root[CC_HEIGHT]; CV_Assert( origWinSize.height > 0 && origWinSize.width > 0 ); - isStumpBased = (int)(root[CC_STAGE_PARAMS][CC_MAX_DEPTH]) == 1 ? true : false; - // load feature params FileNode fn = root[CC_FEATURE_PARAMS]; if( fn.empty() ) @@ -1460,6 +1473,7 @@ bool CascadeClassifierImpl::Data::read(const FileNode &root) nodes.clear(); FileNodeIterator it = fn.begin(), it_end = fn.end(); + isStumpBased = true; for( int si = 0; it != it_end; si++, ++it ) { @@ -1485,6 +1499,9 @@ bool CascadeClassifierImpl::Data::read(const FileNode &root) DTree tree; tree.nodeCount = (int)internalNodes.size()/nodeStep; + if( tree.nodeCount > 1 ) + isStumpBased = false; + classifiers.push_back(tree); nodes.reserve(nodes.size() + tree.nodeCount); diff --git a/modules/objdetect/src/cascadedetect.hpp b/modules/objdetect/src/cascadedetect.hpp index 9841748454..bbe4f083e0 100644 --- a/modules/objdetect/src/cascadedetect.hpp +++ b/modules/objdetect/src/cascadedetect.hpp @@ -63,8 +63,8 @@ protected: double scaleFactor, Size minObjectSize, Size maxObjectSize, bool outputRejectLevels = false ); - enum { BOOST = 0 - }; + enum { MAX_FACES = 10000 }; + enum { BOOST = 0 }; enum { DO_CANNY_PRUNING = CASCADE_DO_CANNY_PRUNING, SCALE_IMAGE = CASCADE_SCALE_IMAGE, FIND_BIGGEST_OBJECT = CASCADE_FIND_BIGGEST_OBJECT, @@ -132,7 +132,7 @@ protected: Ptr maskGenerator; UMat ugrayImage, uimageBuffer; - UMat ufacepos, ustages, uclassifiers, unodes, uleaves, usubsets, uparams; + UMat ufacepos, ustages, uclassifiers, unodes, uleaves, usubsets; ocl::Kernel cascadeKernel; bool tryOpenCL; @@ -327,19 +327,19 @@ inline void HaarEvaluator::OptFeature :: setOffsets( const Feature& _f, int step weight[0] = _f.rect[0].weight; weight[1] = _f.rect[1].weight; weight[2] = _f.rect[2].weight; + + Rect r2 = weight[2] > 0 ? _f.rect[2].r : Rect(0,0,0,0); if (_f.tilted) { CV_TILTED_OFS( ofs[0][0], ofs[0][1], ofs[0][2], ofs[0][3], tofs, _f.rect[0].r, step ); CV_TILTED_OFS( ofs[1][0], ofs[1][1], ofs[1][2], ofs[1][3], tofs, _f.rect[1].r, step ); - if (weight[2]) - CV_TILTED_PTRS( ofs[2][0], ofs[2][1], ofs[2][2], ofs[2][3], tofs, _f.rect[2].r, step ); + CV_TILTED_PTRS( ofs[2][0], ofs[2][1], ofs[2][2], ofs[2][3], tofs, r2, step ); } else { CV_SUM_OFS( ofs[0][0], ofs[0][1], ofs[0][2], ofs[0][3], 0, _f.rect[0].r, step ); CV_SUM_OFS( ofs[1][0], ofs[1][1], ofs[1][2], ofs[1][3], 0, _f.rect[1].r, step ); - if (weight[2]) - CV_SUM_OFS( ofs[2][0], ofs[2][1], ofs[2][2], ofs[2][3], 0, _f.rect[2].r, step ); + CV_SUM_OFS( ofs[2][0], ofs[2][1], ofs[2][2], ofs[2][3], 0, r2, step ); } } diff --git a/modules/objdetect/src/opencl/haarobjectdetect.cl b/modules/objdetect/src/opencl/haarobjectdetect.cl index 980e85dd27..5e46474684 100644 --- a/modules/objdetect/src/opencl/haarobjectdetect.cl +++ b/modules/objdetect/src/opencl/haarobjectdetect.cl @@ -12,6 +12,7 @@ // Nathan, liujun@multicorewareinc.com // Peng Xiao, pengxiao@outlook.com // Erping Pang, erping@multicorewareinc.com +// Vadim Pisarevsky, vadim.pisarevsky@itseez.com // Redistribution and use in source and binary forms, with or without modification, // are permitted provided that the following conditions are met: // @@ -38,559 +39,117 @@ // // -#define CV_HAAR_FEATURE_MAX 3 - -#define calc_sum(rect,offset) (sum[(rect).p0+offset] - sum[(rect).p1+offset] - sum[(rect).p2+offset] + sum[(rect).p3+offset]) -#define calc_sum1(rect,offset,i) (sum[(rect).p0[i]+offset] - sum[(rect).p1[i]+offset] - sum[(rect).p2[i]+offset] + sum[(rect).p3[i]+offset]) - -typedef int sumtype; -typedef float sqsumtype; - -#ifndef STUMP_BASED -#define STUMP_BASED 1 -#endif - -typedef struct __attribute__((aligned (128) )) GpuHidHaarTreeNode +typedef struct __attribute__((aligned(4))) OptFeature { - int p[CV_HAAR_FEATURE_MAX][4] __attribute__((aligned (64))); - float weight[CV_HAAR_FEATURE_MAX]; - float threshold; - float alpha[3] __attribute__((aligned (16))); + int4 ofs[3] __attribute__((aligned (4))); + float4 weight __attribute__((aligned (4))); +} +OptFeature; + +typedef struct __attribute__((aligned(4))) DTreeNode +{ + int featureIdx __attribute__((aligned (4))); + float threshold __attribute__((aligned (4))); // for ordered features only int left __attribute__((aligned (4))); int right __attribute__((aligned (4))); } -GpuHidHaarTreeNode; +DTreeNode; - -//typedef struct __attribute__((aligned (32))) GpuHidHaarClassifier -//{ -// int count __attribute__((aligned (4))); -// GpuHidHaarTreeNode* node __attribute__((aligned (8))); -// float* alpha __attribute__((aligned (8))); -//} -//GpuHidHaarClassifier; - - -typedef struct __attribute__((aligned (64))) GpuHidHaarStageClassifier +typedef struct __attribute__((aligned (4))) DTree { - int count __attribute__((aligned (4))); + int nodeCount __attribute__((aligned (4))); +} +DTree; + +typedef struct __attribute__((aligned (4))) Stage +{ + int first __attribute__((aligned (4))); + int ntrees __attribute__((aligned (4))); float threshold __attribute__((aligned (4))); - int two_rects __attribute__((aligned (4))); - int reserved0 __attribute__((aligned (8))); - int reserved1 __attribute__((aligned (8))); - int reserved2 __attribute__((aligned (8))); - int reserved3 __attribute__((aligned (8))); } -GpuHidHaarStageClassifier; +Stage; +__kernel void runHaarClassifierStump( + __global const int* sum, + int sumstep, int sumoffset, + __global const int* sqsum, + int sqsumstep, int sqsumoffset, + __global const OptFeature* optfeatures, -//typedef struct __attribute__((aligned (64))) GpuHidHaarClassifierCascade -//{ -// int count __attribute__((aligned (4))); -// int is_stump_based __attribute__((aligned (4))); -// int has_tilted_features __attribute__((aligned (4))); -// int is_tree __attribute__((aligned (4))); -// int pq0 __attribute__((aligned (4))); -// int pq1 __attribute__((aligned (4))); -// int pq2 __attribute__((aligned (4))); -// int pq3 __attribute__((aligned (4))); -// int p0 __attribute__((aligned (4))); -// int p1 __attribute__((aligned (4))); -// int p2 __attribute__((aligned (4))); -// int p3 __attribute__((aligned (4))); -// float inv_window_area __attribute__((aligned (4))); -//} GpuHidHaarClassifierCascade; - - -#ifdef PACKED_CLASSIFIER -// this code is scalar, one pixel -> one workitem -__kernel void gpuRunHaarClassifierCascadePacked( - global const GpuHidHaarStageClassifier * stagecascadeptr, - global const int4 * info, - global const GpuHidHaarTreeNode * nodeptr, - global const int * restrict sum, - global const float * restrict sqsum, - volatile global int4 * candidate, - const int pixelstep, - const int loopcount, - const int start_stage, - const int split_stage, - const int end_stage, - const int startnode, - const int splitnode, - const int4 p, - const int4 pq, - const float correction, - global const int* pNodesPK, - global const int4* pWGInfo - ) + int nstages, + __global const Stage* stages, + __global const DTree* trees, + __global const DTreeNode* nodes, + __global const float* leaves, + volatile __global int* facepos, + int2 imgsize, int xyscale, float factor, + int4 normrect, int2 windowsize) { -// this version used information provided for each workgroup -// no empty WG - int gid = (int)get_group_id(0); - int lid_x = (int)get_local_id(0); - int lid_y = (int)get_local_id(1); - int lid = lid_y*LSx+lid_x; - int4 WGInfo = pWGInfo[gid]; - int GroupX = (WGInfo.y >> 16)&0xFFFF; - int GroupY = (WGInfo.y >> 0 )& 0xFFFF; - int Width = (WGInfo.x >> 16)&0xFFFF; - int Height = (WGInfo.x >> 0 )& 0xFFFF; - int ImgOffset = WGInfo.z; - float ScaleFactor = as_float(WGInfo.w); - -#define DATA_SIZE_X (LSx+WND_SIZE_X) -#define DATA_SIZE_Y (LSy+WND_SIZE_Y) -#define DATA_SIZE (DATA_SIZE_X*DATA_SIZE_Y) - - local int SumL[DATA_SIZE]; - - // read input data window into local mem - for(int i = 0; i 0 ? nf : 1.f; + + for( stageIdx = 0; stageIdx < nstages; stageIdx++ ) + { + ntrees = stages[stageIdx].ntrees; + s = 0.f; + for( i = 0; i < ntrees; i++, nodeOfs++, leafOfs += 2 ) + { + node = nodes + nodeOfs; + f = optfeatures + node->featureIdx; + + weight = f->weight; + + ofs = f->ofs[0]; + sval = (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.x; + ofs = f->ofs[1]; + sval += (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.y; + if( weight.z > 0 ) + { + ofs = f->ofs[2]; + sval += (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.z; + } + s += leaves[ sval < node->threshold*nf ? leafOfs : leafOfs + 1 ]; + } + + if( s < stages[stageIdx].threshold ) + break; + } + + if( stageIdx == nstages ) + { + int nfaces = atomic_inc(facepos); + //printf("detected face #d!!!!\n", nfaces); + if( nfaces < MAX_FACES ) + { + volatile __global int* face = facepos + 1 + nfaces*4; + face[0] = convert_int_rte(ix*factor); + face[1] = convert_int_rte(iy*factor); + face[2] = convert_int_rte(windowsize.x*factor); + face[3] = convert_int_rte(windowsize.y*factor); + } } } - barrier(CLK_LOCAL_MEM_FENCE); - - // calc variance_norm_factor for all stages - float variance_norm_factor; - int nodecounter= startnode; - int4 info1 = p; - int4 info2 = pq; - - { - int xl = lid_x; - int yl = lid_y; - int OffsetLocal = yl * DATA_SIZE_X + xl; - int OffsetGlobal = (GroupY+yl)* pixelstep + (GroupX+xl); - - // add shift to get position on scaled image - OffsetGlobal += ImgOffset; - - float mean = - SumL[info1.y*DATA_SIZE_X+info1.x+OffsetLocal] - - SumL[info1.y*DATA_SIZE_X+info1.z+OffsetLocal] - - SumL[info1.w*DATA_SIZE_X+info1.x+OffsetLocal] + - SumL[info1.w*DATA_SIZE_X+info1.z+OffsetLocal]; - float sq = - sqsum[info2.y*pixelstep+info2.x+OffsetGlobal] - - sqsum[info2.y*pixelstep+info2.z+OffsetGlobal] - - sqsum[info2.w*pixelstep+info2.x+OffsetGlobal] + - sqsum[info2.w*pixelstep+info2.z+OffsetGlobal]; - - mean *= correction; - sq *= correction; - - variance_norm_factor = sq - mean * mean; - variance_norm_factor = (variance_norm_factor >=0.f) ? sqrt(variance_norm_factor) : 1.f; - }// end calc variance_norm_factor for all stages - - int result = (1.0f>0.0f); - for(int stageloop = start_stage; (stageloop < end_stage) && result; stageloop++ ) - {// iterate until candidate is exist - float stage_sum = 0.0f; - __global GpuHidHaarStageClassifier* stageinfo = (__global GpuHidHaarStageClassifier*) - ((__global uchar*)stagecascadeptr+stageloop*sizeof(GpuHidHaarStageClassifier)); - int stagecount = stageinfo->count; - float stagethreshold = stageinfo->threshold; - int lcl_off = (lid_y*DATA_SIZE_X)+(lid_x); - for(int nodeloop = 0; nodeloop < stagecount; nodecounter++,nodeloop++ ) - { - // simple macro to extract shorts from int -#define M0(_t) ((_t)&0xFFFF) -#define M1(_t) (((_t)>>16)&0xFFFF) - // load packed node data from global memory (L3) into registers - global const int4* pN = (__global int4*)(pNodesPK+nodecounter*NODE_SIZE); - int4 n0 = pN[0]; - int4 n1 = pN[1]; - int4 n2 = pN[2]; - float nodethreshold = as_float(n2.y) * variance_norm_factor; - // calc sum of intensity pixels according to node information - float classsum = - (SumL[M0(n0.x)+lcl_off] - SumL[M1(n0.x)+lcl_off] - SumL[M0(n0.y)+lcl_off] + SumL[M1(n0.y)+lcl_off]) * as_float(n1.z) + - (SumL[M0(n0.z)+lcl_off] - SumL[M1(n0.z)+lcl_off] - SumL[M0(n0.w)+lcl_off] + SumL[M1(n0.w)+lcl_off]) * as_float(n1.w) + - (SumL[M0(n1.x)+lcl_off] - SumL[M1(n1.x)+lcl_off] - SumL[M0(n1.y)+lcl_off] + SumL[M1(n1.y)+lcl_off]) * as_float(n2.x); - //accumulate stage responce - stage_sum += (classsum >= nodethreshold) ? as_float(n2.w) : as_float(n2.z); - } - result = (stage_sum >= stagethreshold); - }// next stage if needed - - if(result) - {// all stages will be passed and there is a detected face on the tested position - int index = 1+atomic_inc((volatile global int*)candidate); //get index to write global data with face info - if(index>2)<<2; - int readheight = grpszy-1+WINDOWSIZE; - int read_horiz_cnt = readwidth >> 2;//each read int4 - int total_read = mul24(read_horiz_cnt,readheight); - int read_loop = (total_read + lcl_sz - 1) >> 6; - candidate[outputoff+(lcl_id<<2)] = (int4)0; - candidate[outputoff+(lcl_id<<2)+1] = (int4)0; - candidate[outputoff+(lcl_id<<2)+2] = (int4)0; - candidate[outputoff+(lcl_id<<2)+3] = (int4)0; - for(int scalei = 0; scalei > 16; - int totalgrp = scaleinfo1.y & 0xffff; - int imgoff = scaleinfo1.z; - float factor = as_float(scaleinfo1.w); - - __global const int * sum = sum1 + imgoff; - __global const float * sqsum = sqsum1 + imgoff; - for(int grploop=grpidx; grploop=0.f ? sqrt(variance_norm_factor) : 1.f; - - for(int stageloop = start_stage; (stageloop < split_stage) && result; stageloop++ ) - { - float stage_sum = 0.f; - __global GpuHidHaarStageClassifier* stageinfo = (__global GpuHidHaarStageClassifier*) - ((__global uchar*)stagecascadeptr+stageloop*sizeof(GpuHidHaarStageClassifier)); - int stagecount = stageinfo->count; - float stagethreshold = stageinfo->threshold; - for(int nodeloop = 0; nodeloop < stagecount; ) - { - __global GpuHidHaarTreeNode* currentnodeptr = (__global GpuHidHaarTreeNode*) - (((__global uchar*)nodeptr) + nodecounter * sizeof(GpuHidHaarTreeNode)); - - 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])); - float3 alpha3 = *(__global float3*)(&(currentnodeptr->alpha[0])); - - float nodethreshold = w.w * variance_norm_factor; - - info1.x +=lcl_off; - info1.z +=lcl_off; - info2.x +=lcl_off; - info2.z +=lcl_off; - - float classsum = (lcldata[mad24(info1.y,readwidth,info1.x)] - lcldata[mad24(info1.y,readwidth,info1.z)] - - lcldata[mad24(info1.w,readwidth,info1.x)] + lcldata[mad24(info1.w,readwidth,info1.z)]) * w.x; - - classsum += (lcldata[mad24(info2.y,readwidth,info2.x)] - lcldata[mad24(info2.y,readwidth,info2.z)] - - lcldata[mad24(info2.w,readwidth,info2.x)] + lcldata[mad24(info2.w,readwidth,info2.z)]) * w.y; - - info3.x +=lcl_off; - info3.z +=lcl_off; - classsum += (lcldata[mad24(info3.y,readwidth,info3.x)] - lcldata[mad24(info3.y,readwidth,info3.z)] - - lcldata[mad24(info3.w,readwidth,info3.x)] + lcldata[mad24(info3.w,readwidth,info3.z)]) * w.z; - - bool passThres = classsum >= nodethreshold; -#if STUMP_BASED - stage_sum += passThres ? alpha3.y : alpha3.x; - nodecounter++; - nodeloop++; -#else - bool isRootNode = (nodecounter & 1) == 0; - if(isRootNode) - { - if( (passThres && currentnodeptr->right) || - (!passThres && currentnodeptr->left)) - { - nodecounter ++; - } - else - { - stage_sum += alpha3.x; - nodecounter += 2; - nodeloop ++; - } - } - else - { - stage_sum += passThres ? alpha3.z : alpha3.y; - nodecounter ++; - nodeloop ++; - } -#endif - } - - result = (stage_sum >= stagethreshold) ? 1 : 0; - } - if(factor < 2) - { - if(result && lclidx %2 ==0 && lclidy %2 ==0 ) - { - int queueindex = atomic_inc(lclcount); - lcloutindex[queueindex<<1] = (lclidy << 16) | lclidx; - lcloutindex[(queueindex<<1)+1] = as_int((float)variance_norm_factor); - } - } - else - { - if(result) - { - int queueindex = atomic_inc(lclcount); - lcloutindex[queueindex<<1] = (lclidy << 16) | lclidx; - lcloutindex[(queueindex<<1)+1] = as_int((float)variance_norm_factor); - } - } - barrier(CLK_LOCAL_MEM_FENCE); - int queuecount = lclcount[0]; - barrier(CLK_LOCAL_MEM_FENCE); - nodecounter = splitnode; - 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); - __global GpuHidHaarStageClassifier* stageinfo = (__global GpuHidHaarStageClassifier*) - ((__global uchar*)stagecascadeptr+stageloop*sizeof(GpuHidHaarStageClassifier)); - int stagecount = stageinfo->count; - float stagethreshold = stageinfo->threshold; - - int perfscale = queuecount > 4 ? 3 : 2; - int queuecount_loop = (queuecount + (1<> perfscale; - int lcl_compute_win = lcl_sz >> perfscale; - int lcl_compute_win_id = (lcl_id >>(6-perfscale)); - int lcl_loops = (stagecount + lcl_compute_win -1) >> (6-perfscale); - int lcl_compute_id = lcl_id - (lcl_compute_win_id << (6-perfscale)); - for(int queueloop=0; queueloop>16),readwidth,temp_coord & 0xffff); - - if(lcl_compute_win_id < queuecount) - { - int tempnodecounter = lcl_compute_id; - float part_sum = 0.f; - const int stump_factor = STUMP_BASED ? 1 : 2; - int root_offset = 0; - for(int lcl_loop=0; lcl_loopp[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])); - float3 alpha3 = *(__global float3*)(&(currentnodeptr->alpha[0])); - float nodethreshold = w.w * variance_norm_factor; - - info1.x +=queue_pixel; - info1.z +=queue_pixel; - info2.x +=queue_pixel; - info2.z +=queue_pixel; - - float classsum = (lcldata[mad24(info1.y,readwidth,info1.x)] - lcldata[mad24(info1.y,readwidth,info1.z)] - - lcldata[mad24(info1.w,readwidth,info1.x)] + lcldata[mad24(info1.w,readwidth,info1.z)]) * w.x; - - - classsum += (lcldata[mad24(info2.y,readwidth,info2.x)] - lcldata[mad24(info2.y,readwidth,info2.z)] - - lcldata[mad24(info2.w,readwidth,info2.x)] + lcldata[mad24(info2.w,readwidth,info2.z)]) * w.y; - - info3.x +=queue_pixel; - info3.z +=queue_pixel; - classsum += (lcldata[mad24(info3.y,readwidth,info3.x)] - lcldata[mad24(info3.y,readwidth,info3.z)] - - lcldata[mad24(info3.w,readwidth,info3.x)] + lcldata[mad24(info3.w,readwidth,info3.z)]) * w.z; - - bool passThres = classsum >= nodethreshold; -#if STUMP_BASED - part_sum += passThres ? alpha3.y : alpha3.x; - tempnodecounter += lcl_compute_win; - lcl_loop++; -#else - if(root_offset == 0) - { - if( (passThres && currentnodeptr->right) || - (!passThres && currentnodeptr->left)) - { - root_offset = 1; - } - else - { - part_sum += alpha3.x; - tempnodecounter += lcl_compute_win; - lcl_loop++; - } - } - else - { - part_sum += passThres ? alpha3.z : alpha3.y; - tempnodecounter += lcl_compute_win; - lcl_loop++; - root_offset = 0; - } -#endif - }//end for(int lcl_loop=0;lcl_loop= stagethreshold && (lcl_compute_id==0)) - { - int queueindex = atomic_inc(lclcount); - lcloutindex[queueindex<<1] = temp_coord; - lcloutindex[(queueindex<<1)+1] = as_int(variance_norm_factor); - } - lcl_compute_win_id +=(1<0;stageloop++) - - if(lcl_id> 16)); - temp = glboutindex[0]; - int4 candidate_result; - candidate_result.zw = (int2)convert_int_rte(factor*20.f); - candidate_result.x = convert_int_rte(x*factor); - candidate_result.y = convert_int_rte(y*factor); - atomic_inc(glboutindex); - - int i = outputoff+temp+lcl_id; - if(candidate[i].z == 0) - { - candidate[i] = candidate_result; - } - else - { - for(i=i+1;;i++) - { - if(candidate[i].z == 0) - { - candidate[i] = candidate_result; - break; - } - } - } - } - barrier(CLK_LOCAL_MEM_FENCE); - }//end for(int grploop=grpidx;grploop