| // This file is auto-generated. Do not edit! |
| |
| #include "precomp.hpp" |
| #include "opencl_kernels_objdetect.hpp" |
| |
| namespace cv |
| { |
| namespace ocl |
| { |
| namespace objdetect |
| { |
| |
| const struct ProgramEntry cascadedetect={"cascadedetect", |
| "#ifdef HAAR\n" |
| "typedef struct __attribute__((aligned(4))) OptHaarFeature\n" |
| "{\n" |
| "int4 ofs[3] __attribute__((aligned (4)));\n" |
| "float4 weight __attribute__((aligned (4)));\n" |
| "}\n" |
| "OptHaarFeature;\n" |
| "#endif\n" |
| "#ifdef LBP\n" |
| "typedef struct __attribute__((aligned(4))) OptLBPFeature\n" |
| "{\n" |
| "int16 ofs __attribute__((aligned (4)));\n" |
| "}\n" |
| "OptLBPFeature;\n" |
| "#endif\n" |
| "typedef struct __attribute__((aligned(4))) Stump\n" |
| "{\n" |
| "float4 st __attribute__((aligned (4)));\n" |
| "}\n" |
| "Stump;\n" |
| "typedef struct __attribute__((aligned(4))) Node\n" |
| "{\n" |
| "int4 n __attribute__((aligned (4)));\n" |
| "}\n" |
| "Node;\n" |
| "typedef struct __attribute__((aligned (4))) Stage\n" |
| "{\n" |
| "int first __attribute__((aligned (4)));\n" |
| "int ntrees __attribute__((aligned (4)));\n" |
| "float threshold __attribute__((aligned (4)));\n" |
| "}\n" |
| "Stage;\n" |
| "typedef struct __attribute__((aligned (4))) ScaleData\n" |
| "{\n" |
| "float scale __attribute__((aligned (4)));\n" |
| "int szi_width __attribute__((aligned (4)));\n" |
| "int szi_height __attribute__((aligned (4)));\n" |
| "int layer_ofs __attribute__((aligned (4)));\n" |
| "int ystep __attribute__((aligned (4)));\n" |
| "}\n" |
| "ScaleData;\n" |
| "#ifndef SUM_BUF_SIZE\n" |
| "#define SUM_BUF_SIZE 0\n" |
| "#endif\n" |
| "#ifndef NODE_COUNT\n" |
| "#define NODE_COUNT 1\n" |
| "#endif\n" |
| "#ifdef HAAR\n" |
| "__kernel __attribute__((reqd_work_group_size(LOCAL_SIZE_X,LOCAL_SIZE_Y,1)))\n" |
| "void runHaarClassifier(\n" |
| "int nscales, __global const ScaleData* scaleData,\n" |
| "__global const int* sum,\n" |
| "int _sumstep, int sumoffset,\n" |
| "__global const OptHaarFeature* optfeatures,\n" |
| "__global const Stage* stages,\n" |
| "__global const Node* nodes,\n" |
| "__global const float* leaves0,\n" |
| "volatile __global int* facepos,\n" |
| "int4 normrect, int sqofs, int2 windowsize)\n" |
| "{\n" |
| "int lx = get_local_id(0);\n" |
| "int ly = get_local_id(1);\n" |
| "int groupIdx = get_group_id(0);\n" |
| "int i, ngroups = get_global_size(0)/LOCAL_SIZE_X;\n" |
| "int scaleIdx, tileIdx, stageIdx;\n" |
| "int sumstep = (int)(_sumstep/sizeof(int));\n" |
| "int4 nofs0 = (int4)(mad24(normrect.y, sumstep, normrect.x),\n" |
| "mad24(normrect.y, sumstep, normrect.x + normrect.z),\n" |
| "mad24(normrect.y + normrect.w, sumstep, normrect.x),\n" |
| "mad24(normrect.y + normrect.w, sumstep, normrect.x + normrect.z));\n" |
| "int normarea = normrect.z * normrect.w;\n" |
| "float invarea = 1.f/normarea;\n" |
| "int lidx = ly*LOCAL_SIZE_X + lx;\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "int4 nofs = (int4)(mad24(normrect.y, SUM_BUF_STEP, normrect.x),\n" |
| "mad24(normrect.y, SUM_BUF_STEP, normrect.x + normrect.z),\n" |
| "mad24(normrect.y + normrect.w, SUM_BUF_STEP, normrect.x),\n" |
| "mad24(normrect.y + normrect.w, SUM_BUF_STEP, normrect.x + normrect.z));\n" |
| "#else\n" |
| "int4 nofs = nofs0;\n" |
| "#endif\n" |
| "#define LOCAL_SIZE (LOCAL_SIZE_X*LOCAL_SIZE_Y)\n" |
| "__local int lstore[SUM_BUF_SIZE + LOCAL_SIZE*5/2+1];\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "__local int* ibuf = lstore;\n" |
| "__local int* lcount = ibuf + SUM_BUF_SIZE;\n" |
| "#else\n" |
| "__local int* lcount = lstore;\n" |
| "#endif\n" |
| "__local float* lnf = (__local float*)(lcount + 1);\n" |
| "__local float* lpartsum = lnf + LOCAL_SIZE;\n" |
| "__local short* lbuf = (__local short*)(lpartsum + LOCAL_SIZE);\n" |
| "for( scaleIdx = nscales-1; scaleIdx >= 0; scaleIdx-- )\n" |
| "{\n" |
| "__global const ScaleData* s = scaleData + scaleIdx;\n" |
| "int ystep = s->ystep;\n" |
| "int2 worksize = (int2)(max(s->szi_width - windowsize.x, 0), max(s->szi_height - windowsize.y, 0));\n" |
| "int2 ntiles = (int2)((worksize.x + LOCAL_SIZE_X-1)/LOCAL_SIZE_X,\n" |
| "(worksize.y + LOCAL_SIZE_Y-1)/LOCAL_SIZE_Y);\n" |
| "int totalTiles = ntiles.x*ntiles.y;\n" |
| "for( tileIdx = groupIdx; tileIdx < totalTiles; tileIdx += ngroups )\n" |
| "{\n" |
| "int ix0 = (tileIdx % ntiles.x)*LOCAL_SIZE_X;\n" |
| "int iy0 = (tileIdx / ntiles.x)*LOCAL_SIZE_Y;\n" |
| "int ix = lx, iy = ly;\n" |
| "__global const int* psum0 = sum + mad24(iy0, sumstep, ix0) + s->layer_ofs;\n" |
| "__global const int* psum1 = psum0 + mad24(iy, sumstep, ix);\n" |
| "if( ix0 >= worksize.x || iy0 >= worksize.y )\n" |
| "continue;\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "for( i = lidx*4; i < SUM_BUF_SIZE; i += LOCAL_SIZE_X*LOCAL_SIZE_Y*4 )\n" |
| "{\n" |
| "int dy = i/SUM_BUF_STEP, dx = i - dy*SUM_BUF_STEP;\n" |
| "vstore4(vload4(0, psum0 + mad24(dy, sumstep, dx)), 0, ibuf+i);\n" |
| "}\n" |
| "#endif\n" |
| "if( lidx == 0 )\n" |
| "lcount[0] = 0;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if( ix0 + ix < worksize.x && iy0 + iy < worksize.y )\n" |
| "{\n" |
| "#if NODE_COUNT==1\n" |
| "__global const Stump* stump = (__global const Stump*)nodes;\n" |
| "#else\n" |
| "__global const Node* node = nodes;\n" |
| "__global const float* leaves = leaves0;\n" |
| "#endif\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "__local const int* psum = ibuf + mad24(iy, SUM_BUF_STEP, ix);\n" |
| "#else\n" |
| "__global const int* psum = psum1;\n" |
| "#endif\n" |
| "__global const int* psqsum = (__global const int*)(psum1 + sqofs);\n" |
| "float sval = (psum[nofs.x] - psum[nofs.y] - psum[nofs.z] + psum[nofs.w])*invarea;\n" |
| "float sqval = (psqsum[nofs0.x] - psqsum[nofs0.y] - psqsum[nofs0.z] + psqsum[nofs0.w])*invarea;\n" |
| "float nf = (float)normarea * sqrt(max(sqval - sval * sval, 0.f));\n" |
| "nf = nf > 0 ? nf : 1.f;\n" |
| "for( stageIdx = 0; stageIdx < SPLIT_STAGE; stageIdx++ )\n" |
| "{\n" |
| "int ntrees = stages[stageIdx].ntrees;\n" |
| "float s = 0.f;\n" |
| "#if NODE_COUNT==1\n" |
| "for( i = 0; i < ntrees; i++ )\n" |
| "{\n" |
| "float4 st = stump[i].st;\n" |
| "__global const OptHaarFeature* f = optfeatures + as_int(st.x);\n" |
| "float4 weight = f->weight;\n" |
| "int4 ofs = f->ofs[0];\n" |
| "sval = (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.x;\n" |
| "ofs = f->ofs[1];\n" |
| "sval = mad((psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w]), weight.y, sval);\n" |
| "if( weight.z > 0 )\n" |
| "{\n" |
| "ofs = f->ofs[2];\n" |
| "sval = mad((psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w]), weight.z, sval);\n" |
| "}\n" |
| "s += (sval < st.y*nf) ? st.z : st.w;\n" |
| "}\n" |
| "stump += ntrees;\n" |
| "#else\n" |
| "for( i = 0; i < ntrees; i++, node += NODE_COUNT, leaves += NODE_COUNT+1 )\n" |
| "{\n" |
| "int idx = 0;\n" |
| "do\n" |
| "{\n" |
| "int4 n = node[idx].n;\n" |
| "__global const OptHaarFeature* f = optfeatures + n.x;\n" |
| "float4 weight = f->weight;\n" |
| "int4 ofs = f->ofs[0];\n" |
| "sval = (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.x;\n" |
| "ofs = f->ofs[1];\n" |
| "sval = mad((psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w]), weight.y, sval);\n" |
| "if( weight.z > 0 )\n" |
| "{\n" |
| "ofs = f->ofs[2];\n" |
| "sval = mad((psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w]), weight.z, sval);\n" |
| "}\n" |
| "idx = (sval < as_float(n.y)*nf) ? n.z : n.w;\n" |
| "}\n" |
| "while(idx > 0);\n" |
| "s += leaves[-idx];\n" |
| "}\n" |
| "#endif\n" |
| "if( s < stages[stageIdx].threshold )\n" |
| "break;\n" |
| "}\n" |
| "if( stageIdx == SPLIT_STAGE && (ystep == 1 || ((ix | iy) & 1) == 0) )\n" |
| "{\n" |
| "int count = atomic_inc(lcount);\n" |
| "lbuf[count] = (int)(ix | (iy << 8));\n" |
| "lnf[count] = nf;\n" |
| "}\n" |
| "}\n" |
| "for( stageIdx = SPLIT_STAGE; stageIdx < N_STAGES; stageIdx++ )\n" |
| "{\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "int nrects = lcount[0];\n" |
| "if( nrects == 0 )\n" |
| "break;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if( lidx == 0 )\n" |
| "lcount[0] = 0;\n" |
| "{\n" |
| "#if NODE_COUNT == 1\n" |
| "__global const Stump* stump = (__global const Stump*)nodes + stages[stageIdx].first;\n" |
| "#else\n" |
| "__global const Node* node = nodes + stages[stageIdx].first*NODE_COUNT;\n" |
| "__global const float* leaves = leaves0 + stages[stageIdx].first*(NODE_COUNT+1);\n" |
| "#endif\n" |
| "int nparts = LOCAL_SIZE / nrects;\n" |
| "int ntrees = stages[stageIdx].ntrees;\n" |
| "int ntrees_p = (ntrees + nparts - 1)/nparts;\n" |
| "int nr = lidx / nparts;\n" |
| "int partidx = -1, idxval = 0;\n" |
| "float partsum = 0.f, nf = 0.f;\n" |
| "if( nr < nrects )\n" |
| "{\n" |
| "partidx = lidx % nparts;\n" |
| "idxval = lbuf[nr];\n" |
| "nf = lnf[nr];\n" |
| "{\n" |
| "int ntrees0 = ntrees_p*partidx;\n" |
| "int ntrees1 = min(ntrees0 + ntrees_p, ntrees);\n" |
| "int ix1 = idxval & 255, iy1 = idxval >> 8;\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "__local const int* psum = ibuf + mad24(iy1, SUM_BUF_STEP, ix1);\n" |
| "#else\n" |
| "__global const int* psum = psum0 + mad24(iy1, sumstep, ix1);\n" |
| "#endif\n" |
| "#if NODE_COUNT == 1\n" |
| "for( i = ntrees0; i < ntrees1; i++ )\n" |
| "{\n" |
| "float4 st = stump[i].st;\n" |
| "__global const OptHaarFeature* f = optfeatures + as_int(st.x);\n" |
| "float4 weight = f->weight;\n" |
| "int4 ofs = f->ofs[0];\n" |
| "float sval = (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.x;\n" |
| "ofs = f->ofs[1];\n" |
| "sval = mad((psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w]), weight.y, sval);\n" |
| "if( fabs(weight.z) > 0 )\n" |
| "{\n" |
| "ofs = f->ofs[2];\n" |
| "sval = mad((psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w]), weight.z, sval);\n" |
| "}\n" |
| "partsum += (sval < st.y*nf) ? st.z : st.w;\n" |
| "}\n" |
| "#else\n" |
| "for( i = ntrees0; i < ntrees1; i++ )\n" |
| "{\n" |
| "int idx = 0;\n" |
| "do\n" |
| "{\n" |
| "int4 n = node[i*2 + idx].n;\n" |
| "__global const OptHaarFeature* f = optfeatures + n.x;\n" |
| "float4 weight = f->weight;\n" |
| "int4 ofs = f->ofs[0];\n" |
| "float sval = (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.x;\n" |
| "ofs = f->ofs[1];\n" |
| "sval = mad((psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w]), weight.y, sval);\n" |
| "if( weight.z > 0 )\n" |
| "{\n" |
| "ofs = f->ofs[2];\n" |
| "sval = mad((psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w]), weight.z, sval);\n" |
| "}\n" |
| "idx = (sval < as_float(n.y)*nf) ? n.z : n.w;\n" |
| "}\n" |
| "while(idx > 0);\n" |
| "partsum += leaves[i*3-idx];\n" |
| "}\n" |
| "#endif\n" |
| "}\n" |
| "}\n" |
| "lpartsum[lidx] = partsum;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if( partidx == 0 )\n" |
| "{\n" |
| "float s = lpartsum[nr*nparts];\n" |
| "for( i = 1; i < nparts; i++ )\n" |
| "s += lpartsum[i + nr*nparts];\n" |
| "if( s >= stages[stageIdx].threshold )\n" |
| "{\n" |
| "int count = atomic_inc(lcount);\n" |
| "lbuf[count] = idxval;\n" |
| "lnf[count] = nf;\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if( stageIdx == N_STAGES )\n" |
| "{\n" |
| "int nrects = lcount[0];\n" |
| "if( lidx < nrects )\n" |
| "{\n" |
| "int nfaces = atomic_inc(facepos);\n" |
| "if( nfaces < MAX_FACES )\n" |
| "{\n" |
| "volatile __global int* face = facepos + 1 + nfaces*3;\n" |
| "int val = lbuf[lidx];\n" |
| "face[0] = scaleIdx;\n" |
| "face[1] = ix0 + (val & 255);\n" |
| "face[2] = iy0 + (val >> 8);\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "#endif\n" |
| "#ifdef LBP\n" |
| "#undef CALC_SUM_OFS_\n" |
| "#define CALC_SUM_OFS_(p0, p1, p2, p3, ptr) \\\n" |
| "((ptr)[p0] - (ptr)[p1] - (ptr)[p2] + (ptr)[p3])\n" |
| "__kernel void runLBPClassifierStumpSimple(\n" |
| "int nscales, __global const ScaleData* scaleData,\n" |
| "__global const int* sum,\n" |
| "int _sumstep, int sumoffset,\n" |
| "__global const OptLBPFeature* optfeatures,\n" |
| "__global const Stage* stages,\n" |
| "__global const Stump* stumps,\n" |
| "__global const int* bitsets,\n" |
| "int bitsetSize,\n" |
| "volatile __global int* facepos,\n" |
| "int2 windowsize)\n" |
| "{\n" |
| "int lx = get_local_id(0);\n" |
| "int ly = get_local_id(1);\n" |
| "int local_size_x = get_local_size(0);\n" |
| "int local_size_y = get_local_size(1);\n" |
| "int groupIdx = get_group_id(1)*get_num_groups(0) + get_group_id(0);\n" |
| "int ngroups = get_num_groups(0)*get_num_groups(1);\n" |
| "int scaleIdx, tileIdx, stageIdx;\n" |
| "int sumstep = (int)(_sumstep/sizeof(int));\n" |
| "for( scaleIdx = nscales-1; scaleIdx >= 0; scaleIdx-- )\n" |
| "{\n" |
| "__global const ScaleData* s = scaleData + scaleIdx;\n" |
| "int ystep = s->ystep;\n" |
| "int2 worksize = (int2)(max(s->szi_width - windowsize.x, 0), max(s->szi_height - windowsize.y, 0));\n" |
| "int2 ntiles = (int2)((worksize.x/ystep + local_size_x-1)/local_size_x,\n" |
| "(worksize.y/ystep + local_size_y-1)/local_size_y);\n" |
| "int totalTiles = ntiles.x*ntiles.y;\n" |
| "for( tileIdx = groupIdx; tileIdx < totalTiles; tileIdx += ngroups )\n" |
| "{\n" |
| "int iy = mad24((tileIdx / ntiles.x), local_size_y, ly) * ystep;\n" |
| "int ix = mad24((tileIdx % ntiles.x), local_size_x, lx) * ystep;\n" |
| "if( ix < worksize.x && iy < worksize.y )\n" |
| "{\n" |
| "__global const int* p = sum + mad24(iy, sumstep, ix) + s->layer_ofs;\n" |
| "__global const Stump* stump = stumps;\n" |
| "__global const int* bitset = bitsets;\n" |
| "for( stageIdx = 0; stageIdx < N_STAGES; stageIdx++ )\n" |
| "{\n" |
| "int i, ntrees = stages[stageIdx].ntrees;\n" |
| "float s = 0.f;\n" |
| "for( i = 0; i < ntrees; i++, stump++, bitset += bitsetSize )\n" |
| "{\n" |
| "float4 st = stump->st;\n" |
| "__global const OptLBPFeature* f = optfeatures + as_int(st.x);\n" |
| "int16 ofs = f->ofs;\n" |
| "int cval = CALC_SUM_OFS_( ofs.s5, ofs.s6, ofs.s9, ofs.sa, p );\n" |
| "int mask, idx = (CALC_SUM_OFS_( ofs.s0, ofs.s1, ofs.s4, ofs.s5, p ) >= cval ? 4 : 0);\n" |
| "idx |= (CALC_SUM_OFS_( ofs.s1, ofs.s2, ofs.s5, ofs.s6, p ) >= cval ? 2 : 0);\n" |
| "idx |= (CALC_SUM_OFS_( ofs.s2, ofs.s3, ofs.s6, ofs.s7, p ) >= cval ? 1 : 0);\n" |
| "mask = (CALC_SUM_OFS_( ofs.s6, ofs.s7, ofs.sa, ofs.sb, p ) >= cval ? 16 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.sa, ofs.sb, ofs.se, ofs.sf, p ) >= cval ? 8 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s9, ofs.sa, ofs.sd, ofs.se, p ) >= cval ? 4 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s8, ofs.s9, ofs.sc, ofs.sd, p ) >= cval ? 2 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s4, ofs.s5, ofs.s8, ofs.s9, p ) >= cval ? 1 : 0);\n" |
| "s += (bitset[idx] & (1 << mask)) ? st.z : st.w;\n" |
| "}\n" |
| "if( s < stages[stageIdx].threshold )\n" |
| "break;\n" |
| "}\n" |
| "if( stageIdx == N_STAGES )\n" |
| "{\n" |
| "int nfaces = atomic_inc(facepos);\n" |
| "if( nfaces < MAX_FACES )\n" |
| "{\n" |
| "volatile __global int* face = facepos + 1 + nfaces*3;\n" |
| "face[0] = scaleIdx;\n" |
| "face[1] = ix;\n" |
| "face[2] = iy;\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "__kernel __attribute__((reqd_work_group_size(LOCAL_SIZE_X,LOCAL_SIZE_Y,1)))\n" |
| "void runLBPClassifierStump(\n" |
| "int nscales, __global const ScaleData* scaleData,\n" |
| "__global const int* sum,\n" |
| "int _sumstep, int sumoffset,\n" |
| "__global const OptLBPFeature* optfeatures,\n" |
| "__global const Stage* stages,\n" |
| "__global const Stump* stumps,\n" |
| "__global const int* bitsets,\n" |
| "int bitsetSize,\n" |
| "volatile __global int* facepos,\n" |
| "int2 windowsize)\n" |
| "{\n" |
| "int lx = get_local_id(0);\n" |
| "int ly = get_local_id(1);\n" |
| "int groupIdx = get_group_id(0);\n" |
| "int i, ngroups = get_global_size(0)/LOCAL_SIZE_X;\n" |
| "int scaleIdx, tileIdx, stageIdx;\n" |
| "int sumstep = (int)(_sumstep/sizeof(int));\n" |
| "int lidx = ly*LOCAL_SIZE_X + lx;\n" |
| "#define LOCAL_SIZE (LOCAL_SIZE_X*LOCAL_SIZE_Y)\n" |
| "__local int lstore[SUM_BUF_SIZE + LOCAL_SIZE*3/2+1];\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "__local int* ibuf = lstore;\n" |
| "__local int* lcount = ibuf + SUM_BUF_SIZE;\n" |
| "#else\n" |
| "__local int* lcount = lstore;\n" |
| "#endif\n" |
| "__local float* lpartsum = (__local float*)(lcount + 1);\n" |
| "__local short* lbuf = (__local short*)(lpartsum + LOCAL_SIZE);\n" |
| "for( scaleIdx = nscales-1; scaleIdx >= 0; scaleIdx-- )\n" |
| "{\n" |
| "__global const ScaleData* s = scaleData + scaleIdx;\n" |
| "int ystep = s->ystep;\n" |
| "int2 worksize = (int2)(max(s->szi_width - windowsize.x, 0), max(s->szi_height - windowsize.y, 0));\n" |
| "int2 ntiles = (int2)((worksize.x + LOCAL_SIZE_X-1)/LOCAL_SIZE_X,\n" |
| "(worksize.y + LOCAL_SIZE_Y-1)/LOCAL_SIZE_Y);\n" |
| "int totalTiles = ntiles.x*ntiles.y;\n" |
| "for( tileIdx = groupIdx; tileIdx < totalTiles; tileIdx += ngroups )\n" |
| "{\n" |
| "int ix0 = (tileIdx % ntiles.x)*LOCAL_SIZE_X;\n" |
| "int iy0 = (tileIdx / ntiles.x)*LOCAL_SIZE_Y;\n" |
| "int ix = lx, iy = ly;\n" |
| "__global const int* psum0 = sum + mad24(iy0, sumstep, ix0) + s->layer_ofs;\n" |
| "if( ix0 >= worksize.x || iy0 >= worksize.y )\n" |
| "continue;\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "for( i = lidx*4; i < SUM_BUF_SIZE; i += LOCAL_SIZE_X*LOCAL_SIZE_Y*4 )\n" |
| "{\n" |
| "int dy = i/SUM_BUF_STEP, dx = i - dy*SUM_BUF_STEP;\n" |
| "vstore4(vload4(0, psum0 + mad24(dy, sumstep, dx)), 0, ibuf+i);\n" |
| "}\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "#endif\n" |
| "if( lidx == 0 )\n" |
| "lcount[0] = 0;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if( ix0 + ix < worksize.x && iy0 + iy < worksize.y )\n" |
| "{\n" |
| "__global const Stump* stump = stumps;\n" |
| "__global const int* bitset = bitsets;\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "__local const int* p = ibuf + mad24(iy, SUM_BUF_STEP, ix);\n" |
| "#else\n" |
| "__global const int* p = psum0 + mad24(iy, sumstep, ix);\n" |
| "#endif\n" |
| "for( stageIdx = 0; stageIdx < SPLIT_STAGE; stageIdx++ )\n" |
| "{\n" |
| "int ntrees = stages[stageIdx].ntrees;\n" |
| "float s = 0.f;\n" |
| "for( i = 0; i < ntrees; i++, stump++, bitset += bitsetSize )\n" |
| "{\n" |
| "float4 st = stump->st;\n" |
| "__global const OptLBPFeature* f = optfeatures + as_int(st.x);\n" |
| "int16 ofs = f->ofs;\n" |
| "int cval = CALC_SUM_OFS_( ofs.s5, ofs.s6, ofs.s9, ofs.sa, p );\n" |
| "int mask, idx = (CALC_SUM_OFS_( ofs.s0, ofs.s1, ofs.s4, ofs.s5, p ) >= cval ? 4 : 0);\n" |
| "idx |= (CALC_SUM_OFS_( ofs.s1, ofs.s2, ofs.s5, ofs.s6, p ) >= cval ? 2 : 0);\n" |
| "idx |= (CALC_SUM_OFS_( ofs.s2, ofs.s3, ofs.s6, ofs.s7, p ) >= cval ? 1 : 0);\n" |
| "mask = (CALC_SUM_OFS_( ofs.s6, ofs.s7, ofs.sa, ofs.sb, p ) >= cval ? 16 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.sa, ofs.sb, ofs.se, ofs.sf, p ) >= cval ? 8 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s9, ofs.sa, ofs.sd, ofs.se, p ) >= cval ? 4 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s8, ofs.s9, ofs.sc, ofs.sd, p ) >= cval ? 2 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s4, ofs.s5, ofs.s8, ofs.s9, p ) >= cval ? 1 : 0);\n" |
| "s += (bitset[idx] & (1 << mask)) ? st.z : st.w;\n" |
| "}\n" |
| "if( s < stages[stageIdx].threshold )\n" |
| "break;\n" |
| "}\n" |
| "if( stageIdx == SPLIT_STAGE && (ystep == 1 || ((ix | iy) & 1) == 0) )\n" |
| "{\n" |
| "int count = atomic_inc(lcount);\n" |
| "lbuf[count] = (int)(ix | (iy << 8));\n" |
| "}\n" |
| "}\n" |
| "for( stageIdx = SPLIT_STAGE; stageIdx < N_STAGES; stageIdx++ )\n" |
| "{\n" |
| "int nrects = lcount[0];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if( nrects == 0 )\n" |
| "break;\n" |
| "if( lidx == 0 )\n" |
| "lcount[0] = 0;\n" |
| "{\n" |
| "__global const Stump* stump = stumps + stages[stageIdx].first;\n" |
| "__global const int* bitset = bitsets + stages[stageIdx].first*bitsetSize;\n" |
| "int nparts = LOCAL_SIZE / nrects;\n" |
| "int ntrees = stages[stageIdx].ntrees;\n" |
| "int ntrees_p = (ntrees + nparts - 1)/nparts;\n" |
| "int nr = lidx / nparts;\n" |
| "int partidx = -1, idxval = 0;\n" |
| "float partsum = 0.f, nf = 0.f;\n" |
| "if( nr < nrects )\n" |
| "{\n" |
| "partidx = lidx % nparts;\n" |
| "idxval = lbuf[nr];\n" |
| "{\n" |
| "int ntrees0 = ntrees_p*partidx;\n" |
| "int ntrees1 = min(ntrees0 + ntrees_p, ntrees);\n" |
| "int ix1 = idxval & 255, iy1 = idxval >> 8;\n" |
| "#if SUM_BUF_SIZE > 0\n" |
| "__local const int* p = ibuf + mad24(iy1, SUM_BUF_STEP, ix1);\n" |
| "#else\n" |
| "__global const int* p = psum0 + mad24(iy1, sumstep, ix1);\n" |
| "#endif\n" |
| "for( i = ntrees0; i < ntrees1; i++ )\n" |
| "{\n" |
| "float4 st = stump[i].st;\n" |
| "__global const OptLBPFeature* f = optfeatures + as_int(st.x);\n" |
| "int16 ofs = f->ofs;\n" |
| "#define CALC_SUM_OFS_(p0, p1, p2, p3, ptr) \\\n" |
| "((ptr)[p0] - (ptr)[p1] - (ptr)[p2] + (ptr)[p3])\n" |
| "int cval = CALC_SUM_OFS_( ofs.s5, ofs.s6, ofs.s9, ofs.sa, p );\n" |
| "int mask, idx = (CALC_SUM_OFS_( ofs.s0, ofs.s1, ofs.s4, ofs.s5, p ) >= cval ? 4 : 0);\n" |
| "idx |= (CALC_SUM_OFS_( ofs.s1, ofs.s2, ofs.s5, ofs.s6, p ) >= cval ? 2 : 0);\n" |
| "idx |= (CALC_SUM_OFS_( ofs.s2, ofs.s3, ofs.s6, ofs.s7, p ) >= cval ? 1 : 0);\n" |
| "mask = (CALC_SUM_OFS_( ofs.s6, ofs.s7, ofs.sa, ofs.sb, p ) >= cval ? 16 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.sa, ofs.sb, ofs.se, ofs.sf, p ) >= cval ? 8 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s9, ofs.sa, ofs.sd, ofs.se, p ) >= cval ? 4 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s8, ofs.s9, ofs.sc, ofs.sd, p ) >= cval ? 2 : 0);\n" |
| "mask |= (CALC_SUM_OFS_( ofs.s4, ofs.s5, ofs.s8, ofs.s9, p ) >= cval ? 1 : 0);\n" |
| "partsum += (bitset[i*bitsetSize + idx] & (1 << mask)) ? st.z : st.w;\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "lpartsum[lidx] = partsum;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if( partidx == 0 )\n" |
| "{\n" |
| "float s = lpartsum[nr*nparts];\n" |
| "for( i = 1; i < nparts; i++ )\n" |
| "s += lpartsum[i + nr*nparts];\n" |
| "if( s >= stages[stageIdx].threshold )\n" |
| "{\n" |
| "int count = atomic_inc(lcount);\n" |
| "lbuf[count] = idxval;\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if( stageIdx == N_STAGES )\n" |
| "{\n" |
| "int nrects = lcount[0];\n" |
| "if( lidx < nrects )\n" |
| "{\n" |
| "int nfaces = atomic_inc(facepos);\n" |
| "if( nfaces < MAX_FACES )\n" |
| "{\n" |
| "volatile __global int* face = facepos + 1 + nfaces*3;\n" |
| "int val = lbuf[lidx];\n" |
| "face[0] = scaleIdx;\n" |
| "face[1] = ix0 + (val & 255);\n" |
| "face[2] = iy0 + (val >> 8);\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "}\n" |
| "#endif\n" |
| , "06c037755dc15a7796c0f3bc7ff05233"}; |
| ProgramSource cascadedetect_oclsrc(cascadedetect.programStr); |
| const struct ProgramEntry objdetect_hog={"objdetect_hog", |
| "#define CELL_WIDTH 8\n" |
| "#define CELL_HEIGHT 8\n" |
| "#define CELLS_PER_BLOCK_X 2\n" |
| "#define CELLS_PER_BLOCK_Y 2\n" |
| "#define NTHREADS 256\n" |
| "#define CV_PI_F M_PI_F\n" |
| "#ifdef INTEL_DEVICE\n" |
| "#define QANGLE_TYPE int\n" |
| "#define QANGLE_TYPE2 int2\n" |
| "#else\n" |
| "#define QANGLE_TYPE uchar\n" |
| "#define QANGLE_TYPE2 uchar2\n" |
| "#endif\n" |
| "__kernel void compute_hists_lut_kernel(\n" |
| "const int cblock_stride_x, const int cblock_stride_y,\n" |
| "const int cnbins, const int cblock_hist_size, const int img_block_width,\n" |
| "const int blocks_in_group, const int blocks_total,\n" |
| "const int grad_quadstep, const int qangle_step,\n" |
| "__global const float* grad, __global const QANGLE_TYPE* qangle,\n" |
| "__global const float* gauss_w_lut,\n" |
| "__global float* block_hists, __local float* smem)\n" |
| "{\n" |
| "const int lx = get_local_id(0);\n" |
| "const int lp = lx / 24; \n" |
| "const int gid = get_group_id(0) * blocks_in_group + lp;\n" |
| "const int gidY = gid / img_block_width;\n" |
| "const int gidX = gid - gidY * img_block_width;\n" |
| "const int lidX = lx - lp * 24;\n" |
| "const int lidY = get_local_id(1);\n" |
| "const int cell_x = lidX / 12;\n" |
| "const int cell_y = lidY;\n" |
| "const int cell_thread_x = lidX - cell_x * 12;\n" |
| "__local float* hists = smem + lp * cnbins * (CELLS_PER_BLOCK_X *\n" |
| "CELLS_PER_BLOCK_Y * 12 + CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y);\n" |
| "__local float* final_hist = hists + cnbins *\n" |
| "(CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y * 12);\n" |
| "const int offset_x = gidX * cblock_stride_x + (cell_x << 2) + cell_thread_x;\n" |
| "const int offset_y = gidY * cblock_stride_y + (cell_y << 2);\n" |
| "__global const float* grad_ptr = (gid < blocks_total) ?\n" |
| "grad + offset_y * grad_quadstep + (offset_x << 1) : grad;\n" |
| "__global const QANGLE_TYPE* qangle_ptr = (gid < blocks_total) ?\n" |
| "qangle + offset_y * qangle_step + (offset_x << 1) : qangle;\n" |
| "__local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) +\n" |
| "cell_thread_x;\n" |
| "for (int bin_id = 0; bin_id < cnbins; ++bin_id)\n" |
| "hist[bin_id * 48] = 0.f;\n" |
| "const int dist_x = -4 + cell_thread_x - 4 * cell_x;\n" |
| "const int dist_center_x = dist_x - 4 * (1 - 2 * cell_x);\n" |
| "const int dist_y_begin = -4 - 4 * lidY;\n" |
| "for (int dist_y = dist_y_begin; dist_y < dist_y_begin + 12; ++dist_y)\n" |
| "{\n" |
| "float2 vote = (float2) (grad_ptr[0], grad_ptr[1]);\n" |
| "QANGLE_TYPE2 bin = (QANGLE_TYPE2) (qangle_ptr[0], qangle_ptr[1]);\n" |
| "grad_ptr += grad_quadstep;\n" |
| "qangle_ptr += qangle_step;\n" |
| "int dist_center_y = dist_y - 4 * (1 - 2 * cell_y);\n" |
| "int idx = (dist_center_y + 8) * 16 + (dist_center_x + 8);\n" |
| "float gaussian = gauss_w_lut[idx];\n" |
| "idx = (dist_y + 8) * 16 + (dist_x + 8);\n" |
| "float interp_weight = gauss_w_lut[256+idx];\n" |
| "hist[bin.x * 48] += gaussian * interp_weight * vote.x;\n" |
| "hist[bin.y * 48] += gaussian * interp_weight * vote.y;\n" |
| "}\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "volatile __local float* hist_ = hist;\n" |
| "for (int bin_id = 0; bin_id < cnbins; ++bin_id, hist_ += 48)\n" |
| "{\n" |
| "if (cell_thread_x < 6)\n" |
| "hist_[0] += hist_[6];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (cell_thread_x < 3)\n" |
| "hist_[0] += hist_[3];\n" |
| "#ifdef CPU\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "#endif\n" |
| "if (cell_thread_x == 0)\n" |
| "final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] =\n" |
| "hist_[0] + hist_[1] + hist_[2];\n" |
| "}\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 12 + cell_thread_x;\n" |
| "if ((tid < cblock_hist_size) && (gid < blocks_total))\n" |
| "{\n" |
| "__global float* block_hist = block_hists +\n" |
| "(gidY * img_block_width + gidX) * cblock_hist_size;\n" |
| "block_hist[tid] = final_hist[tid];\n" |
| "}\n" |
| "}\n" |
| "__kernel void normalize_hists_36_kernel(__global float* block_hists,\n" |
| "const float threshold, __local float *squares)\n" |
| "{\n" |
| "const int tid = get_local_id(0);\n" |
| "const int gid = get_global_id(0);\n" |
| "const int bid = tid / 36; \n" |
| "const int boffset = bid * 36; \n" |
| "const int hid = tid - boffset; \n" |
| "float elem = block_hists[gid];\n" |
| "squares[tid] = elem * elem;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "__local float* smem = squares + boffset;\n" |
| "float sum = smem[hid];\n" |
| "if (hid < 18)\n" |
| "smem[hid] = sum = sum + smem[hid + 18];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (hid < 9)\n" |
| "smem[hid] = sum = sum + smem[hid + 9];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (hid < 4)\n" |
| "smem[hid] = sum + smem[hid + 4];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "sum = smem[0] + smem[1] + smem[2] + smem[3] + smem[8];\n" |
| "elem = elem / (sqrt(sum) + 3.6f);\n" |
| "elem = min(elem, threshold);\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "squares[tid] = elem * elem;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "sum = smem[hid];\n" |
| "if (hid < 18)\n" |
| "smem[hid] = sum = sum + smem[hid + 18];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (hid < 9)\n" |
| "smem[hid] = sum = sum + smem[hid + 9];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (hid < 4)\n" |
| "smem[hid] = sum + smem[hid + 4];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "sum = smem[0] + smem[1] + smem[2] + smem[3] + smem[8];\n" |
| "block_hists[gid] = elem / (sqrt(sum) + 1e-3f);\n" |
| "}\n" |
| "inline float reduce_smem(volatile __local float* smem, int size)\n" |
| "{\n" |
| "unsigned int tid = get_local_id(0);\n" |
| "float sum = smem[tid];\n" |
| "if (size >= 512) { if (tid < 256) smem[tid] = sum = sum + smem[tid + 256];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "if (size >= 256) { if (tid < 128) smem[tid] = sum = sum + smem[tid + 128];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "if (size >= 128) { if (tid < 64) smem[tid] = sum = sum + smem[tid + 64];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "#ifdef CPU\n" |
| "if (size >= 64) { if (tid < 32) smem[tid] = sum = sum + smem[tid + 32];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "if (size >= 32) { if (tid < 16) smem[tid] = sum = sum + smem[tid + 16];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "if (size >= 16) { if (tid < 8) smem[tid] = sum = sum + smem[tid + 8];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "if (size >= 8) { if (tid < 4) smem[tid] = sum = sum + smem[tid + 4];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "if (size >= 4) { if (tid < 2) smem[tid] = sum = sum + smem[tid + 2];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "if (size >= 2) { if (tid < 1) smem[tid] = sum = sum + smem[tid + 1];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE); }\n" |
| "#else\n" |
| "if (tid < 32)\n" |
| "{\n" |
| "if (size >= 64) smem[tid] = sum = sum + smem[tid + 32];\n" |
| "#if WAVE_SIZE < 32\n" |
| "} barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 16) {\n" |
| "#endif\n" |
| "if (size >= 32) smem[tid] = sum = sum + smem[tid + 16];\n" |
| "if (size >= 16) smem[tid] = sum = sum + smem[tid + 8];\n" |
| "if (size >= 8) smem[tid] = sum = sum + smem[tid + 4];\n" |
| "if (size >= 4) smem[tid] = sum = sum + smem[tid + 2];\n" |
| "if (size >= 2) smem[tid] = sum = sum + smem[tid + 1];\n" |
| "}\n" |
| "#endif\n" |
| "return sum;\n" |
| "}\n" |
| "__kernel void normalize_hists_kernel(\n" |
| "const int nthreads, const int block_hist_size, const int img_block_width,\n" |
| "__global float* block_hists, const float threshold, __local float *squares)\n" |
| "{\n" |
| "const int tid = get_local_id(0);\n" |
| "const int gidX = get_group_id(0);\n" |
| "const int gidY = get_group_id(1);\n" |
| "__global float* hist = block_hists + (gidY * img_block_width + gidX) *\n" |
| "block_hist_size + tid;\n" |
| "float elem = 0.f;\n" |
| "if (tid < block_hist_size)\n" |
| "elem = hist[0];\n" |
| "squares[tid] = elem * elem;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "float sum = reduce_smem(squares, nthreads);\n" |
| "float scale = 1.0f / (sqrt(sum) + 0.1f * block_hist_size);\n" |
| "elem = min(elem * scale, threshold);\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "squares[tid] = elem * elem;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "sum = reduce_smem(squares, nthreads);\n" |
| "scale = 1.0f / (sqrt(sum) + 1e-3f);\n" |
| "if (tid < block_hist_size)\n" |
| "hist[0] = elem * scale;\n" |
| "}\n" |
| "__kernel void classify_hists_180_kernel(\n" |
| "const int cdescr_width, const int cdescr_height, const int cblock_hist_size,\n" |
| "const int img_win_width, const int img_block_width,\n" |
| "const int win_block_stride_x, const int win_block_stride_y,\n" |
| "__global const float * block_hists, __global const float* coefs,\n" |
| "float free_coef, float threshold, __global uchar* labels)\n" |
| "{\n" |
| "const int tid = get_local_id(0);\n" |
| "const int gidX = get_group_id(0);\n" |
| "const int gidY = get_group_id(1);\n" |
| "__global const float* hist = block_hists + (gidY * win_block_stride_y *\n" |
| "img_block_width + gidX * win_block_stride_x) * cblock_hist_size;\n" |
| "float product = 0.f;\n" |
| "for (int i = 0; i < cdescr_height; i++)\n" |
| "{\n" |
| "product += coefs[i * cdescr_width + tid] *\n" |
| "hist[i * img_block_width * cblock_hist_size + tid];\n" |
| "}\n" |
| "__local float products[180];\n" |
| "products[tid] = product;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 90) products[tid] = product = product + products[tid + 90];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 45) products[tid] = product = product + products[tid + 45];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "volatile __local float* smem = products;\n" |
| "#ifdef CPU\n" |
| "if (tid < 13) smem[tid] = product = product + smem[tid + 32];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 16) smem[tid] = product = product + smem[tid + 16];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<8) smem[tid] = product = product + smem[tid + 8];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<4) smem[tid] = product = product + smem[tid + 4];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<2) smem[tid] = product = product + smem[tid + 2];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "#else\n" |
| "if (tid < 13)\n" |
| "{\n" |
| "smem[tid] = product = product + smem[tid + 32];\n" |
| "}\n" |
| "#if WAVE_SIZE < 32\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "#endif\n" |
| "if (tid < 16)\n" |
| "{\n" |
| "smem[tid] = product = product + smem[tid + 16];\n" |
| "smem[tid] = product = product + smem[tid + 8];\n" |
| "smem[tid] = product = product + smem[tid + 4];\n" |
| "smem[tid] = product = product + smem[tid + 2];\n" |
| "}\n" |
| "#endif\n" |
| "if (tid == 0){\n" |
| "product = product + smem[tid + 1];\n" |
| "labels[gidY * img_win_width + gidX] = (product + free_coef >= threshold);\n" |
| "}\n" |
| "}\n" |
| "__kernel void classify_hists_252_kernel(\n" |
| "const int cdescr_width, const int cdescr_height, const int cblock_hist_size,\n" |
| "const int img_win_width, const int img_block_width,\n" |
| "const int win_block_stride_x, const int win_block_stride_y,\n" |
| "__global const float * block_hists, __global const float* coefs,\n" |
| "float free_coef, float threshold, __global uchar* labels)\n" |
| "{\n" |
| "const int tid = get_local_id(0);\n" |
| "const int gidX = get_group_id(0);\n" |
| "const int gidY = get_group_id(1);\n" |
| "__global const float* hist = block_hists + (gidY * win_block_stride_y *\n" |
| "img_block_width + gidX * win_block_stride_x) * cblock_hist_size;\n" |
| "float product = 0.f;\n" |
| "if (tid < cdescr_width)\n" |
| "{\n" |
| "for (int i = 0; i < cdescr_height; i++)\n" |
| "product += coefs[i * cdescr_width + tid] *\n" |
| "hist[i * img_block_width * cblock_hist_size + tid];\n" |
| "}\n" |
| "__local float products[NTHREADS];\n" |
| "products[tid] = product;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 128) products[tid] = product = product + products[tid + 128];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 64) products[tid] = product = product + products[tid + 64];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "volatile __local float* smem = products;\n" |
| "#ifdef CPU\n" |
| "if(tid<32) smem[tid] = product = product + smem[tid + 32];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<16) smem[tid] = product = product + smem[tid + 16];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<8) smem[tid] = product = product + smem[tid + 8];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<4) smem[tid] = product = product + smem[tid + 4];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<2) smem[tid] = product = product + smem[tid + 2];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "#else\n" |
| "if (tid < 32)\n" |
| "{\n" |
| "smem[tid] = product = product + smem[tid + 32];\n" |
| "#if WAVE_SIZE < 32\n" |
| "} barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 16) {\n" |
| "#endif\n" |
| "smem[tid] = product = product + smem[tid + 16];\n" |
| "smem[tid] = product = product + smem[tid + 8];\n" |
| "smem[tid] = product = product + smem[tid + 4];\n" |
| "smem[tid] = product = product + smem[tid + 2];\n" |
| "}\n" |
| "#endif\n" |
| "if (tid == 0){\n" |
| "product = product + smem[tid + 1];\n" |
| "labels[gidY * img_win_width + gidX] = (product + free_coef >= threshold);\n" |
| "}\n" |
| "}\n" |
| "__kernel void classify_hists_kernel(\n" |
| "const int cdescr_size, const int cdescr_width, const int cblock_hist_size,\n" |
| "const int img_win_width, const int img_block_width,\n" |
| "const int win_block_stride_x, const int win_block_stride_y,\n" |
| "__global const float * block_hists, __global const float* coefs,\n" |
| "float free_coef, float threshold, __global uchar* labels)\n" |
| "{\n" |
| "const int tid = get_local_id(0);\n" |
| "const int gidX = get_group_id(0);\n" |
| "const int gidY = get_group_id(1);\n" |
| "__global const float* hist = block_hists + (gidY * win_block_stride_y *\n" |
| "img_block_width + gidX * win_block_stride_x) * cblock_hist_size;\n" |
| "float product = 0.f;\n" |
| "for (int i = tid; i < cdescr_size; i += NTHREADS)\n" |
| "{\n" |
| "int offset_y = i / cdescr_width;\n" |
| "int offset_x = i - offset_y * cdescr_width;\n" |
| "product += coefs[i] *\n" |
| "hist[offset_y * img_block_width * cblock_hist_size + offset_x];\n" |
| "}\n" |
| "__local float products[NTHREADS];\n" |
| "products[tid] = product;\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 128) products[tid] = product = product + products[tid + 128];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 64) products[tid] = product = product + products[tid + 64];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "volatile __local float* smem = products;\n" |
| "#ifdef CPU\n" |
| "if(tid<32) smem[tid] = product = product + smem[tid + 32];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<16) smem[tid] = product = product + smem[tid + 16];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<8) smem[tid] = product = product + smem[tid + 8];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<4) smem[tid] = product = product + smem[tid + 4];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if(tid<2) smem[tid] = product = product + smem[tid + 2];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "#else\n" |
| "if (tid < 32)\n" |
| "{\n" |
| "smem[tid] = product = product + smem[tid + 32];\n" |
| "#if WAVE_SIZE < 32\n" |
| "} barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (tid < 16) {\n" |
| "#endif\n" |
| "smem[tid] = product = product + smem[tid + 16];\n" |
| "smem[tid] = product = product + smem[tid + 8];\n" |
| "smem[tid] = product = product + smem[tid + 4];\n" |
| "smem[tid] = product = product + smem[tid + 2];\n" |
| "}\n" |
| "#endif\n" |
| "if (tid == 0){\n" |
| "smem[tid] = product = product + smem[tid + 1];\n" |
| "labels[gidY * img_win_width + gidX] = (product + free_coef >= threshold);\n" |
| "}\n" |
| "}\n" |
| "__kernel void extract_descrs_by_rows_kernel(\n" |
| "const int cblock_hist_size, const int descriptors_quadstep,\n" |
| "const int cdescr_size, const int cdescr_width, const int img_block_width,\n" |
| "const int win_block_stride_x, const int win_block_stride_y,\n" |
| "__global const float* block_hists, __global float* descriptors)\n" |
| "{\n" |
| "int tid = get_local_id(0);\n" |
| "int gidX = get_group_id(0);\n" |
| "int gidY = get_group_id(1);\n" |
| "__global const float* hist = block_hists + (gidY * win_block_stride_y *\n" |
| "img_block_width + gidX * win_block_stride_x) * cblock_hist_size;\n" |
| "__global float* descriptor = descriptors +\n" |
| "(gidY * get_num_groups(0) + gidX) * descriptors_quadstep;\n" |
| "for (int i = tid; i < cdescr_size; i += NTHREADS)\n" |
| "{\n" |
| "int offset_y = i / cdescr_width;\n" |
| "int offset_x = i - offset_y * cdescr_width;\n" |
| "descriptor[i] = hist[offset_y * img_block_width * cblock_hist_size + offset_x];\n" |
| "}\n" |
| "}\n" |
| "__kernel void extract_descrs_by_cols_kernel(\n" |
| "const int cblock_hist_size, const int descriptors_quadstep, const int cdescr_size,\n" |
| "const int cnblocks_win_x, const int cnblocks_win_y, const int img_block_width,\n" |
| "const int win_block_stride_x, const int win_block_stride_y,\n" |
| "__global const float* block_hists, __global float* descriptors)\n" |
| "{\n" |
| "int tid = get_local_id(0);\n" |
| "int gidX = get_group_id(0);\n" |
| "int gidY = get_group_id(1);\n" |
| "__global const float* hist = block_hists + (gidY * win_block_stride_y *\n" |
| "img_block_width + gidX * win_block_stride_x) * cblock_hist_size;\n" |
| "__global float* descriptor = descriptors +\n" |
| "(gidY * get_num_groups(0) + gidX) * descriptors_quadstep;\n" |
| "for (int i = tid; i < cdescr_size; i += NTHREADS)\n" |
| "{\n" |
| "int block_idx = i / cblock_hist_size;\n" |
| "int idx_in_block = i - block_idx * cblock_hist_size;\n" |
| "int y = block_idx / cnblocks_win_x;\n" |
| "int x = block_idx - y * cnblocks_win_x;\n" |
| "descriptor[(x * cnblocks_win_y + y) * cblock_hist_size + idx_in_block] =\n" |
| "hist[(y * img_block_width + x) * cblock_hist_size + idx_in_block];\n" |
| "}\n" |
| "}\n" |
| "__kernel void compute_gradients_8UC4_kernel(\n" |
| "const int height, const int width,\n" |
| "const int img_step, const int grad_quadstep, const int qangle_step,\n" |
| "const __global uchar4 * img, __global float * grad, __global QANGLE_TYPE * qangle,\n" |
| "const float angle_scale, const char correct_gamma, const int cnbins)\n" |
| "{\n" |
| "const int x = get_global_id(0);\n" |
| "const int tid = get_local_id(0);\n" |
| "const int gSizeX = get_local_size(0);\n" |
| "const int gidY = get_group_id(1);\n" |
| "__global const uchar4* row = img + gidY * img_step;\n" |
| "__local float sh_row[(NTHREADS + 2) * 3];\n" |
| "uchar4 val;\n" |
| "if (x < width)\n" |
| "val = row[x];\n" |
| "else\n" |
| "val = row[width - 2];\n" |
| "sh_row[tid + 1] = val.x;\n" |
| "sh_row[tid + 1 + (NTHREADS + 2)] = val.y;\n" |
| "sh_row[tid + 1 + 2 * (NTHREADS + 2)] = val.z;\n" |
| "if (tid == 0)\n" |
| "{\n" |
| "val = row[max(x - 1, 1)];\n" |
| "sh_row[0] = val.x;\n" |
| "sh_row[(NTHREADS + 2)] = val.y;\n" |
| "sh_row[2 * (NTHREADS + 2)] = val.z;\n" |
| "}\n" |
| "if (tid == gSizeX - 1)\n" |
| "{\n" |
| "val = row[min(x + 1, width - 2)];\n" |
| "sh_row[gSizeX + 1] = val.x;\n" |
| "sh_row[gSizeX + 1 + (NTHREADS + 2)] = val.y;\n" |
| "sh_row[gSizeX + 1 + 2 * (NTHREADS + 2)] = val.z;\n" |
| "}\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (x < width)\n" |
| "{\n" |
| "float4 a = (float4) (sh_row[tid], sh_row[tid + (NTHREADS + 2)],\n" |
| "sh_row[tid + 2 * (NTHREADS + 2)], 0);\n" |
| "float4 b = (float4) (sh_row[tid + 2], sh_row[tid + 2 + (NTHREADS + 2)],\n" |
| "sh_row[tid + 2 + 2 * (NTHREADS + 2)], 0);\n" |
| "float4 dx;\n" |
| "if (correct_gamma == 1)\n" |
| "dx = sqrt(b) - sqrt(a);\n" |
| "else\n" |
| "dx = b - a;\n" |
| "float4 dy = (float4) 0.f;\n" |
| "if (gidY > 0 && gidY < height - 1)\n" |
| "{\n" |
| "a = convert_float4(img[(gidY - 1) * img_step + x].xyzw);\n" |
| "b = convert_float4(img[(gidY + 1) * img_step + x].xyzw);\n" |
| "if (correct_gamma == 1)\n" |
| "dy = sqrt(b) - sqrt(a);\n" |
| "else\n" |
| "dy = b - a;\n" |
| "}\n" |
| "float4 mag = hypot(dx, dy);\n" |
| "float best_dx = dx.x;\n" |
| "float best_dy = dy.x;\n" |
| "float mag0 = mag.x;\n" |
| "if (mag0 < mag.y)\n" |
| "{\n" |
| "best_dx = dx.y;\n" |
| "best_dy = dy.y;\n" |
| "mag0 = mag.y;\n" |
| "}\n" |
| "if (mag0 < mag.z)\n" |
| "{\n" |
| "best_dx = dx.z;\n" |
| "best_dy = dy.z;\n" |
| "mag0 = mag.z;\n" |
| "}\n" |
| "float ang = (atan2(best_dy, best_dx) + CV_PI_F) * angle_scale - 0.5f;\n" |
| "int hidx = (int)floor(ang);\n" |
| "ang -= hidx;\n" |
| "hidx = (hidx + cnbins) % cnbins;\n" |
| "qangle[(gidY * qangle_step + x) << 1] = hidx;\n" |
| "qangle[((gidY * qangle_step + x) << 1) + 1] = (hidx + 1) % cnbins;\n" |
| "grad[(gidY * grad_quadstep + x) << 1] = mag0 * (1.f - ang);\n" |
| "grad[((gidY * grad_quadstep + x) << 1) + 1] = mag0 * ang;\n" |
| "}\n" |
| "}\n" |
| "__kernel void compute_gradients_8UC1_kernel(\n" |
| "const int height, const int width,\n" |
| "const int img_step, const int grad_quadstep, const int qangle_step,\n" |
| "__global const uchar * img, __global float * grad, __global QANGLE_TYPE * qangle,\n" |
| "const float angle_scale, const char correct_gamma, const int cnbins)\n" |
| "{\n" |
| "const int x = get_global_id(0);\n" |
| "const int tid = get_local_id(0);\n" |
| "const int gSizeX = get_local_size(0);\n" |
| "const int gidY = get_group_id(1);\n" |
| "__global const uchar* row = img + gidY * img_step;\n" |
| "__local float sh_row[NTHREADS + 2];\n" |
| "if (x < width)\n" |
| "sh_row[tid + 1] = row[x];\n" |
| "else\n" |
| "sh_row[tid + 1] = row[width - 2];\n" |
| "if (tid == 0)\n" |
| "sh_row[0] = row[max(x - 1, 1)];\n" |
| "if (tid == gSizeX - 1)\n" |
| "sh_row[gSizeX + 1] = row[min(x + 1, width - 2)];\n" |
| "barrier(CLK_LOCAL_MEM_FENCE);\n" |
| "if (x < width)\n" |
| "{\n" |
| "float dx;\n" |
| "if (correct_gamma == 1)\n" |
| "dx = sqrt(sh_row[tid + 2]) - sqrt(sh_row[tid]);\n" |
| "else\n" |
| "dx = sh_row[tid + 2] - sh_row[tid];\n" |
| "float dy = 0.f;\n" |
| "if (gidY > 0 && gidY < height - 1)\n" |
| "{\n" |
| "float a = (float) img[ (gidY + 1) * img_step + x ];\n" |
| "float b = (float) img[ (gidY - 1) * img_step + x ];\n" |
| "if (correct_gamma == 1)\n" |
| "dy = sqrt(a) - sqrt(b);\n" |
| "else\n" |
| "dy = a - b;\n" |
| "}\n" |
| "float mag = hypot(dx, dy);\n" |
| "float ang = (atan2(dy, dx) + CV_PI_F) * angle_scale - 0.5f;\n" |
| "int hidx = (int)floor(ang);\n" |
| "ang -= hidx;\n" |
| "hidx = (hidx + cnbins) % cnbins;\n" |
| "qangle[ (gidY * qangle_step + x) << 1 ] = hidx;\n" |
| "qangle[ ((gidY * qangle_step + x) << 1) + 1 ] = (hidx + 1) % cnbins;\n" |
| "grad[ (gidY * grad_quadstep + x) << 1 ] = mag * (1.f - ang);\n" |
| "grad[ ((gidY * grad_quadstep + x) << 1) + 1 ] = mag * ang;\n" |
| "}\n" |
| "}\n" |
| , "bc57f4f75fb81bae73bfe73cc4ca15e4"}; |
| ProgramSource objdetect_hog_oclsrc(objdetect_hog.programStr); |
| } |
| }} |
