llvm-project/libclc/ptx-nvidiacl/lib/synchronization/barrier.cl

#include <clc/clc.h>

_CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
  __syncthreads();
}
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>`

libclc: Make all built-ins overloadable The SPIR spec states that all OpenCL built-in functions should be overloadable and mangled, to ensure consistency. Add the overload attribute to functions which were missing them: work dimensions, memory barriers and fences, and events. Reviewed By: tstellar, jenatali Differential Revision: https://reviews.llvm.org/D82078 2020-08-18 04:44:49 +08:00			`_CLC_DEF _CLC_OVERLOAD 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			`}`