mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-29 05:29:54 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Fix ocl::HOG crash on Intel OCL

Browse Source
This commit is contained in:
yao 2013-02-28 16:56:39 +08:00
parent 620c699456
commit f12369a53c
1 changed files with 9 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
modules/ocl/src/kernels/objdetect_hog.cl
View File

@ -140,6 +140,10 @@ float reduce_smem(volatile __local float* smem, int size)
if (tid < 32) if (tid < 32)
{ {
if (size >= 64) smem[tid] = sum = sum + smem[tid + 32]; if (size >= 64) smem[tid] = sum = sum + smem[tid + 32];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
{
if (size >= 32) smem[tid] = sum = sum + smem[tid + 16]; if (size >= 32) smem[tid] = sum = sum + smem[tid + 16];
if (size >= 16) smem[tid] = sum = sum + smem[tid + 8]; if (size >= 16) smem[tid] = sum = sum + smem[tid + 8];
if (size >= 8) smem[tid] = sum = sum + smem[tid + 4]; if (size >= 8) smem[tid] = sum = sum + smem[tid + 4];
@ -224,6 +228,11 @@ __kernel void classify_hists_kernel(const int cblock_hist_size, const int cdescr
{ {
volatile __local float* smem = products; volatile __local float* smem = products;
smem[tid] = product = product + smem[tid + 32]; smem[tid] = product = product + smem[tid + 32];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
{
volatile __local float* smem = products;
smem[tid] = product = product + smem[tid + 16]; smem[tid] = product = product + smem[tid + 16];
smem[tid] = product = product + smem[tid + 8]; smem[tid] = product = product + smem[tid + 8];
smem[tid] = product = product + smem[tid + 4]; smem[tid] = product = product + smem[tid + 4];