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

SWDEV-1 - Directly call __builtin_amdgcn_fence

Browse Source 
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
This commit is contained in:
Matt Arsenault
2022-12-23 09:26:22 -05:00
parent 91b851fad7
commit 3e54bad5b3
2 changed files with 7 additions and 32 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_device_functions.h
+7 -10
View File
@@ -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__
hipamd/include/hip/amd_detail/device_library_decls.h
-22
View File
@@ -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