//M*/
#include "precomp.hpp"
-
using namespace cv;
using namespace cv::ocl;
using namespace std;
}
}
-
void cv::ocl::HOGDescriptor::computeBlockHistograms(const oclMat &img)
{
computeGradient(img, grad, qangle);
cdescr_size = descr_size;
}
+static inline int divUp(int total, int grain)
+{
+ return (total + grain - 1) / grain;
+}
+
+static void openCLExecuteKernel_hog(Context *clCxt , const char **source, string kernelName,
+ size_t globalThreads[3], size_t localThreads[3],
+ vector< pair<size_t, const void *> > &args)
+{
+ size_t wave_size = 0;
+ queryDeviceInfo(WAVEFRONT_SIZE, &wave_size);
+ if (wave_size <= 16)
+ {
+ char build_options[64];
+ sprintf(build_options, (wave_size == 16) ? "-D WAVE_SIZE_16" : "-D WAVE_SIZE_1");
+ openCLExecuteKernel(clCxt, source, kernelName, globalThreads, localThreads, args, -1, -1, build_options);
+ }
+ else
+ openCLExecuteKernel(clCxt, source, kernelName, globalThreads, localThreads, args, -1, -1);
+}
+
void cv::ocl::device::hog::compute_hists(int nbins, int block_stride_x, int block_stride_y,
int height, int width, const cv::ocl::oclMat &grad,
const cv::ocl::oclMat &qangle, float sigma, cv::ocl::oclMat &block_hists)
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) / block_stride_x;
int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y) / block_stride_y;
- size_t globalThreads[3] = { img_block_width * 32, img_block_height * 2, 1 };
- size_t localThreads[3] = { 32, 2, 1 };
+ int blocks_total = img_block_width * img_block_height;
+ int blocks_in_group = 4;
+ size_t localThreads[3] = { blocks_in_group * 24, 2, 1 };
+ size_t globalThreads[3] = { divUp(blocks_total, blocks_in_group) * localThreads[0], 2, 1 };
int grad_quadstep = grad.step >> 2;
int qangle_step = qangle.step;
int hists_size = (nbins * CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y * 12) * sizeof(float);
int final_hists_size = (nbins * CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y) * sizeof(float);
- int smem = hists_size + final_hists_size;
+ int smem = (hists_size + final_hists_size) * blocks_in_group;
- args.push_back( make_pair( sizeof(cl_int), (void *)&width));
args.push_back( make_pair( sizeof(cl_int), (void *)&cblock_stride_x));
args.push_back( make_pair( sizeof(cl_int), (void *)&cblock_stride_y));
args.push_back( make_pair( sizeof(cl_int), (void *)&cnbins));
args.push_back( make_pair( sizeof(cl_int), (void *)&cblock_hist_size));
args.push_back( make_pair( sizeof(cl_int), (void *)&img_block_width));
+ args.push_back( make_pair( sizeof(cl_int), (void *)&blocks_in_group));
+ args.push_back( make_pair( sizeof(cl_int), (void *)&blocks_total));
args.push_back( make_pair( sizeof(cl_int), (void *)&grad_quadstep));
args.push_back( make_pair( sizeof(cl_int), (void *)&qangle_step));
args.push_back( make_pair( sizeof(cl_mem), (void *)&grad.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&block_hists.data));
args.push_back( make_pair( smem, (void *)NULL));
- openCLExecuteKernel2(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
+ openCLExecuteKernel_hog(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args);
}
void cv::ocl::device::hog::normalize_hists(int nbins, int block_stride_x, int block_stride_y,
args.push_back( make_pair( sizeof(cl_float), (void *)&threshold));
args.push_back( make_pair( nthreads * sizeof(float), (void *)NULL));
- openCLExecuteKernel2(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
+ openCLExecuteKernel_hog(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args);
}
void cv::ocl::device::hog::classify_hists(int win_height, int win_width, int block_stride_y,
args.push_back( make_pair( sizeof(cl_float), (void *)&threshold));
args.push_back( make_pair( sizeof(cl_mem), (void *)&labels.data));
- openCLExecuteKernel2(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
+ openCLExecuteKernel_hog(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args);
}
void cv::ocl::device::hog::extract_descrs_by_rows(int win_height, int win_width, int block_stride_y, int block_stride_x,
args.push_back( make_pair( sizeof(cl_mem), (void *)&block_hists.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&descriptors.data));
- openCLExecuteKernel2(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
+ openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
}
void cv::ocl::device::hog::extract_descrs_by_cols(int win_height, int win_width, int block_stride_y, int block_stride_x,
args.push_back( make_pair( sizeof(cl_mem), (void *)&block_hists.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&descriptors.data));
- openCLExecuteKernel2(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
-}
-
-static inline int divUp(int total, int grain)
-{
- return (total + grain - 1) / grain;
+ openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
}
void cv::ocl::device::hog::compute_gradients_8UC1(int height, int width, const cv::ocl::oclMat &img,
args.push_back( make_pair( sizeof(cl_char), (void *)&correctGamma));
args.push_back( make_pair( sizeof(cl_int), (void *)&cnbins));
- openCLExecuteKernel2(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
+ openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
}
void cv::ocl::device::hog::compute_gradients_8UC4(int height, int width, const cv::ocl::oclMat &img,
args.push_back( make_pair( sizeof(cl_char), (void *)&correctGamma));
args.push_back( make_pair( sizeof(cl_int), (void *)&cnbins));
- openCLExecuteKernel2(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
+ openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
}
void cv::ocl::device::hog::resize( const oclMat &src, oclMat &dst, const Size sz)
float ifx = (float)src.cols / sz.width;
float ify = (float)src.rows / sz.height;
+ int src_step = static_cast<int>(src.step);
+ int dst_step = static_cast<int>(dst.step);
vector< pair<size_t, const void *> > args;
args.push_back( make_pair(sizeof(cl_mem), (void *)&dst.data));
args.push_back( make_pair(sizeof(cl_mem), (void *)&src.data));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.offset));
args.push_back( make_pair(sizeof(cl_int), (void *)&src.offset));
- args.push_back( make_pair(sizeof(cl_int), (void *)&dst.step));
- args.push_back( make_pair(sizeof(cl_int), (void *)&src.step));
+ args.push_back( make_pair(sizeof(cl_int), (void *)&dst_step));
+ args.push_back( make_pair(sizeof(cl_int), (void *)&src_step));
args.push_back( make_pair(sizeof(cl_int), (void *)&src.cols));
args.push_back( make_pair(sizeof(cl_int), (void *)&src.rows));
args.push_back( make_pair(sizeof(cl_int), (void *)&sz.width));
args.push_back( make_pair(sizeof(cl_float), (void *)&ifx));
args.push_back( make_pair(sizeof(cl_float), (void *)&ify));
- openCLExecuteKernel2(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
+ openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, localThreads, args, -1, -1);
}
//----------------------------------------------------------------------------
// Histogram computation
-
-__kernel void compute_hists_kernel(const int width, 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 grad_quadstep, const int qangle_step,
- __global const float* grad, __global const uchar* qangle,
- const float scale, __global float* block_hists, __local float* smem)
+// 12 threads for a cell, 12x4 threads per block
+__kernel void compute_hists_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 uchar* qangle,
+ const float scale, __global float* block_hists, __local float* smem)
{
- const int lidX = get_local_id(0);
+ const int lx = get_local_id(0);
+ const int lp = lx / 24; /* local group id */
+ const int gid = get_group_id(0) * blocks_in_group + lp;/* global group id */
+ 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 gidX = get_group_id(0);
- const int gidY = get_group_id(1);
- const int cell_x = lidX / 16;
+ const int cell_x = lidX / 12;
const int cell_y = lidY;
- const int cell_thread_x = lidX & 0xF;
+ const int cell_thread_x = lidX - cell_x * 12;
- __local float* hists = smem;
- __local float* final_hist = smem + cnbins * 48;
+ __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 = grad + offset_y * grad_quadstep + (offset_x << 1);
- __global const uchar* qangle_ptr = qangle + offset_y * qangle_step + (offset_x << 1);
-
- // 12 means that 12 pixels affect on block's cell (in one row)
- if (cell_thread_x < 12)
- {
- __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;
+ __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) ?
+ qangle + offset_y * qangle_step + (offset_x << 1) : qangle;
- const int dist_x = -4 + cell_thread_x - 4 * cell_x;
+ __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_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]);
- uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]);
+ const int dist_x = -4 + cell_thread_x - 4 * cell_x;
+ const int dist_center_x = dist_x - 4 * (1 - 2 * cell_x);
- grad_ptr += grad_quadstep;
- qangle_ptr += qangle_step;
+ 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]);
+ uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]);
- int dist_center_y = dist_y - 4 * (1 - 2 * cell_y);
- int dist_center_x = dist_x - 4 * (1 - 2 * cell_x);
+ grad_ptr += grad_quadstep;
+ qangle_ptr += qangle_step;
- float gaussian = exp(-(dist_center_y * dist_center_y + dist_center_x * dist_center_x) * scale);
- float interp_weight = (8.f - fabs(dist_y + 0.5f)) * (8.f - fabs(dist_x + 0.5f)) / 64.f;
+ int dist_center_y = dist_y - 4 * (1 - 2 * cell_y);
- hist[bin.x * 48] += gaussian * interp_weight * vote.x;
- hist[bin.y * 48] += gaussian * interp_weight * vote.y;
- }
+ float gaussian = exp(-(dist_center_y * dist_center_y + dist_center_x *
+ dist_center_x) * scale);
+ float interp_weight = (8.f - fabs(dist_y + 0.5f)) *
+ (8.f - fabs(dist_x + 0.5f)) / 64.f;
- 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];
- if (cell_thread_x < 3) hist_[0] += hist_[3];
- if (cell_thread_x == 0)
- final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = hist_[0] + hist_[1] + hist_[2];
- }
+ hist[bin.x * 48] += gaussian * interp_weight * vote.x;
+ hist[bin.y * 48] += gaussian * interp_weight * vote.y;
}
-
barrier(CLK_LOCAL_MEM_FENCE);
- __global float* block_hist = block_hists + (gidY * img_block_width + gidX) * cblock_hist_size;
+ 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];
+#ifdef WAVE_SIZE_1
+ barrier(CLK_LOCAL_MEM_FENCE);
+#endif
+ if (cell_thread_x == 0)
+ final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] =
+ hist_[0] + hist_[1] + hist_[2];
+ }
+#ifdef WAVE_SIZE_1
+ barrier(CLK_LOCAL_MEM_FENCE);
+#endif
- int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 16 + cell_thread_x;
- if (tid < cblock_hist_size)
+ 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];
+ }
}
//-------------------------------------------------------------
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 >= 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 (tid < 32)
{
if (size >= 64) smem[tid] = sum = sum + smem[tid + 32];
+#if defined(WAVE_SIZE_16) || defined(WAVE_SIZE_1)
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
{
+#endif
if (size >= 32) smem[tid] = sum = sum + smem[tid + 16];
+#ifdef WAVE_SIZE_1
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (tid < 8)
+ {
+#endif
if (size >= 16) smem[tid] = sum = sum + smem[tid + 8];
+#ifdef WAVE_SIZE_1
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (tid < 4)
+ {
+#endif
if (size >= 8) smem[tid] = sum = sum + smem[tid + 4];
+#ifdef WAVE_SIZE_1
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (tid < 2)
+ {
+#endif
if (size >= 4) smem[tid] = sum = sum + smem[tid + 2];
+#ifdef WAVE_SIZE_1
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (tid < 1)
+ {
+#endif
if (size >= 2) smem[tid] = sum = sum + smem[tid + 1];
}
if (tid < 64) products[tid] = product = product + products[tid + 64];
barrier(CLK_LOCAL_MEM_FENCE);
+ volatile __local float* smem = products;
if (tid < 32)
{
- volatile __local float* smem = products;
smem[tid] = product = product + smem[tid + 32];
+#if defined(WAVE_SIZE_16) || defined(WAVE_SIZE_1)
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
{
- volatile __local float* smem = products;
+#endif
smem[tid] = product = product + smem[tid + 16];
+#ifdef WAVE_SIZE_1
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (tid < 8)
+ {
+#endif
smem[tid] = product = product + smem[tid + 8];
+#ifdef WAVE_SIZE_1
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (tid < 4)
+ {
+#endif
smem[tid] = product = product + smem[tid + 4];
+#ifdef WAVE_SIZE_1
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (tid < 2)
+ {
+#endif
smem[tid] = product = product + smem[tid + 2];
+#ifdef WAVE_SIZE_1
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (tid < 1)
+ {
+#endif
smem[tid] = product = product + smem[tid + 1];
}
// Extract descriptors
__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)
+ 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);
}
__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)
+ 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);
// Gradients computation
__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 float angle_scale, const char correct_gamma, const int cnbins)
+ const __global uchar4 * img, __global float * grad, __global uchar * 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);
}
__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,
- const float angle_scale, const char correct_gamma, const int cnbins)
+ __global const uchar * img, __global float * grad, __global uchar * 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);
projects / platform / upstream / opencv.git / commitdiff