summaryrefslogtreecommitdiff
path: root/kernels/compiler_local_memory_two_ptr.cl
diff options
context:
space:
mode:
authorYang Rong <rong.r.yang@intel.com>2013-06-08 12:33:37 +0800
committerZhigang Gong <zhigang.gong@linux.intel.com>2013-06-09 15:47:09 +0800
commit751ab1a6baa22022404ed4e0f992e1db84c62300 (patch)
tree4704d4bb2d94c88c08126f6bbf99ec8a01405327 /kernels/compiler_local_memory_two_ptr.cl
parente85d116985d4e9335f805717e4b8aa69e6e659d7 (diff)
downloadbeignet-751ab1a6baa22022404ed4e0f992e1db84c62300.tar.gz
Fix two tests fail when OCL_SIMD_WIDTH=8.
Add barrier for compiler_local_memory and compiler_local_memory_two_ptr, otherwise tests may fail if work group size bigger than thread's simd size. After add barrier, the test compiler_local_memory is same as compiler_local_memory_barrier, so delete test compiler_local_memory. Signed-off-by: Yang Rong <rong.r.yang@intel.com> Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Diffstat (limited to 'kernels/compiler_local_memory_two_ptr.cl')
-rw-r--r--kernels/compiler_local_memory_two_ptr.cl1
1 files changed, 1 insertions, 0 deletions
diff --git a/kernels/compiler_local_memory_two_ptr.cl b/kernels/compiler_local_memory_two_ptr.cl
index f410406b..46589ba6 100644
--- a/kernels/compiler_local_memory_two_ptr.cl
+++ b/kernels/compiler_local_memory_two_ptr.cl
@@ -4,6 +4,7 @@ __kernel void compiler_local_memory_two_ptr(__global int *dst,
{
src0[get_local_id(0)] = get_local_id(0);
src1[get_local_id(0)] = get_global_id(0);
+ barrier(CLK_LOCAL_MEM_FENCE);
dst[get_global_id(0)] = src0[15 - get_local_id(0)] + src1[15 - get_local_id(0)];
}