Repository navigation
fix(rccl): LL/LL128: system-scope comm-FIFO load to work around uncached VMM behavior on gfx1250 - #10969
Merged
Conversation
A cuMem/VMM FIFO is cacheable to the reader despite requesting uncached, so the flag poll needs scope:SCOPE_SYS. Fixes ~75x on gfx1250 with CUMEM enabled.
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
mustafabar
changed the title
Contributor
There was a problem hiding this comment.
Pull request overview
This PR adds a gfx1250-specific workaround for LL/LL128 comm-FIFO polling stalls when the FIFO is allocated via cuMem/VMM and becomes cacheable to the reader, by switching the poll-side FIFO loads to system-scope global loads.
Changes:
- Introduce
RCCL_LL_FIFO_SYS_SCOPE_LOAD(gfx1250-only) to enable the workaround when DWORDx4 global load/store builtins are available. - Update LL128 FIFO non-temporal load path to use a system-scope
global_load_b128on gfx1250. - Update LL FIFO polling reads (
readLL,readLLBeginAll,readLLFinish) to use a system-scope 128-bit load on gfx1250, while keeping stores plain.
Reviewed changes
Copilot reviewed 3 out of 3 changed files in this pull request and generated 1 comment.
| File | Description |
|---|---|
| projects/rccl/src/include/nccl_device/rccl_ptr.h | Adds a gfx1250-only compile-time gate (RCCL_LL_FIFO_SYS_SCOPE_LOAD) for enabling system-scope FIFO loads. |
| projects/rccl/src/device/prims_ll128.h | Switches LL128 comm-FIFO load128NT to a system-scope 128-bit global load on gfx1250 to avoid stale polling. |
| projects/rccl/src/device/prims_ll.h | Adds a gfx1250-only system-scope FIFO line load helper and routes LL polling reads through it. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
PedramAlizadeh
approved these changes
Aug 31, 2026
garzaran
approved these changes
Sep 1, 2026
garzaran
left a comment
garzaran
left a comment
There was a problem hiding this comment.
Approving to unblock this, but I assume there will be a followup on this
Contributor
TheRock Submodule Bump ActivityNewest first
|
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)
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
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 withNCCL_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
hipMemAllocationTypeUncachedbehaves indistinguishably from ordinary cacheable memory on this configuration, even thoughhipMemGetAllocationPropertiesFromHandlereads 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_loademitsglobal_load_b64 ... th:TH_LOAD_NTwith 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)emitsglobal_load_b128 ... scope:SCOPE_SYS, which does.A clean-build A/B over {non-temporal, system-scope} x {load, store}, with the
.sochecksummed per variant to confirm each build actually differed, isolates the load as the entire fix: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_LOADinrccl_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/readLLFinishinprims_ll.handload128NTinprims_ll128.h.Issue Tracking
JIRA ID: AICOMRCCL-2230
Upstream defect (filed): ROCM-30384 -
hipMemAllocationTypeUncachedhas no observable effect on VMM allocations on gfx1250Test 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:
# Out of bounds values : 0 OKin every arm. Non-gfx1250 targets are unchanged by construction.Submission Checklist
🤖 Generated with Claude Code