From e8fb1335fd0bd609d727ca589bf65a026bcb74a5 Mon Sep 17 00:00:00 2001 From: Dingming Wu Date: Mon, 10 Feb 2025 08:53:25 -0800 Subject: [PATCH] Replace atomicAdd with _hip_atmoc_fetch_add in getting colltrace tail position (#1539) --- src/device/common.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/device/common.h b/src/device/common.h index a9ecd732e6..49d3ee1c02 100644 --- a/src/device/common.h +++ b/src/device/common.h @@ -43,7 +43,7 @@ #endif #ifdef ENABLE_COLLTRACE #define INC_COLL_TRACE \ - uint32_t pos = atomicAdd(&ncclShmem.collTraceTail->tail, 1)%COLLTRACE_NUM_ITEMS; \ + uint32_t pos = __hip_atomic_fetch_add(&ncclShmem.collTraceTail->tail, 1, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_WORKGROUP)%COLLTRACE_NUM_ITEMS; \ struct ncclCollTrace* collTrace = ncclShmem.collTrace+pos; \ collTrace->timeStamp = wall_clock64(); \ collTrace->bid = blockIdx.x; \