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

SWDEV-1 - Reapply "SWDEV-1 - Directly call __builtin_amdgcn_fence"

Browse Source 
This reverts commit 7b3cf2e9d7.

Resubmit after https://github.com/ROCmSoftwarePlatform/rocPRIM/commit/ba8a86166ec9c4cf3a40ecde7017feeb1484e7f6

Change-Id: I194559484a5a0d7f4443381fe77218ea73d2def9


[ROCm/clr commit: 8c7391398a]
This commit is contained in:
Matthew Arsenault
2023-01-25 14:47:48 -05:00
committed by Brian Sumner
parent 5981bd5c22
commit 716ddd1d98
2 changed files with 7 additions and 32 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h
+7 -10
View File
@@ -777,21 +777,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
@@ -864,14 +864,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();
}
@@ -881,7 +878,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__
projects/clr/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