llvm-project/libclc/amdgcn/lib/synchronization/barrier.cl

#include <clc/clc.h>

_CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
  mem_fence(flags);
  __builtin_amdgcn_s_barrier();
}
amdgcn: rewrite barrier() using fence and clang __builtin_amdgcn_s_barrier Specs require using fences when barrier() is invoked: "The barrier function will either flush any variables stored in local memory or queue a memory fence to ensure correct ordering of memory operations to local memory." and "The barrier function will queue a memory fence to ensure correct ordering of memory operations to global memory." Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu> Reviewed-by: Aaron Watry <awatry@gmail.com> Tested-by: Aaron Watry <awatry@gmail.com> llvm-svn: 311022 2017-08-17 01:09:00 +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) {`
			`mem_fence(flags);`
			`__builtin_amdgcn_s_barrier();`
amdgcn: rewrite barrier() using fence and clang __builtin_amdgcn_s_barrier Specs require using fences when barrier() is invoked: "The barrier function will either flush any variables stored in local memory or queue a memory fence to ensure correct ordering of memory operations to local memory." and "The barrier function will queue a memory fence to ensure correct ordering of memory operations to global memory." Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu> Reviewed-by: Aaron Watry <awatry@gmail.com> Tested-by: Aaron Watry <awatry@gmail.com> llvm-svn: 311022 2017-08-17 01:09:00 +08:00			`}`