diff options
author | Yang Rong <rong.r.yang@intel.com> | 2013-06-08 12:33:37 +0800 |
---|---|---|
committer | Zhigang Gong <zhigang.gong@linux.intel.com> | 2013-06-09 15:47:09 +0800 |
commit | 751ab1a6baa22022404ed4e0f992e1db84c62300 (patch) | |
tree | 4704d4bb2d94c88c08126f6bbf99ec8a01405327 /kernels/compiler_local_memory_two_ptr.cl | |
parent | e85d116985d4e9335f805717e4b8aa69e6e659d7 (diff) | |
download | beignet-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.cl | 1 |
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)]; } |