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