diff options
Diffstat (limited to 'libc/src/__support/GPU/amdgpu/utils.h')
| -rw-r--r-- | libc/src/__support/GPU/amdgpu/utils.h | 5 |
1 files changed, 5 insertions, 0 deletions
diff --git a/libc/src/__support/GPU/amdgpu/utils.h b/libc/src/__support/GPU/amdgpu/utils.h index 9b520a6bcf38..5f8ad74f6aea 100644 --- a/libc/src/__support/GPU/amdgpu/utils.h +++ b/libc/src/__support/GPU/amdgpu/utils.h @@ -140,6 +140,11 @@ 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_fence(__ATOMIC_ACQ_REL, ""); +} + /// Wait for all threads in the wavefront to converge, this is a noop on AMDGPU. [[clang::convergent]] LIBC_INLINE void sync_lane(uint64_t) { __builtin_amdgcn_wave_barrier(); |