From c04355bbe3971e7737dd2c6adaf91b74d29ce8ec Mon Sep 17 00:00:00 2001 From: "Belton-Schure, Aidan" Date: Thu, 15 May 2025 09:50:13 +0100 Subject: [PATCH] Add __syncwarp operation (#160) Change-Id: I6a3783beafdbb9f11a3b37333f4ff3f5be27ea54 [ROCm/clr commit: c50610b44daf5d3a4e5761b90f9f13606d102b5d] --- .../hip/amd_detail/amd_warp_sync_functions.h | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h index c4885868e8..84bfcb4320 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h @@ -164,6 +164,20 @@ T __hip_readfirstlane(T val) { } \ } while(0) +__device__ inline void __syncwarp() { + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "wavefront"); + __builtin_amdgcn_wave_barrier(); + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "wavefront"); +} + +template __device__ inline void __syncwarp(MaskT mask) { + static_assert(__hip_internal::is_integral::value && sizeof(MaskT) == 8, + "The mask must be a 64-bit integer. " + "Implicitly promoting a smaller integer is almost always an error."); + __hip_check_mask(mask); + return __syncwarp(); +} + // __all_sync, __any_sync, __ballot_sync template