From f12369a53c86239579ebf533f276d6fe9992abd3 Mon Sep 17 00:00:00 2001 From: yao Date: Thu, 28 Feb 2013 16:56:39 +0800 Subject: [PATCH] Fix ocl::HOG crash on Intel OCL --- modules/ocl/src/kernels/objdetect_hog.cl | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/modules/ocl/src/kernels/objdetect_hog.cl b/modules/ocl/src/kernels/objdetect_hog.cl index 076c22c..db11ed1 100644 --- a/modules/ocl/src/kernels/objdetect_hog.cl +++ b/modules/ocl/src/kernels/objdetect_hog.cl @@ -140,6 +140,10 @@ float reduce_smem(volatile __local float* smem, int size) if (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 >= 16) smem[tid] = sum = sum + smem[tid + 8]; 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; 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 + 8]; smem[tid] = product = product + smem[tid + 4]; -- 2.7.4