AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-548838 Add local and global fence support for barrier function (#437)

Browse Source 
* 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 <Tim.Gu@Amd.com>
Co-authored-by: systems-assistant[bot] <systems-assistant[bot]@users.noreply.github.com>
Co-authored-by: Tim Gu <timgu102@amd.com>
This commit is contained in:
systems-assistant[bot]
2025-09-16 14:20:57 -04:00
committed by GitHub
parent 288bca17ea
commit d5fc1b3703
2 changed files with 17 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h
+15 -3
View File
@@ -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) {
projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h
+2 -1
View File
@@ -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