From 30a8308f8e66fd46483cb193cac66665666a77fd Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Wed, 19 Mar 2014 17:30:13 +0400 Subject: [PATCH] Enabled Intel-specific optimizations for HOG detector. --- modules/objdetect/src/hog.cpp | 16 +++++++++------- modules/objdetect/src/opencl/objdetect_hog.cl | 18 +++++++++++++----- 2 files changed, 22 insertions(+), 12 deletions(-) diff --git a/modules/objdetect/src/hog.cpp b/modules/objdetect/src/hog.cpp index 18bb7af..0f4456a 100644 --- a/modules/objdetect/src/hog.cpp +++ b/modules/objdetect/src/hog.cpp @@ -1085,8 +1085,8 @@ static bool ocl_compute_gradients_8UC1(int height, int width, InputArray _img, f size_t globalThreads[3] = { width, height, 1 }; char correctGamma = (correct_gamma) ? 1 : 0; int grad_quadstep = (int)grad.step >> 3; - int qangle_step_shift = 0; - int qangle_step = (int)qangle.step >> (1 + qangle_step_shift); + int qangle_elem_size = CV_ELEM_SIZE1(qangle.type()); + int qangle_step = (int)qangle.step / (2 * qangle_elem_size); int idx = 0; idx = k.set(idx, height); @@ -1137,9 +1137,9 @@ static bool ocl_compute_hists(int nbins, int block_stride_x, int block_stride_y, int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y)/block_stride_y; int blocks_total = img_block_width * img_block_height; - int qangle_step_shift = 0; + int qangle_elem_size = CV_ELEM_SIZE1(qangle.type()); int grad_quadstep = (int)grad.step >> 2; - int qangle_step = (int)qangle.step >> qangle_step_shift; + int qangle_step = (int)qangle.step / qangle_elem_size; int blocks_in_group = 4; size_t localThreads[3] = { blocks_in_group * 24, 2, 1 }; @@ -1316,11 +1316,12 @@ static bool ocl_extract_descrs_by_cols(int win_height, int win_width, int block_ static bool ocl_compute(InputArray _img, Size win_stride, std::vector& _descriptors, int descr_format, Size blockSize, Size cellSize, int nbins, Size blockStride, Size winSize, float sigma, bool gammaCorrection, double L2HysThreshold) { - Size imgSize = _img.size(); + Size imgSize = _img.size(); Size effect_size = imgSize; UMat grad(imgSize, CV_32FC2); - UMat qangle(imgSize, CV_8UC2); + int qangle_type = ocl::Device::getDefault().isIntel() ? CV_32SC2 : CV_8UC2; + UMat qangle(imgSize, qangle_type); const size_t block_hist_size = getBlockHistogramSize(blockSize, cellSize, nbins); const Size blocks_per_img = numPartsWithin(imgSize, blockSize, blockStride); @@ -1720,7 +1721,8 @@ static bool ocl_detect(InputArray img, std::vector &hits, double hit_thre Size imgSize = img.size(); Size effect_size = imgSize; UMat grad(imgSize, CV_32FC2); - UMat qangle(imgSize, CV_8UC2); + int qangle_type = ocl::Device::getDefault().isIntel() ? CV_32SC2 : CV_8UC2; + UMat qangle(imgSize, qangle_type); const size_t block_hist_size = getBlockHistogramSize(blockSize, cellSize, nbins); const Size blocks_per_img = numPartsWithin(imgSize, blockSize, blockStride); diff --git a/modules/objdetect/src/opencl/objdetect_hog.cl b/modules/objdetect/src/opencl/objdetect_hog.cl index 082f9ab..5c71aa1 100644 --- a/modules/objdetect/src/opencl/objdetect_hog.cl +++ b/modules/objdetect/src/opencl/objdetect_hog.cl @@ -50,6 +50,14 @@ #define NTHREADS 256 #define CV_PI_F 3.1415926535897932384626433832795f +#ifdef INTEL_DEVICE +#define QANGLE_TYPE int +#define QANGLE_TYPE2 int2 +#else +#define QANGLE_TYPE uchar +#define QANGLE_TYPE2 uchar2 +#endif + //---------------------------------------------------------------------------- // Histogram computation // 12 threads for a cell, 12x4 threads per block @@ -59,7 +67,7 @@ __kernel void compute_hists_lut_kernel( 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 uchar* qangle, + __global const float* grad, __global const QANGLE_TYPE* qangle, __global const float* gauss_w_lut, __global float* block_hists, __local float* smem) { @@ -86,7 +94,7 @@ __kernel void compute_hists_lut_kernel( __global const float* grad_ptr = (gid < blocks_total) ? grad + offset_y * grad_quadstep + (offset_x << 1) : grad; - __global const uchar* qangle_ptr = (gid < blocks_total) ? + __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) + @@ -101,7 +109,7 @@ __kernel void compute_hists_lut_kernel( for (int dist_y = dist_y_begin; dist_y < dist_y_begin + 12; ++dist_y) { float2 vote = (float2) (grad_ptr[0], grad_ptr[1]); - uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]); + QANGLE_TYPE2 bin = (QANGLE_TYPE2) (qangle_ptr[0], qangle_ptr[1]); grad_ptr += grad_quadstep; qangle_ptr += qangle_step; @@ -558,7 +566,7 @@ __kernel void extract_descrs_by_cols_kernel( __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 uchar * qangle, + 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); @@ -660,7 +668,7 @@ __kernel void compute_gradients_8UC4_kernel( __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 uchar * qangle, + __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); -- 2.7.4