[libc-commits] [libc] [libc] Add memory fence utility to the GPU utilities (PR #91756)
Joseph Huber via libc-commits
libc-commits at lists.llvm.org
Fri May 10 14:43:14 PDT 2024
================
@@ -140,6 +140,12 @@ LIBC_INLINE uint32_t get_lane_size() {
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
}
+/// Waits for all pending memory operations to complete in program order.
+[[clang::convergent]] LIBC_INLINE void memory_fence() {
+ __builtin_amdgcn_s_waitcnt(0);
+ __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "");
----------------
jhuber6 wrote:
I don't need it for any other use-case right now. If I were to start exporting these as generic utilties (Which I might do in compiler-rt eventually) then definitely.
https://github.com/llvm/llvm-project/pull/91756
More information about the libc-commits
mailing list