Skip to content

fix(rccl): replace __hip_atomic builtins with Clang scoped atomics #11321

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
alex-breslow-amd merged 5 commits into from
Sep 9, 2026
Merged

fix(rccl): replace __hip_atomic builtins with Clang scoped atomics #11321

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
84dc12a
Add scoped atomics to replace deprecated hip atomics
alex-breslow-amd Sep 8, 2026
35f73af
Add poison pill where use of __hip_atomic builtins causes the code to…
alex-breslow-amd Sep 8, 2026
f2c2901
Fix compilation issue
alex-breslow-amd Sep 8, 2026
8320c7c
Add poison tests to host tests
alex-breslow-amd Sep 8, 2026
f3e6526
Address Argus review feedback on the scoped-atomics port
alex-breslow-amd Sep 8, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions projects/rccl/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -110,6 +110,7 @@ option(TRACE "Enable additional tracing"
option(QP_TRACKING "Enable per-device QP tracking" OFF)
option(FAULT_INJECTION "Enable fault injection" ON)
option(QUIET_WARNINGS "Supress compiler warnings" OFF)
option(RCCL_POISON_HIP_ATOMICS "Reject __hip_atomic_* builtins in RCCL sources" ON)
Comment thread
alex-breslow-amd marked this conversation as resolved.
Resolved
option(DWORDX4_INTRINSICS "Turn off compilation with dwordx4 intrinsics" ON)
if(ENABLE_TDM_SIMPLE)
add_compile_definitions(ENABLE_TDM_SIMPLE=1)
Expand Down
17 changes: 17 additions & 0 deletions projects/rccl/src/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1209,6 +1209,23 @@ if(RCCL_ROCPROFILER_REGISTER)
target_link_libraries(
rccl PRIVATE rocprofiler-register::rocprofiler-register)
endif()

# Poison the __hip_atomic_* builtins
#==================================================================================================
# RCCL spells device atomics with the Clang scoped-atomic builtins so that memory
# scope is explicit at every call site. Force-include (rather than #include) the
# poison header so it also covers hipified and generated sources, and set it here,
# before DeviceLinker.cmake snapshots COMPILE_OPTIONS, so the per-kernel device
# compiles inherit it too. See src/include/poison_hip_atomics.h for the coverage
# this does and does not give.
if(RCCL_POISON_HIP_ATOMICS)
# Joined --include=<file> rather than separate-argument -include <file>: the

@amd-argus amd-argus Bot Sep 8, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This comment says the joined form was needed because rccl-device-compile "would take the path for a source file", but the same PR adds explicit -include/-imacros handling at tools/rccl-device-compile:722-725, which makes that no longer true. The other direction also holds: since CMake emits only the joined form, and the pre-existing catch-all at rccl-device-compile:726 already forwards it, the new parser branch is unreachable in-tree. Could we reconcile the two, keeping whichever you intended?

@alex-breslow-amd alex-breslow-amd Sep 8, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Good catch -- the comment was stale the moment I added that branch.

Kept the branch, rewrote the comment. It now states the actual reason: --include=<file> is a single token, so it needs no cooperation from anything that forwards flags one at a time, and it rides the pre-existing startswith('-') catch-all unchanged.

On the branch being unreachable, you are right that it is, from CMake's emission. I would rather keep it than delete it: a flag forwarder that silently turns -include foo.h into a source file is a sharp edge for anyone adding a flag by hand later. I added a comment in rccl-device-compile saying exactly that and naming which spelling the build actually emits, so the next reader does not have to rediscover the relationship.

Verified that --include=<path>, -include <path> and -imacros <path> all forward the path and none of them leak into sources.

# flag has to survive rccl-device-compile, which forwards unrecognized tokens
# one at a time and would take the path for a source file.
target_compile_options(rccl PRIVATE --include=${RCCL_SOURCE_DIR}/src/include/poison_hip_atomics.h)
Comment thread
alex-breslow-amd marked this conversation as resolved.
Resolved
message(STATUS "RCCL_POISON_HIP_ATOMICS enabled - __hip_atomic_* builtins rejected in RCCL sources")
endif()

# Device Linker: include pipeline and wire up dependencies
# (placed after all target_compile_definitions so DeviceLinker.cmake
# can read the rccl target's definitions directly)
Expand Down
4 changes: 2 additions & 2 deletions projects/rccl/src/algorithms/dda/device/CollCommon_ll128.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,10 +65,10 @@ __host__ __device__ __forceinline__ size_t ddaLL128NumLines(size_t nWords) {
// comes from gfx1250 preserving per-thread program order of system-scope stores,
// not from an atomic fence. Matches the validated microbenchmark primitives.
__device__ __forceinline__ void ddaLL128StoreWord(uint64_t* p, uint64_t v) {
__hip_atomic_store((u64_gptr)p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((u64_gptr)p, v, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
}
__device__ __forceinline__ uint64_t ddaLL128LoadWord(const uint64_t* p) {
return __hip_atomic_load((u64_gptr) const_cast<uint64_t*>(p), __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
return __scoped_atomic_load_n((u64_gptr) const_cast<uint64_t*>(p), __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
}

// Element-wise add of the T-elements packed into two 8B payload words. An 8B
Expand Down
8 changes: 4 additions & 4 deletions projects/rccl/src/device/all_gather.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -195,7 +195,7 @@ struct RunWorkColl<ncclFuncAllGather, T, RedOp, NCCL_ALGO_PAT, NCCL_PROTO_SIMPLE
while (1) {
struct ncclPatStep* ps = shmem->patSteps + (step % NCCL_SHMEM_PAT_STEPS);
int* poll = &ps->flags;
while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) != 0) {
while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) != 0) {
pollCount++;// Wait for workers to be done with step 'step-NCCL_SHMEM_PAT_STEPS'
}
patAlgo.getNextOp(ps);
Expand Down Expand Up @@ -224,14 +224,14 @@ struct RunWorkColl<ncclFuncAllGather, T, RedOp, NCCL_ALGO_PAT, NCCL_PROTO_SIMPLE
while (1) {
struct ncclPatStep* ps = shmem->patSteps + (step % NCCL_SHMEM_PAT_STEPS);
int* poll = &ps->flags;
while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) == 0) {
while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) == 0) {
pollCount++; // Wait for compute thread
}
int last = ps->last;
prims.patCopy(ps, shmem);
if (tidInGroup == 0)
__hip_atomic_store(poll, 0, __ATOMIC_RELEASE,
__HIP_MEMORY_SCOPE_WORKGROUP); // Return element to compute thread
__scoped_atomic_store_n(poll, 0, __ATOMIC_RELEASE,
__MEMORY_SCOPE_WRKGRP); // Return element to compute thread
if (last) break;
step += nGroups;
}
Expand Down
16 changes: 8 additions & 8 deletions projects/rccl/src/device/ce_reduce/ce_reduce_impl.h.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,12 +38,12 @@ __device__ __forceinline__ void ncclCeGlobalBlockBarrier(uint32_t* d_barrierSync
uint32_t* doneGen = d_barrierSync + 1;
__syncthreads();
if (threadIdx.x == 0) {
uint32_t prev = __hip_atomic_fetch_add(arrival, 1, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
uint32_t prev = __scoped_atomic_fetch_add(arrival, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
if (prev + 1 == (uint32_t)numBlocks) {
__hip_atomic_store(arrival, 0, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
__hip_atomic_store(doneGen, gen, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_AGENT);
__scoped_atomic_store_n(arrival, 0, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
__scoped_atomic_store_n(doneGen, gen, __ATOMIC_RELEASE, __MEMORY_SCOPE_DEVICE);
} else {
while (__hip_atomic_load(doneGen, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_AGENT) != gen) {
while (__scoped_atomic_load_n(doneGen, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_DEVICE) != gen) {
__builtin_amdgcn_s_sleep(1);
}
}
Expand Down Expand Up @@ -249,8 +249,8 @@ __global__ __launch_bounds__(256) void ncclCeLocalReduceKernelVec(
while (true) {
int ready = 0;
for (int r = 0; r < nRanks; r++) {
uint32_t sv = __hip_atomic_load((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], __ATOMIC_ACQUIRE,
__HIP_MEMORY_SCOPE_SYSTEM);
uint32_t sv = __scoped_atomic_load_n((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], __ATOMIC_ACQUIRE,
__MEMORY_SCOPE_SYSTEM);
if (sv == 1) {
ready++;
}
Expand Down Expand Up @@ -354,8 +354,8 @@ __global__ __launch_bounds__(256) void ncclCeLocalReduceKernelVec(
}
if (blockIdx.x == 0) {
for (int r = (int)threadIdx.x; r < nRanks; r += (int)blockDim.x) {
__hip_atomic_store((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], 0, __ATOMIC_RELEASE,
__HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], 0, __ATOMIC_RELEASE,
__MEMORY_SCOPE_SYSTEM);
}
}
}
Expand Down
4 changes: 2 additions & 2 deletions projects/rccl/src/device/common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,8 @@
#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
#define STORE(DST, SRC) \
{ \
__hip_atomic_store((__attribute__((address_space(1))) __typeof__(*(DST))*)(DST), (SRC), __ATOMIC_RELAXED, \
__HIP_MEMORY_SCOPE_SYSTEM); \
__scoped_atomic_store_n((__attribute__((address_space(1))) __typeof__(*(DST))*)(DST), (SRC), __ATOMIC_RELAXED, \
__MEMORY_SCOPE_SYSTEM); \
}
#elif defined(__GFX9__)
#define STORE(DST, SRC) \
Expand Down
32 changes: 16 additions & 16 deletions projects/rccl/src/device/op128.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -272,7 +272,7 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value)
// template<> __device__ __forceinline__ void st_relaxed_gpu_global<0>(uintptr_t addr, BytePack<0> value) {}

// Used to define implementations for above prototypes.
#define DEFINE_ld_st__size_space_hip_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, \
#define DEFINE_ld_st__size_space_scoped_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, \
addr_reg_ty) \
template <> \
__device__ __forceinline__ BytePack<bytes> ld_##space<bytes>(addr_cxx_ty addr) { \
Expand All @@ -285,16 +285,16 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value)
template <> \
__device__ __forceinline__ BytePack<bytes> ld_volatile_##space<bytes>(addr_cxx_ty addr) { \
data_cxx_ty tmp; \
tmp = __hip_atomic_load((__attribute__((address_space(1))) data_cxx_ty*)addr, __ATOMIC_RELAXED, \
__HIP_MEMORY_SCOPE_SYSTEM); \
tmp = __scoped_atomic_load_n((__attribute__((address_space(1))) data_cxx_ty*)addr, __ATOMIC_RELAXED, \
__MEMORY_SCOPE_SYSTEM); \
BytePack<bytes> ans; \
ans.native = tmp; \
return ans; \
} \
template <> \
__device__ __forceinline__ void st_##space<bytes>(addr_cxx_ty addr, BytePack<bytes> value) { \
__hip_atomic_store((__attribute__((address_space(1))) data_cxx_ty*)addr, value.native, __ATOMIC_RELAXED, \
__HIP_MEMORY_SCOPE_SYSTEM); \
__scoped_atomic_store_n((__attribute__((address_space(1))) data_cxx_ty*)addr, value.native, __ATOMIC_RELAXED, \
__MEMORY_SCOPE_SYSTEM); \
}

#define DEFINE_ld_st__size_space_fallback(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, \
Expand Down Expand Up @@ -322,7 +322,7 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value)

#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
#define DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \
DEFINE_ld_st__size_space_hip_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty)
DEFINE_ld_st__size_space_scoped_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty)
#else
#define DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \
DEFINE_ld_st__size_space_fallback(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty)
Expand Down Expand Up @@ -359,7 +359,7 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value)
DEFINE_ld_st__size(1, uint8_t, b8, r) DEFINE_ld_st__size(2, uint16_t, b16, h) DEFINE_ld_st__size(4, uint32_t, b32, r)
DEFINE_ld_st__size(8, uint64_t, b64, l)
#undef DEFINE_ld_st__size_space
#undef DEFINE_ld_st__size_space_hip_atomic
#undef DEFINE_ld_st__size_space_scoped_atomic
#undef DEFINE_ld_st__size_space_fallback
#undef DEFINE_ld_st__size

