Skip to content

O
opencv
Commit af33c118 authored by perping's avatar perping
Browse files
Options

fixed a bug of haar.

parent dfa5a27b
Show whitespace changes
Inline Side-by-side
Showing with 64 additions and 17 deletions
modules/ocl/src/haar.cpp
View file @ af33c118
......@@ -1676,9 +1676,9 @@ void cv::ocl::OclCascadeClassifierBuf::CreateFactorRelatedBufs(
{
sz = sizev[i];
factor = scalev[i];
int ystep = cvRound(std::max(2., factor));
int width = (cols - 1 - sz.width + ystep - 1) / ystep;
int height = (rows - 1 - sz.height + ystep - 1) / ystep;
double ystep = cv::max(2.,factor);
int width = cvRound((cols - 1 - sz.width + ystep - 1) / ystep);
int height = cvRound((rows - 1 - sz.height + ystep - 1) / ystep);
int grpnumperline = (width + localThreads[0] - 1) / localThreads[0];
int totalgrp = ((height + localThreads[1] - 1) / localThreads[1]) * grpnumperline;
......
modules/ocl/src/opencl/haarobjectdetect.cl
View file @ af33c118
......@@ -11,6 +11,7 @@
// Jia Haipeng, jiahaipeng95@gmail.com
// Nathan, liujun@multicorewareinc.com
// Peng Xiao, pengxiao@outlook.com
// Erping Pang, erping@multicorewareinc.com
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
......@@ -321,7 +322,7 @@ __kernel void __attribute__((reqd_work_group_size(8,8,1)))gpuRunHaarClassifierCa
int glb_x = grpoffx + (lcl_x<<2);
int glb_y = grpoffy + lcl_y;
int glb_off = mad24(min(glb_y, height - 1),pixelstep,glb_x);
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);
......@@ -421,12 +422,25 @@ __kernel void __attribute__((reqd_work_group_size(8,8,1)))gpuRunHaarClassifierCa
result = (stage_sum >= stagethreshold);
}
if(factor < 2)
{
if(result && lclidx %2 ==0 && lclidy %2 ==0 )
{
if(result && (x < width) && (y < height))
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(variance_norm_factor);
lcloutindex[(queueindex<<1)+1] = as_int((float)variance_norm_factor);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
int queuecount = lclcount[0];
......@@ -549,11 +563,27 @@ __kernel void __attribute__((reqd_work_group_size(8,8,1)))gpuRunHaarClassifierCa
int y = mad24(grpidy,grpszy,((temp & (int)0xffff0000) >> 16));
temp = glboutindex[0];
int4 candidate_result;
candidate_result.zw = (int2)convert_int_rtn(factor*20.f);
candidate_result.x = convert_int_rtn(x*factor);
candidate_result.y = convert_int_rtn(y*factor);
candidate_result.zw = (int2)convert_int_rtn(round(factor*20.f));
candidate_result.x = convert_int_rtn(round(x*factor));
candidate_result.y = convert_int_rtn(round(y*factor));
atomic_inc(glboutindex);
candidate[outputoff+temp+lcl_id] = candidate_result;
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<totalgrp;grploop+=grpnumx)
......
modules/ocl/src/opencl/haarobjectdetect_scaled2.cl
View file @ af33c118
......@@ -18,6 +18,7 @@
// Wu Xinglong, wxl370@126.com
// Sen Liu, swjtuls1987@126.com
// Peng Xiao, pengxiao@outlook.com
// Erping Pang, erping@multicorewareinc.com
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
......@@ -142,7 +143,7 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
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);
float ystep = max(2.0f, factor);
for (int grploop = get_group_id(0); grploop < totalgrp; grploop += grpnumx)
{
......@@ -151,8 +152,8 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
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;
int x = round(ix * ystep);
int y = round(iy * ystep);
lcloutindex[lcl_id] = 0;
lclcount[0] = 0;
int nodecounter;
......@@ -243,7 +244,7 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
barrier(CLK_LOCAL_MEM_FENCE);
if (result && (ix < width) && (iy < height))
if (result)
{
int queueindex = atomic_inc(lclcount);
lcloutindex[queueindex] = (y << 16) | x;
......@@ -258,10 +259,26 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
int y = (temp & (int)0xffff0000) >> 16;
temp = atomic_inc(glboutindex);
int4 candidate_result;
candidate_result.zw = (int2)convert_int_rtn(factor * 20.f);
candidate_result.zw = (int2)convert_int_rtn(round(factor * 20.f));
candidate_result.x = x;
candidate_result.y = y;
candidate[outputoff + temp + lcl_id] = candidate_result;
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);
......@@ -284,7 +301,7 @@ __kernel void gpuscaleclassifier(global GpuHidHaarTreeNode *orinode, global GpuH
tr_h[i] = (int)(t1.p[i][3] * 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]);
t1.weight[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]);
counter += nodenum;
#pragma unroll
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment