From d5fc1b37038db39678053b10194837d600e7e06a Mon Sep 17 00:00:00 2001 From: "systems-assistant[bot]" <221163467+systems-assistant[bot]@users.noreply.github.com> Date: Tue, 16 Sep 2025 14:20:57 -0400 Subject: [PATCH] SWDEV-548838 Add local and global fence support for barrier function (#437) * SWDEV-548838 Add local and global fence support for barrier function The original barrier function didn't distinct between local and global scope. There was only __CLK_LOCAL_MEM_FENCE which triggers both local and global fence. This commit introduces __CLK_LOCAL_MEM_FENCE and __CLK_GLOBAL_MEM_FENCE that properly distinguish the scopes. --------- Co-authored-by: Tim Co-authored-by: systems-assistant[bot] Co-authored-by: Tim Gu --- .../hip/amd_detail/amd_device_functions.h | 18 +++++++++++++++--- .../hip/amd_detail/device_library_decls.h | 3 ++- 2 files changed, 17 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h index 16131ab920..cf33122570 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -687,10 +687,18 @@ __device__ inline static void __threadfence_system() { __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, ""); } __device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags) { - if (flags) { + if (flags == (__CLK_GLOBAL_MEM_FENCE | __CLK_LOCAL_MEM_FENCE)) { __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup"); __builtin_amdgcn_s_barrier(); __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup"); + } else if (flags & (__CLK_GLOBAL_MEM_FENCE)) { + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup", "global"); + __builtin_amdgcn_s_barrier(); + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup", "global"); + } else if (flags & (__CLK_LOCAL_MEM_FENCE)) { + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup", "local"); + __builtin_amdgcn_s_barrier(); + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup", "local"); } else { __builtin_amdgcn_s_barrier(); } @@ -698,8 +706,12 @@ __device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags) { __device__ inline static void __barrier(int n) { __work_group_barrier((__cl_mem_fence_flags)n); } -__device__ inline __attribute__((convergent)) void __syncthreads() { - __barrier(__CLK_LOCAL_MEM_FENCE); +__device__ +inline +__attribute__((convergent)) +void __syncthreads() +{ + __barrier(__CLK_GLOBAL_MEM_FENCE | __CLK_LOCAL_MEM_FENCE); } __device__ inline __attribute__((convergent)) int __syncthreads_count(int predicate) { diff --git a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h index 00a98f5750..33623f5881 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h +++ b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h @@ -128,7 +128,8 @@ __device__ inline static __local void* __to_local(unsigned x) { return (__local #endif //__HIP_DEVICE_COMPILE__ // Using hip.amdgcn.bc - sync threads -#define __CLK_LOCAL_MEM_FENCE 0x01 +#define __CLK_LOCAL_MEM_FENCE 0x01 +#define __CLK_GLOBAL_MEM_FENCE 0x02 typedef unsigned __cl_mem_fence_flags; #endif