[libclc] 47/92: Make ptx barrier work irrespective of the cl_mem_fence_flags

[Andreas Boll]

This is an automated email from the git hooks/post-receive script.

aboll-guest pushed a commit to branch master
in repository libclc.

commit ba746aa086d03a09dae870cdfa0ab46c61eaa0a0
Author: Jeroen Ketema <j.ketema at xs4all.nl>
Date:   Mon Oct 9 18:36:48 2017 +0000

    Make ptx barrier work irrespective of the cl_mem_fence_flags
    
    This generates a "bar.sync 0” instruction, which not only causes the
    threads to wait, but does acts as a memory fence, as required by
    OpenCL. The fence does not differentiate between local and global
    memory. Unfortunately, there is no similar instruction which does
    not include a memory fence. Hence, we cannot optimize the case
    where neither CLK_LOCAL_MEM_FENCE nor CLK_GLOBAL_MEM_FENCE is
    passed.
    
    
    git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@315228 91177308-0d34-0410-b5e6-96231b3b80d8
---
 ptx-nvidiacl/lib/synchronization/barrier.cl | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/ptx-nvidiacl/lib/synchronization/barrier.cl b/ptx-nvidiacl/lib/synchronization/barrier.cl
index 88e1493..930e36a 100644
--- a/ptx-nvidiacl/lib/synchronization/barrier.cl
+++ b/ptx-nvidiacl/lib/synchronization/barrier.cl
@@ -1,8 +1,6 @@
 #include <clc/clc.h>
 
 _CLC_DEF void barrier(cl_mem_fence_flags flags) {
-  if (flags & CLK_LOCAL_MEM_FENCE) {
-    __syncthreads();
-  }
+  __syncthreads();
 }
 

-- 
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/libclc.git