Skip to content

fix(rccl): drop the hipHostMallocUncached fence from LL store<> #13451

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
i-kosarev merged 2 commits into from
Oct 16, 2026
Merged

fix(rccl): drop the hipHostMallocUncached fence from LL store<> #13451

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
d3794a3
fix(rccl): drop the hipHostMallocUncached fence from LL store<>
i-kosarev Oct 14, 2026
a570e33
docs(rccl): state the LL store<> contract instead of the removed fence
i-kosarev Oct 15, 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/CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,6 +56,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
* Fixed virtual address space exhaustion when symmetric windows backed by the same physical allocation are registered repeatedly.
* Fixed Copy Engine `ncclAlltoAllv` hanging on intra-node communicators of more than 8 ranks with uneven send sizes. Each rank decided from its own send sizes whether to add intra-batch barriers, so ranks could issue different barrier counts. The decision now uses the full send-size matrix, which every rank shares.
* Fixed `NCCL_P2P_DISABLE` and `NCCL_P2P_LEVEL` being ignored on Intel, ARM and Zhaoxin hosts, where the CPU-based default P2P level overwrote the user setting. On such hosts `NCCL_P2P_DISABLE=1` now falls back to SHM as documented, and `NCCL_P2P_LEVEL` is honored.
* Fixed LL-protocol latency on gfx950 when RCCL is built against a ROCm older than 7.0.2 without the 128-bit global load/store builtins (for example ROCm 7.0.1, or `-DDWORDX4_INTRINSICS=OFF`). Every store to the user output buffer was followed by a release fence and an L2 writeback, which made 4 KB to 64 KB `ncclBroadcast` about 30% slower on MI355X. Builds against ROCm 7.0.2 or later were not affected and are unchanged.

### Known issues
* The FP8 ReduceSum and ReduceCopy device APIs are not exported in the LLVM bitcode library.
Expand Down
7 changes: 2 additions & 5 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 @@ -367,6 +367,8 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
return elt;
}

// User buffers only. Comm FIFO stores must go through storeLL, which keeps
// the gfx950 hipHostMallocUncached fence for ROCm < 7.0.2.
template <typename U>
__device__ static void store(U* dst, U val) {
union {
Expand Down Expand Up @@ -396,11 +398,6 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
else if (sizeof(U) == 2) __builtin_nontemporal_store(u2, (uint16_t*)dst);
else if (sizeof(U) == 4) __builtin_nontemporal_store(u4, (uint32_t*)dst);
else __builtin_nontemporal_store(u8, (uint64_t*)dst);
#if defined(__gfx950__) && ROCM_VERSION < 70002
__builtin_amdgcn_fence(
__ATOMIC_RELEASE,
""); // flush cache on gfx950 if ROCr fix for hipHostMallocUncached is not available (ROCm version < 7.0.2)
#endif
Comment thread
alex-breslow-amd marked this conversation as resolved.
Resolved
#endif
#else
if (sizeof(U) == 1) asm volatile("st.volatile.global.b8 [%0],%1;" ::"l"(dst), "r"(u4) : "memory");
Expand Down
Loading