maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/libclc/ptx-nvidiacl/lib/synchronization/barrier.cl

7 lines
94 B
Common Lisp
Raw Normal View History

PTX: move implementations of work-item and synchronisation functions to lib, and add header files in generic. Incorporates a patch by Tom Stellard! llvm-svn: 161313
2012-08-06 06:25:37 +08:00
#include <clc/clc.h>
_CLC_DEF void barrier(cl_mem_fence_flags flags) {
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. llvm-svn: 315228
2017-10-10 02:36:48 +08:00
__syncthreads();
Initial commit. llvm-svn: 147756
2012-01-09 06:09:58 +08:00
}