From 1de57c06b311e89cd30a3d5103b4859a1e8f2efb Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Fri, 23 Dec 2022 09:26:22 -0500 Subject: [PATCH] SWDEV-1 - Directly call __builtin_amdgcn_fence The device libraries provide hip.bc, which only contains __atomic_work_item_fence. This is identical to OpenCL's atomic_work_item_fence, except prefixed with __. This only existed to switch over the unknown values the OpenCL function provides. There's no need to go through the wrapper here since it's using constant values, and without it we can drop one of the bitcode libraries. I'm guessing this was a holdover from before the fence builtin existed. Change-Id: I8b04614f96194eafddeba295215c12c3856e6c42 [ROCm/clr commit: 3e54bad5b30f4d57d45afd96d7918f021b084a39] --- .../hip/amd_detail/amd_device_functions.h | 17 ++++++-------- .../hip/amd_detail/device_library_decls.h | 22 ------------------- 2 files changed, 7 insertions(+), 32 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 ab1e1f9e97..664f425396 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 @@ -826,21 +826,21 @@ __device__ inline static void __threadfence() { - __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_device); + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent"); } __device__ inline static void __threadfence_block() { - __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_work_group); + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup"); } __device__ inline static void __threadfence_system() { - __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_all_svm_devices); + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, ""); } // abort @@ -913,14 +913,11 @@ void __assertfail() } #endif /* defined(_WIN32) || defined(_WIN64) */ -__device__ -inline -static void __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope) -{ +__device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags) { if (flags) { - __atomic_work_item_fence(flags, __memory_order_release, scope); + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup"); __builtin_amdgcn_s_barrier(); - __atomic_work_item_fence(flags, __memory_order_acquire, scope); + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup"); } else { __builtin_amdgcn_s_barrier(); } @@ -930,7 +927,7 @@ __device__ inline static void __barrier(int n) { - __work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group); + __work_group_barrier((__cl_mem_fence_flags)n); } __device__ 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 8add4fa279..0222870590 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,26 +128,4 @@ __device__ inline static __local void* __to_local(unsigned x) { return (__local #define __CLK_LOCAL_MEM_FENCE 0x01 typedef unsigned __cl_mem_fence_flags; -typedef enum __memory_scope { - __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM, - __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP, - __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE, - __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES, - __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP -} __memory_scope; - -// enum values aligned with what clang uses in EmitAtomicExpr() -typedef enum __memory_order -{ - __memory_order_relaxed = __ATOMIC_RELAXED, - __memory_order_acquire = __ATOMIC_ACQUIRE, - __memory_order_release = __ATOMIC_RELEASE, - __memory_order_acq_rel = __ATOMIC_ACQ_REL, - __memory_order_seq_cst = __ATOMIC_SEQ_CST -} __memory_order; - -// Linked from hip.amdgcn.bc -extern "C" __device__ void -__atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope); - #endif