[libc-commits] [libc] [libc] Add memory fence utility to the GPU utilities (PR #91756)

[Joseph Huber via libc-commits]

================
@@ -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);
----------------
jhuber6 wrote:

I looked at the IR for the fence and it didn't emit any, so my understanding is that in order to be sure all the preceding operations are done we'd need one prior. Is there an equivalent instruction that has the save effect?

https://github.com/llvm/llvm-project/pull/91756