Repository navigation
fix(rccl) JIRA ID : ROCM-25479 LL: cache-bypass user buffers, plain/non-temporal comm FIFO - #9336
Merged
Merged
Conversation
Retarget the LL protocol memory-access policy on gfx942/gfx950 to match the LL128 rationale: registered user buffers are cacheable and need system-scope cache-bypass, while the comm FIFO buffers are allocated uncached and gain nothing from bypass (it only adds register pressure). - User output store (storeData): switch from plain stores to the system-scope store<> helper (cache-bypass), for both the aligned fast path and the element loop. User input/acc load (load<>) was already system-scope. - Comm FIFO reads (readLL/readLLBeginAll/readLLFinish): drop the system-scope global_load_b128 builtin in favor of non-temporal loads, so the flag poll never observes a stale cache line. - Comm FIFO store (storeLL): drop the system-scope global_store_b128 builtin in favor of plain stores; the gfx950/ROCm<7.0.2 release fence is preserved. Only the RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS path is retargeted; GFX11 and CUDA/PTX paths are unchanged. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 92be5e5)
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
pvallem
reviewed
Jul 28, 2026
pvallem
reviewed
Jul 28, 2026
pvallem
reviewed
Jul 28, 2026
Contributor
|
A few testing suggestions:
|
Contributor
Author
non-temporal load/store have been used for 6+ years. Only recently cache bypassing load/store became available. This PR basically reverting to old instructions for comm buffers. The chance of having functional issue is very low. |
Contributor
Author
|
Single node direct AG w and w/o LL fix. Saving was about 1us.
|
wenkaidu
changed the title
pvallem
approved these changes
Jul 28, 2026
1 task
mustafabar
added a commit
that referenced
this pull request
Sep 1, 2026
…hed VMM behavior on gfx1250 (#10969) ## Motivation The LL/LL128 comm-FIFO access policy set in #9336 rests on the FIFO being allocated uncached. That holds for the legacy IPC allocator (`hipDeviceMallocUncached`) but not for the cuMem/VMM path. On gfx1250 with `NCCL_CUMEM_ENABLE=1`, the flag poll stops observing peer writes and the collective stalls: 3474 us at 1 MB against 45 us once the load carries system scope, a ~75x regression on a 4-rank single-node all_gather. Results stay correct throughout, since the poll simply waits, so this presents as a throughput cliff rather than a wrong answer. This is a workaround, not a fix. A VMM allocation created with `hipMemAllocationTypeUncached` behaves indistinguishably from ordinary cacheable memory on this configuration, even though `hipMemGetAllocationPropertiesFromHandle` reads the property back correctly. A defect has been filed against the HIP runtime (ROCM-30384) with a standalone two-GPU reproducer that shows the same stall outside RCCL entirely, so this is not specific to the collective path. ## Technical Details `__builtin_nontemporal_load` emits `global_load_b64 ... th:TH_LOAD_NT` with no scope modifier. On gfx12-family targets that is an eviction-priority hint, not a coherence qualifier, so it does not force a miss and the poll can sit on a stale line indefinitely. `__builtin_amdgcn_global_load_b128(p, RCCL_SYSTEM_SYNCSCOPE)` emits `global_load_b128 ... scope:SCOPE_SYS`, which does. A clean-build A/B over {non-temporal, system-scope} x {load, store}, with the `.so` checksummed per variant to confirm each build actually differed, isolates the load as the entire fix: | load | store | busbw (GB/s) | |---|---|---| | non-temporal | plain | 0.0389 | | non-temporal | system-scope | 0.0384 | | **system-scope** | **plain** | **2.902** | | system-scope | system-scope | 2.884 | Only the reader's poll changes. The store side keeps the plain cacheable path from #9336, so its store throughput and register pressure are unaffected. The behavior is gated by `RCCL_LL_FIFO_SYS_SCOPE_LOAD` in `rccl_ptr.h`, restricted to gfx1250 - the only architecture where this was observed and measured, and the only one currently using cuMem P2P. Other targets compile byte-identically to before. The single macro is the removal point once ROCM-30384 is resolved. Applies to `readLL` / `readLLBeginAll` / `readLLFinish` in `prims_ll.h` and `load128NT` in `prims_ll128.h`. ## Issue Tracking JIRA ID: AICOMRCCL-2230 Upstream defect (filed): ROCM-30384 - `hipMemAllocationTypeUncached` has no observable effect on VMM allocations on gfx1250 ## Test Plan Exhaustive sweeps over message size across the priority collectives, for both LL and LL128 protocols, on gfx1250 with cuMem enabled, scaling up to 32 ranks. Each configuration is compared against a control build with the workaround disabled. ## Test Result 4-rank all_gather, 1 MB: | Proto | gate off | gate on | |---|---|---| | LL128 | 3452 us / 0.038 GB/s | 45.7 us / 2.904 GB/s | | LL | 3133 us / 0.043 GB/s | 32.4 us / 4.059 GB/s | `# Out of bounds values : 0 OK` in every arm. Non-gfx1250 targets are unchanged by construction. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
1 task done
i-kosarev
added a commit
that referenced
this pull request
Oct 14, 2026
storeData has called store<> since #9336. When RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS is 0 (compiler without __builtin_amdgcn_global_load_b128, e.g. ROCm 7.0.1, or DWORDX4_INTRINSICS=OFF), store<> falls back to a non-temporal store followed, on gfx950 with ROCM_VERSION < 70002, by a release fence: buffer_wbl2 sc0 sc1 plus s_waitcnt vmcnt(0) lgkmcnt(0) after every user-buffer store. The fence works around the ROCr hipHostMallocUncached bug for the comm FIFO, which storeLL still covers. store<> only writes user buffers, had no callers before #9336, and its builtin path has not carried the fence since #3812. Only gfx950 builds against ROCm < 7.0.2 without the builtins change; every other build preprocesses to the same code. JIRA ID : ROCM-31938 Co-authored-by: Cursor <cursoragent@cursor.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 14, 2026
…hed VMM behavior on gfx1250 (#10969) ## Motivation The LL/LL128 comm-FIFO access policy set in #9336 rests on the FIFO being allocated uncached. That holds for the legacy IPC allocator (`hipDeviceMallocUncached`) but not for the cuMem/VMM path. On gfx1250 with `NCCL_CUMEM_ENABLE=1`, the flag poll stops observing peer writes and the collective stalls: 3474 us at 1 MB against 45 us once the load carries system scope, a ~75x regression on a 4-rank single-node all_gather. Results stay correct throughout, since the poll simply waits, so this presents as a throughput cliff rather than a wrong answer. This is a workaround, not a fix. A VMM allocation created with `hipMemAllocationTypeUncached` behaves indistinguishably from ordinary cacheable memory on this configuration, even though `hipMemGetAllocationPropertiesFromHandle` reads the property back correctly. A defect has been filed against the HIP runtime (ROCM-30384) with a standalone two-GPU reproducer that shows the same stall outside RCCL entirely, so this is not specific to the collective path. ## Technical Details `__builtin_nontemporal_load` emits `global_load_b64 ... th:TH_LOAD_NT` with no scope modifier. On gfx12-family targets that is an eviction-priority hint, not a coherence qualifier, so it does not force a miss and the poll can sit on a stale line indefinitely. `__builtin_amdgcn_global_load_b128(p, RCCL_SYSTEM_SYNCSCOPE)` emits `global_load_b128 ... scope:SCOPE_SYS`, which does. A clean-build A/B over {non-temporal, system-scope} x {load, store}, with the `.so` checksummed per variant to confirm each build actually differed, isolates the load as the entire fix: | load | store | busbw (GB/s) | |---|---|---| | non-temporal | plain | 0.0389 | | non-temporal | system-scope | 0.0384 | | **system-scope** | **plain** | **2.902** | | system-scope | system-scope | 2.884 | Only the reader's poll changes. The store side keeps the plain cacheable path from #9336, so its store throughput and register pressure are unaffected. The behavior is gated by `RCCL_LL_FIFO_SYS_SCOPE_LOAD` in `rccl_ptr.h`, restricted to gfx1250 - the only architecture where this was observed and measured, and the only one currently using cuMem P2P. Other targets compile byte-identically to before. The single macro is the removal point once ROCM-30384 is resolved. Applies to `readLL` / `readLLBeginAll` / `readLLFinish` in `prims_ll.h` and `load128NT` in `prims_ll128.h`. ## Issue Tracking JIRA ID: AICOMRCCL-2230 Upstream defect (filed): ROCM-30384 - `hipMemAllocationTypeUncached` has no observable effect on VMM allocations on gfx1250 ## Test Plan Exhaustive sweeps over message size across the priority collectives, for both LL and LL128 protocols, on gfx1250 with cuMem enabled, scaling up to 32 ranks. Each configuration is compared against a control build with the workaround disabled. ## Test Result 4-rank all_gather, 1 MB: | Proto | gate off | gate on | |---|---|---| | LL128 | 3452 us / 0.038 GB/s | 45.7 us / 2.904 GB/s | | LL | 3133 us / 0.043 GB/s | 32.4 us / 4.059 GB/s | `# Out of bounds values : 0 OK` in every arm. Non-gfx1250 targets are unchanged by construction. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
…hed VMM behavior on gfx1250 (#10969) ## Motivation The LL/LL128 comm-FIFO access policy set in #9336 rests on the FIFO being allocated uncached. That holds for the legacy IPC allocator (`hipDeviceMallocUncached`) but not for the cuMem/VMM path. On gfx1250 with `NCCL_CUMEM_ENABLE=1`, the flag poll stops observing peer writes and the collective stalls: 3474 us at 1 MB against 45 us once the load carries system scope, a ~75x regression on a 4-rank single-node all_gather. Results stay correct throughout, since the poll simply waits, so this presents as a throughput cliff rather than a wrong answer. This is a workaround, not a fix. A VMM allocation created with `hipMemAllocationTypeUncached` behaves indistinguishably from ordinary cacheable memory on this configuration, even though `hipMemGetAllocationPropertiesFromHandle` reads the property back correctly. A defect has been filed against the HIP runtime (ROCM-30384) with a standalone two-GPU reproducer that shows the same stall outside RCCL entirely, so this is not specific to the collective path. ## Technical Details `__builtin_nontemporal_load` emits `global_load_b64 ... th:TH_LOAD_NT` with no scope modifier. On gfx12-family targets that is an eviction-priority hint, not a coherence qualifier, so it does not force a miss and the poll can sit on a stale line indefinitely. `__builtin_amdgcn_global_load_b128(p, RCCL_SYSTEM_SYNCSCOPE)` emits `global_load_b128 ... scope:SCOPE_SYS`, which does. A clean-build A/B over {non-temporal, system-scope} x {load, store}, with the `.so` checksummed per variant to confirm each build actually differed, isolates the load as the entire fix: | load | store | busbw (GB/s) | |---|---|---| | non-temporal | plain | 0.0389 | | non-temporal | system-scope | 0.0384 | | **system-scope** | **plain** | **2.902** | | system-scope | system-scope | 2.884 | Only the reader's poll changes. The store side keeps the plain cacheable path from #9336, so its store throughput and register pressure are unaffected. The behavior is gated by `RCCL_LL_FIFO_SYS_SCOPE_LOAD` in `rccl_ptr.h`, restricted to gfx1250 - the only architecture where this was observed and measured, and the only one currently using cuMem P2P. Other targets compile byte-identically to before. The single macro is the removal point once ROCM-30384 is resolved. Applies to `readLL` / `readLLBeginAll` / `readLLFinish` in `prims_ll.h` and `load128NT` in `prims_ll128.h`. ## Issue Tracking JIRA ID: AICOMRCCL-2230 Upstream defect (filed): ROCM-30384 - `hipMemAllocationTypeUncached` has no observable effect on VMM allocations on gfx1250 ## Test Plan Exhaustive sweeps over message size across the priority collectives, for both LL and LL128 protocols, on gfx1250 with cuMem enabled, scaling up to 32 ranks. Each configuration is compared against a control build with the workaround disabled. ## Test Result 4-rank all_gather, 1 MB: | Proto | gate off | gate on | |---|---|---| | LL128 | 3452 us / 0.038 GB/s | 45.7 us / 2.904 GB/s | | LL | 3133 us / 0.043 GB/s | 32.4 us / 4.059 GB/s | `# Out of bounds values : 0 OK` in every arm. Non-gfx1250 targets are unchanged by construction. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
…hed VMM behavior on gfx1250 (#10969) ## Motivation The LL/LL128 comm-FIFO access policy set in #9336 rests on the FIFO being allocated uncached. That holds for the legacy IPC allocator (`hipDeviceMallocUncached`) but not for the cuMem/VMM path. On gfx1250 with `NCCL_CUMEM_ENABLE=1`, the flag poll stops observing peer writes and the collective stalls: 3474 us at 1 MB against 45 us once the load carries system scope, a ~75x regression on a 4-rank single-node all_gather. Results stay correct throughout, since the poll simply waits, so this presents as a throughput cliff rather than a wrong answer. This is a workaround, not a fix. A VMM allocation created with `hipMemAllocationTypeUncached` behaves indistinguishably from ordinary cacheable memory on this configuration, even though `hipMemGetAllocationPropertiesFromHandle` reads the property back correctly. A defect has been filed against the HIP runtime (ROCM-30384) with a standalone two-GPU reproducer that shows the same stall outside RCCL entirely, so this is not specific to the collective path. ## Technical Details `__builtin_nontemporal_load` emits `global_load_b64 ... th:TH_LOAD_NT` with no scope modifier. On gfx12-family targets that is an eviction-priority hint, not a coherence qualifier, so it does not force a miss and the poll can sit on a stale line indefinitely. `__builtin_amdgcn_global_load_b128(p, RCCL_SYSTEM_SYNCSCOPE)` emits `global_load_b128 ... scope:SCOPE_SYS`, which does. A clean-build A/B over {non-temporal, system-scope} x {load, store}, with the `.so` checksummed per variant to confirm each build actually differed, isolates the load as the entire fix: | load | store | busbw (GB/s) | |---|---|---| | non-temporal | plain | 0.0389 | | non-temporal | system-scope | 0.0384 | | **system-scope** | **plain** | **2.902** | | system-scope | system-scope | 2.884 | Only the reader's poll changes. The store side keeps the plain cacheable path from #9336, so its store throughput and register pressure are unaffected. The behavior is gated by `RCCL_LL_FIFO_SYS_SCOPE_LOAD` in `rccl_ptr.h`, restricted to gfx1250 - the only architecture where this was observed and measured, and the only one currently using cuMem P2P. Other targets compile byte-identically to before. The single macro is the removal point once ROCM-30384 is resolved. Applies to `readLL` / `readLLBeginAll` / `readLLFinish` in `prims_ll.h` and `load128NT` in `prims_ll128.h`. ## Issue Tracking JIRA ID: AICOMRCCL-2230 Upstream defect (filed): ROCM-30384 - `hipMemAllocationTypeUncached` has no observable effect on VMM allocations on gfx1250 ## Test Plan Exhaustive sweeps over message size across the priority collectives, for both LL and LL128 protocols, on gfx1250 with cuMem enabled, scaling up to 32 ranks. Each configuration is compared against a control build with the workaround disabled. ## Test Result 4-rank all_gather, 1 MB: | Proto | gate off | gate on | |---|---|---| | LL128 | 3452 us / 0.038 GB/s | 45.7 us / 2.904 GB/s | | LL | 3133 us / 0.043 GB/s | 32.4 us / 4.059 GB/s | `# Out of bounds values : 0 OK` in every arm. Non-gfx1250 targets are unchanged by construction. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
i-kosarev
added a commit
that referenced
this pull request
Oct 16, 2026
## Motivation Since #9336, the LL `storeData` writes the user output buffer through `store<>`. Without the b128 builtins (`RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS == 0`), `store<>` uses a non-temporal store. On gfx950 with `ROCM_VERSION < 70002` it also issues a release fence (`buffer_wbl2 sc0 sc1` + `s_waitcnt vmcnt(0) lgkmcnt(0)`) after every store. This makes 4 KB–64 KB `ncclBroadcast` on MI355X about 30% slower in RCCL built against ROCm 7.0.1. ## Technical Details - Removes the gfx950 `ROCM_VERSION < 70002` fence from the `store<>` fallback. - The fence is the ROCr `hipHostMallocUncached` workaround for the comm FIFO, and `storeLL` keeps it. `store<>` only writes user buffers, and its builtin path has had no fence since #3812. - Builds against ROCm 7.0.2 or later, and gfx942 builds, compile to the same code as before. ## Issue Tracking JIRA ID: ROCM-31938 ## Test Plan 4 nodes x 8 MI355X, RCCL built against ROCm 7.0.1: - `broadcast_perf` 1K–256K, interleaved with RCCL 2.27.7 and develop. - `all_reduce_perf` and `reduce_perf` 1K–4M. - Correctness with `NCCL_PROTO=LL -d all -u 1`, 8 B–1 MB, for broadcast, all_reduce, all_gather, reduce_scatter and reduce. ## Test Result `broadcast_perf` out-of-place time, ratio vs RCCL 2.27.7: | Size | develop | this PR | | --- | --- | --- | | 8K | 1.23 | 1.08 | | 16K | 1.26 | 1.08 | | 32K | 1.28 | 1.06 | | 64K | 1.31 | 1.02 | | 128K | 0.99 | 1.00 | - `all_reduce` and `reduce`: no change. - Correctness: `#wrong` = 0 for every collective, data type and size. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
1 task
alexander-sannikov
added a commit
that referenced
this pull request
Oct 30, 2026
Backported PRs: - #10670 fix(rccl): LL protocol explicit 128-bit vector loads/stores - #10339 fix(rccl): default cheap post-send fence off for gfx950 - #9336 LL: cache-bypass user buffers, plain/non-temporal comm FIFO Partially backported PRs: - #9347 [RCCL] [JIRA ID : ROCM-25479] LL128: generate separate kernels for UBR Aligned with develop: - RCCL_GFX9_CHEAP_FENCE_OFF default Co-authored-by: Cursor <cursoragent@cursor.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Retarget the LL protocol memory-access policy on gfx942/gfx950 to match the LL128 rationale: registered user buffers are cacheable and need system-scope cache-bypass, while the comm FIFO buffers are allocated uncached and gain nothing from bypass (it only adds register pressure).
Only the RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS path is retargeted; GFX11 and CUDA/PTX paths are unchanged.
Technical Details
Issue Tracking
ROCM-25479
Test Plan
Existing RCCL UT and perf testing
Test Result
Submission Checklist