|
|
|
|
@ -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(4))) OptFeature |
|
|
|
|
{ |
|
|
|
|
int4 ofs[3] __attribute__((aligned (4))); |
|
|
|
|
float4 weight __attribute__((aligned (4))); |
|
|
|
|
} |
|
|
|
|
OptFeature; |
|
|
|
|
|
|
|
|
|
typedef struct __attribute__((aligned (128) )) GpuHidHaarTreeNode |
|
|
|
|
typedef struct __attribute__((aligned(4))) DTreeNode |
|
|
|
|
{ |
|
|
|
|
int p[CV_HAAR_FEATURE_MAX][4] __attribute__((aligned (64))); |
|
|
|
|
float weight[CV_HAAR_FEATURE_MAX]; |
|
|
|
|
float threshold; |
|
|
|
|
float alpha[3] __attribute__((aligned (16))); |
|
|
|
|
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))); |
|
|
|
|
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))); |
|
|
|
|
int nodeCount __attribute__((aligned (4))); |
|
|
|
|
} |
|
|
|
|
GpuHidHaarStageClassifier; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
//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 |
|
|
|
|
) |
|
|
|
|
DTree; |
|
|
|
|
|
|
|
|
|
typedef struct __attribute__((aligned (4))) Stage |
|
|
|
|
{ |
|
|
|
|
// 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<DATA_SIZE; i+=(LSx*LSy)) |
|
|
|
|
{ |
|
|
|
|
int index = i+lid; // index in shared local memory |
|
|
|
|
if(index<DATA_SIZE) |
|
|
|
|
{// calc global x,y coordinat and read data from there |
|
|
|
|
int x = min(GroupX + (index % (DATA_SIZE_X)),Width-1); |
|
|
|
|
int y = min(GroupY + (index / (DATA_SIZE_X)),Height-1); |
|
|
|
|
SumL[index] = sum[ImgOffset+y*pixelstep+x]; |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
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<OUTPUTSZ) |
|
|
|
|
{ |
|
|
|
|
int x = GroupX+lid_x; |
|
|
|
|
int y = GroupY+lid_y; |
|
|
|
|
int4 candidate_result; |
|
|
|
|
candidate_result.x = convert_int_rtn(x*ScaleFactor); |
|
|
|
|
candidate_result.y = convert_int_rtn(y*ScaleFactor); |
|
|
|
|
candidate_result.z = convert_int_rtn(ScaleFactor*WND_SIZE_X); |
|
|
|
|
candidate_result.w = convert_int_rtn(ScaleFactor*WND_SIZE_Y); |
|
|
|
|
candidate[index] = candidate_result; |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
}//end gpuRunHaarClassifierCascade |
|
|
|
|
#else |
|
|
|
|
|
|
|
|
|
__kernel void __attribute__((reqd_work_group_size(8,8,1)))gpuRunHaarClassifierCascade( |
|
|
|
|
global GpuHidHaarStageClassifier * stagecascadeptr, |
|
|
|
|
global int4 * info, |
|
|
|
|
global GpuHidHaarTreeNode * nodeptr, |
|
|
|
|
global const int * restrict sum1, |
|
|
|
|
global const float * restrict sqsum1, |
|
|
|
|
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) |
|
|
|
|
int first __attribute__((aligned (4))); |
|
|
|
|
int ntrees __attribute__((aligned (4))); |
|
|
|
|
float threshold __attribute__((aligned (4))); |
|
|
|
|
} |
|
|
|
|
Stage; |
|
|
|
|
|
|
|
|
|
__kernel void runHaarClassifierStump( |
|
|
|
|
__global const int* sum, |
|
|
|
|
int sumstep, int sumoffset, |
|
|
|
|
__global const int* sqsum, |
|
|
|
|
int sqsumstep, int sqsumoffset, |
|
|
|
|
__global const OptFeature* optfeatures, |
|
|
|
|
|
|
|
|
|
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) |
|
|
|
|
{ |
|
|
|
|
int grpszx = get_local_size(0); |
|
|
|
|
int grpszy = get_local_size(1); |
|
|
|
|
int grpnumx = get_num_groups(0); |
|
|
|
|
int grpidx = get_group_id(0); |
|
|
|
|
int lclidx = get_local_id(0); |
|
|
|
|
int lclidy = get_local_id(1); |
|
|
|
|
|
|
|
|
|
int lcl_sz = mul24(grpszx,grpszy); |
|
|
|
|
int lcl_id = mad24(lclidy,grpszx,lclidx); |
|
|
|
|
|
|
|
|
|
__local int lclshare[1024]; |
|
|
|
|
__local int* lcldata = lclshare;//for save win data |
|
|
|
|
__local int* glboutindex = lcldata + 28*28;//for save global out index |
|
|
|
|
__local int* lclcount = glboutindex + 1;//for save the numuber of temp pass pixel |
|
|
|
|
__local int* lcloutindex = lclcount + 1;//for save info of temp pass pixel |
|
|
|
|
__local float* partialsum = (__local float*)(lcloutindex + (lcl_sz<<1)); |
|
|
|
|
glboutindex[0]=0; |
|
|
|
|
int outputoff = mul24(grpidx,256); |
|
|
|
|
|
|
|
|
|
//assume window size is 20X20 |
|
|
|
|
#define WINDOWSIZE 20+1 |
|
|
|
|
//make sure readwidth is the multiple of 4 |
|
|
|
|
//ystep =1, from host code |
|
|
|
|
int readwidth = ((grpszx-1 + WINDOWSIZE+3)>>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 <loopcount; scalei++) |
|
|
|
|
int ix = get_global_id(0)*xyscale; |
|
|
|
|
int iy = get_global_id(1)*xyscale; |
|
|
|
|
sumstep /= sizeof(int); |
|
|
|
|
sqsumstep /= sizeof(int); |
|
|
|
|
|
|
|
|
|
if( ix < imgsize.x && iy < imgsize.y ) |
|
|
|
|
{ |
|
|
|
|
int4 scaleinfo1= info[scalei]; |
|
|
|
|
int height = scaleinfo1.x & 0xffff; |
|
|
|
|
int grpnumperline =(scaleinfo1.y & 0xffff0000) >> 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<totalgrp; grploop+=grpnumx) |
|
|
|
|
int ntrees, nodeOfs = 0, leafOfs = 0; |
|
|
|
|
int stageIdx, i; |
|
|
|
|
float s = 0.f; |
|
|
|
|
__global const DTreeNode* node; |
|
|
|
|
__global const OptFeature* f; |
|
|
|
|
|
|
|
|
|
__global const int* psum = sum + mad24(iy, sumstep, ix); |
|
|
|
|
__global const int* pnsum = psum + mad24(normrect.y, sumstep, normrect.x); |
|
|
|
|
int normarea = normrect.z * normrect.w; |
|
|
|
|
float invarea = 1.f/normarea; |
|
|
|
|
float sval = (pnsum[0] - pnsum[normrect.z] - pnsum[mul24(normrect.w, sumstep)] + |
|
|
|
|
pnsum[mad24(normrect.w, sumstep, normrect.z)])*invarea; |
|
|
|
|
float sqval = (sqsum[mad24(iy + normrect.y, sqsumstep, ix + normrect.x)])*invarea; |
|
|
|
|
float nf = (float)normarea * sqrt(max(sqval - sval * sval, 0.f)); |
|
|
|
|
float4 weight; |
|
|
|
|
int4 ofs; |
|
|
|
|
nf = nf > 0 ? nf : 1.f; |
|
|
|
|
|
|
|
|
|
for( stageIdx = 0; stageIdx < nstages; stageIdx++ ) |
|
|
|
|
{ |
|
|
|
|
int grpidy = grploop / grpnumperline; |
|
|
|
|
int grpidx = grploop - mul24(grpidy, grpnumperline); |
|
|
|
|
int x = mad24(grpidx,grpszx,lclidx); |
|
|
|
|
int y = mad24(grpidy,grpszy,lclidy); |
|
|
|
|
int grpoffx = x-lclidx; |
|
|
|
|
int grpoffy = y-lclidy; |
|
|
|
|
|
|
|
|
|
for(int i=0; i<read_loop; i++) |
|
|
|
|
ntrees = stages[stageIdx].ntrees; |
|
|
|
|
s = 0.f; |
|
|
|
|
for( i = 0; i < ntrees; i++, nodeOfs++, leafOfs += 2 ) |
|
|
|
|
{ |
|
|
|
|
int pos_id = mad24(i,lcl_sz,lcl_id); |
|
|
|
|
pos_id = pos_id < total_read ? pos_id : 0; |
|
|
|
|
|
|
|
|
|
int lcl_y = pos_id / read_horiz_cnt; |
|
|
|
|
int lcl_x = pos_id - mul24(lcl_y, read_horiz_cnt); |
|
|
|
|
|
|
|
|
|
int glb_x = grpoffx + (lcl_x<<2); |
|
|
|
|
int glb_y = grpoffy + lcl_y; |
|
|
|
|
|
|
|
|
|
int glb_off = mad24(min(glb_y, height + WINDOWSIZE - 1),pixelstep,glb_x); |
|
|
|
|
int4 data = *(__global int4*)&sum[glb_off]; |
|
|
|
|
int lcl_off = mad24(lcl_y, readwidth, lcl_x<<2); |
|
|
|
|
|
|
|
|
|
vstore4(data, 0, &lcldata[lcl_off]); |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
lcloutindex[lcl_id] = 0; |
|
|
|
|
lclcount[0] = 0; |
|
|
|
|
int result = 1; |
|
|
|
|
int nodecounter= startnode; |
|
|
|
|
float mean, variance_norm_factor; |
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE); |
|
|
|
|
|
|
|
|
|
int lcl_off = mad24(lclidy,readwidth,lclidx); |
|
|
|
|
int4 cascadeinfo1, cascadeinfo2; |
|
|
|
|
cascadeinfo1 = p; |
|
|
|
|
cascadeinfo2 = pq; |
|
|
|
|
|
|
|
|
|
cascadeinfo1.x +=lcl_off; |
|
|
|
|
cascadeinfo1.z +=lcl_off; |
|
|
|
|
mean = (lcldata[mad24(cascadeinfo1.y,readwidth,cascadeinfo1.x)] - lcldata[mad24(cascadeinfo1.y,readwidth,cascadeinfo1.z)] - |
|
|
|
|
lcldata[mad24(cascadeinfo1.w,readwidth,cascadeinfo1.x)] + lcldata[mad24(cascadeinfo1.w,readwidth,cascadeinfo1.z)]) |
|
|
|
|
*correction; |
|
|
|
|
|
|
|
|
|
int p_offset = mad24(y, pixelstep, x); |
|
|
|
|
|
|
|
|
|
cascadeinfo2.x +=p_offset; |
|
|
|
|
cascadeinfo2.z +=p_offset; |
|
|
|
|
variance_norm_factor =sqsum[mad24(cascadeinfo2.y, pixelstep, cascadeinfo2.x)] - sqsum[mad24(cascadeinfo2.y, pixelstep, cascadeinfo2.z)] - |
|
|
|
|
sqsum[mad24(cascadeinfo2.w, pixelstep, cascadeinfo2.x)] + sqsum[mad24(cascadeinfo2.w, pixelstep, cascadeinfo2.z)]; |
|
|
|
|
|
|
|
|
|
variance_norm_factor = variance_norm_factor * correction - mean * mean; |
|
|
|
|
variance_norm_factor = variance_norm_factor >=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) |
|
|
|
|
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 ) |
|
|
|
|
{ |
|
|
|
|
int queueindex = atomic_inc(lclcount); |
|
|
|
|
lcloutindex[queueindex<<1] = (lclidy << 16) | lclidx; |
|
|
|
|
lcloutindex[(queueindex<<1)+1] = as_int((float)variance_norm_factor); |
|
|
|
|
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 ]; |
|
|
|
|
} |
|
|
|
|
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)-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<queuecount_loop; 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_pixel = mad24(((temp_coord & (int)0xffff0000)>>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_loop<lcl_loops && tempnodecounter<stagecount;) |
|
|
|
|
{ |
|
|
|
|
__global GpuHidHaarTreeNode* currentnodeptr = (__global GpuHidHaarTreeNode*) |
|
|
|
|
(((__global uchar*)nodeptr) + sizeof(GpuHidHaarTreeNode) * ((nodecounter + tempnodecounter) * stump_factor + root_offset)); |
|
|
|
|
|
|
|
|
|
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 +=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<lcl_loops;lcl_loop++) |
|
|
|
|
partialsum[lcl_id]=part_sum; |
|
|
|
|
} |
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE); |
|
|
|
|
if(lcl_compute_win_id < queuecount) |
|
|
|
|
{ |
|
|
|
|
for(int i=0; i<lcl_compute_win && (lcl_compute_id==0); i++) |
|
|
|
|
{ |
|
|
|
|
stage_sum += partialsum[lcl_id+i]; |
|
|
|
|
} |
|
|
|
|
if(stage_sum >= 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<<perfscale); |
|
|
|
|
} |
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE); |
|
|
|
|
}//end for(int queueloop=0;queueloop<queuecount_loop;queueloop++) |
|
|
|
|
|
|
|
|
|
queuecount = lclcount[0]; |
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE); |
|
|
|
|
nodecounter += stagecount; |
|
|
|
|
}//end for(int stageloop = splitstage; stageloop< endstage && queuecount>0;stageloop++) |
|
|
|
|
|
|
|
|
|
if(lcl_id<queuecount) |
|
|
|
|
|
|
|
|
|
if( s < stages[stageIdx].threshold ) |
|
|
|
|
break; |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
if( stageIdx == nstages ) |
|
|
|
|
{ |
|
|
|
|
int nfaces = atomic_inc(facepos); |
|
|
|
|
//printf("detected face #d!!!!\n", nfaces); |
|
|
|
|
if( nfaces < MAX_FACES ) |
|
|
|
|
{ |
|
|
|
|
int temp = lcloutindex[lcl_id<<1]; |
|
|
|
|
int x = mad24(grpidx,grpszx,temp & 0xffff); |
|
|
|
|
int y = mad24(grpidy,grpszy,((temp & (int)0xffff0000) >> 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; |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
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); |
|
|
|
|
}//end for(int grploop=grpidx;grploop<totalgrp;grploop+=grpnumx) |
|
|
|
|
}//end for(int scalei = 0; scalei <loopcount; scalei++) |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
#endif |
|
|
|
|
|
Vadim Pisarevsky
11 years ago