Expand Down Expand Up @@ -427,13 +427,13 @@ DEFINE_ld_st_16__space(global, uintptr_t, l)
__device__ __forceinline__ uint64_t ld_volatile_global(uint64_t* ptr) {
uint64_t ans;
ans =
__hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
return ans;
}
__device__ __forceinline__ uint64_t ld_relaxed_sys_global(uint64_t* ptr) {
uint64_t ans;
ans =
__hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
return ans;
}

Expand All @@ -450,21 +450,21 @@ __device__ __forceinline__ uint64_t ld_relaxed_sys_global(uint64_t* ptr) {
__device__ __forceinline__ uint64_t ld_acquire_sys_global(uint64_t* ptr) {
uint64_t ans;
ans =
__hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_SYSTEM);
return ans;
}

__device__ __forceinline__ void st_volatile_global(uint64_t* ptr, uint64_t val) {
__hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
}
__device__ __forceinline__ void st_relaxed_sys_global(uint64_t* ptr, uint64_t val) {
__hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
}
__device__ __forceinline__ void st_release_sys_global(uint64_t* ptr, uint64_t val) {
__hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELEASE,
__HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELEASE,
__MEMORY_SCOPE_SYSTEM);
}

__device__ __forceinline__ void fence_acq_rel_sys() {
Expand Down
6 changes: 3 additions & 3 deletions projects/rccl/src/device/primitives.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,10 +27,10 @@
if (wid == 0) { \
(BARRIER_NEXT) += (NWORKERS) / WARP_SIZE; \
__THREAD_FENCE; \
__hip_atomic_fetch_add((BARRIERS_PTR), 1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_WORKGROUP); \
__scoped_atomic_fetch_add((BARRIERS_PTR), 1, __ATOMIC_RELEASE, __MEMORY_SCOPE_WRKGRP); \
int spins = 0; \
int rate_limit = 50; \
while (__hip_atomic_load((BARRIERS_PTR), __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) < (BARRIER_NEXT)) { \
while (__scoped_atomic_load_n((BARRIERS_PTR), __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) < (BARRIER_NEXT)) { \
spins++; \
if (spins == NCCL_SPINS_BEFORE_CHECK_ABORT) { \
if (__atomic_load_n(ncclShmem.comm.abortFlag, __ATOMIC_SEQ_CST)) { \
Expand All @@ -42,7 +42,7 @@
if (spins == 0 && rate_limit > 0) { \
rate_limit--; \
traceData(__LINE__, threadIdx.x, \
__hip_atomic_load((BARRIERS_PTR), __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP), \
__scoped_atomic_load_n((BARRIERS_PTR), __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP), \
(BARRIER_NEXT)); \
} \
__builtin_amdgcn_s_sleep(1); \
Expand Down
34 changes: 17 additions & 17 deletions projects/rccl/src/device/prims_ll.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -288,8 +288,8 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
*((u64_gptr)dst->v + 1) = *((u64_gptr)i4.v + 1);
#else
#if defined(__gfx1200__) || defined(__gfx1201__)
__hip_atomic_store((u64_gptr)dst->v, i4.v[0], __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__hip_atomic_store((u64_gptr)dst->v + 1, i4.v[1], __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((u64_gptr)dst->v, i4.v[0], __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((u64_gptr)dst->v + 1, i4.v[1], __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM);
#else
*((u64_gptr)dst->v) = *((u64_gptr)i4.v);
*((u64_gptr)dst->v + 1) = *((u64_gptr)i4.v + 1);
Expand Down Expand Up @@ -322,16 +322,16 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
if (sizeof(U) == 1)
u1 =
__hip_atomic_load((__attribute__((address_space(1))) uint8_t*)src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_load_n((__attribute__((address_space(1))) uint8_t*)src, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);

@amd-argus amd-argus Bot Sep 8, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This line is now 121 columns against the ColumnLimit: 120 in .clang-format. The rename grows the callee by 5 characters, but every continuation line in the change kept its old leading whitespace, so arguments no longer sit under the open paren:

prims_ll.h            328, 331, 334, 384, 387, 390, 393
op128.h               276, 289, 297, 459, 463, 467
hip_compat.h          195, 198, 255-257, 260-262
all_gather.h          234
reduce_scatter.h      261
common.h              25
gin_anvil_ipc_copy.h  25
ce_reduce_impl.h.in   253, 358

Could we re-run clang-format over the changed files? Note .pre-commit-config.yaml:49 matches \.h$ and not .h.in, so ce_reduce_impl.h.in needs a hand fix. compare_exchange_weak at hip_compat.h:233-235 is coincidentally already correct, since the old and new callee names are both 34 characters.

@alex-breslow-amd alex-breslow-amd Sep 8, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed. Ran clang-format restricted to the ranges this PR changed, and the output matched your line list exactly: op128.h 6 for 6, hip_compat.h 8 for 8, and the single lines in all_gather.h, reduce_scatter.h, common.h and gin_anvil_ipc_copy.h. ce_reduce_impl.h.in went through the same pass with --assume-filename, since as you note the hook's \.h$ does not reach it. prims_ll.h:325 reflowed to two lines; it was the only added line in the PR over 120 columns.

You were right about hip_compat.h:233-235 too -- clang-format leaves those alone, same 34-character coincidence.

Two things I checked that are worth passing back:

  • Whole-file formatting produces exactly the same result as the range-restricted run for all eight files, so there was no collateral churn to weigh. The one file I did revert was llx_latency_bench.c, where clang-format wanted to restyle two function bodies that this PR only touched for the builtin rename.
  • .pre-commit-config.yaml will not catch this class of thing today. It pins mirrors-clang-format v18.1.4, but projects/rccl/.clang-format uses AlignPPAndNotPP, OneLineFormatOffRegex and BreakBeforeCloseBracketBracedList, which only exist in clang-format 21 -- both 18 and ROCm's own 20 refuse to parse the config outright. Pre-existing and out of scope for this PR, but the hook is not the safety net it looks like.

To confirm the reformat was inert I compared whitespace-normalized text for all eight files against the previous revision (token-identical), then compiled real TUs with the poison active: common.h + all_gather.h and common.h + reduce_scatter.h, which pull in primitives.h -> prims_ll.h and op128.h.

else if (sizeof(U) == 2)
u2 = __hip_atomic_load((__attribute__((address_space(1))) uint16_t*)src, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
u2 = __scoped_atomic_load_n((__attribute__((address_space(1))) uint16_t*)src, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
else if (sizeof(U) == 4)
u4 = __hip_atomic_load((__attribute__((address_space(1))) uint32_t*)src, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
u4 = __scoped_atomic_load_n((__attribute__((address_space(1))) uint32_t*)src, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
else
u8 = __hip_atomic_load((__attribute__((address_space(1))) uint64_t*)src, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
u8 = __scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)src, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
#else
if (sizeof(U) == 1)
#ifdef __GFX11__
Expand Down Expand Up @@ -380,17 +380,17 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
if (sizeof(U) == 1)
__hip_atomic_store((__attribute__((address_space(1))) uint8_t*)dst, u1, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((__attribute__((address_space(1))) uint8_t*)dst, u1, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
else if (sizeof(U) == 2)
__hip_atomic_store((__attribute__((address_space(1))) uint16_t*)dst, u2, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((__attribute__((address_space(1))) uint16_t*)dst, u2, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
else if (sizeof(U) == 4)
__hip_atomic_store((__attribute__((address_space(1))) uint32_t*)dst, u4, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((__attribute__((address_space(1))) uint32_t*)dst, u4, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
else
__hip_atomic_store((__attribute__((address_space(1))) uint64_t*)dst, u8, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)dst, u8, __ATOMIC_RELAXED,
__MEMORY_SCOPE_SYSTEM);
#else
if (sizeof(U) == 1) __builtin_nontemporal_store(u1, (uint8_t*)dst);
else if (sizeof(U) == 2) __builtin_nontemporal_store(u2, (uint16_t*)dst);
Expand Down
8 changes: 4 additions & 4 deletions projects/rccl/src/device/reduce_scatter.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -223,7 +223,7 @@ struct RunWorkColl<ncclFuncReduceScatter, T, RedOp, NCCL_ALGO_PAT, NCCL_PROTO_SI
while (1) {
struct ncclPatStep* ps = shmem->patSteps + (step % NCCL_SHMEM_PAT_STEPS);
int* poll = &ps->flags;
while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) != 0) {
while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) != 0) {
// pollCount++;// unused variable - compiler warning // Wait for workers to be done with step 'step-NCCL_SHMEM_PAT_STEPS'
}
patAlgo.getNextOp(ps);
Expand Down Expand Up @@ -251,14 +251,14 @@ struct RunWorkColl<ncclFuncReduceScatter, T, RedOp, NCCL_ALGO_PAT, NCCL_PROTO_SI
while (1) {
struct ncclPatStep* ps = shmem->patSteps + (step % NCCL_SHMEM_PAT_STEPS);
int* poll = &ps->flags;
while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) == 0) {
while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) == 0) {
// pollCount++; // unused variable - compiler warning // Wait for compute thread
}
int last = ps->last;
prims.patReduce(ps, shmem);
if (tidInGroup == 0)
__hip_atomic_store(poll, 0, __ATOMIC_RELEASE,
__HIP_MEMORY_SCOPE_WORKGROUP); // Return element to compute thread
__scoped_atomic_store_n(poll, 0, __ATOMIC_RELEASE,
__MEMORY_SCOPE_WRKGRP); // Return element to compute thread
if (last) break;
step += nGroups;
}
Expand Down
4 changes: 2 additions & 2 deletions projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_ipc_copy.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,8 @@ namespace detail {

// Relaxed: ordering is provided by the caller's fence (signalPeer).
NCCL_DEVICE_INLINE void ipcFlatAtomicAddSys64(uint64_t* dst, uint64_t val) {
__hip_atomic_fetch_add(reinterpret_cast<unsigned long long*>(dst), static_cast<unsigned long long>(val),
__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__scoped_atomic_fetch_add(reinterpret_cast<unsigned long long*>(dst), static_cast<unsigned long long>(val),
__ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
}

} // namespace detail
Expand Down
Loading
Loading