generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Enabled Intel-specific optimizations for HOG detector.

Browse Source
This commit is contained in:
Alexander Karsakov 2014-03-19 17:30:13 +04:00
parent dc28451b8a
commit 30a8308f8e
2 changed files with 22 additions and 12 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
modules/objdetect/src/hog.cpp
View File

@ -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<float>& _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<Point> &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);

18
modules/objdetect/src/opencl/objdetect_hog.cl
View File

@ -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);