Repository navigation
Conversation
## Motivation
This branch introduces tuning infrastructure for extended algo/proto
selection such as DDA/CE/SYMM.
It replaces scattered hard-coded thresholds and magic numbers with a
structured, per-architecture table (`rcclArchThresholds`) that supplies
message-size thresholds fro DDA protocol tier selection, CE, and
symmetric kernel.
AlltoAll is migrated to a unified `rcclSelectAlltoAll()` decision
function.
The gfx942, gfx950 architectures retain their existing behavior through
their own table entries; gfx1250 gets a fully specified set of defaults
tuned from benchmark data.
Tuning-table entries can be ignored when setting threshold-related or
algo-enablement/disablement env. variables.
This branch also enables rccl-tests to report algo/protocol for
alltoall.
## Technical Details
### 1. `rcclArchThresholds` — Per-Architecture Dispatch Table
(`rccl_arch_thresholds.h`, `rccl_arch_thresholds.cc`)
The `rcclArchThresholds` struct is defined in a new dedicated header
`src/include/rccl_arch_thresholds.h` and implemented in a new source
file `src/rccl_arch_thresholds.cc`. Table definitions for `tuning.cc`
per-arch constants continue to live in `tuning.cc`.
The struct centralizes all DDA/CE/symmetric dispatch caps per
architecture and per collective:
- `ddaLLMax[func]` — upper bound for DDA LL protocol tier (fast lane, no
GPU barrier)
- `ddaLL128Max[func]` — upper bound for DDA LL128 tier (128B lines,
mid-message)
- `ddaVmmMax[func]` — upper bound for DDA VMM/IPC (fabric simple or IPC
cap)
- `ddaVmmMaxR2[func]` — alternative VMM cap when buffers are registered
- `ddaVmmMaxGraph[func]` — alternative VMM cap during HIP graph capture
(CE AllReduce is blocked during capture)
- `ceNonRegMin[func]` / `ceNonRegMax[func]` — CE non-registered (scratch
/ 2-shot) window per collective: fires when `ceNonRegMin <= totalBytes
<= ceNonRegMax`; `ceNonRegMax = 0` disables CE-Scratch for that
collective; for AllReduce, this is a selector cap only — the staging
buffer is allocated at default and grown only when this entry (or
`RCCL_CE_AR_MAX_MSG_BYTES`) exceeds it
- `ceRegMax[func]` — CE registered-window upper bound per collective;
for AllReduce: registered CE tuning cap; for AllGather: upper bound for
CE-registered (-R 2); 0 means no upper bound
- `symMaxR2[func]` — symmetric kernel upper-bound when recv is
registered: above this size CE is faster than symk and symk is
suppressed; `kThreshUnlimited` = no suppression (symk may win at any
size); unused collectives default to `kThreshUnlimited`
- `symMaxR2Graph[func]` — graph-mode variant of `symMaxR2`; CE is
graph-unsafe so typically `kThreshUnlimited`
- `symMinR2[func]` — symmetric kernel lower-bound when buffers are
registered: 0 = no suppression
Two new support constants are introduced alongside the struct:
- `RCCL_DDA_FUNC_COUNT = ncclFuncAlltoAll + 1` — index dimension for all
per-collective arrays
- `kThreshUnlimited = SIZE_MAX` — sentinel for "no upper-bound
suppression"
Tables are defined in `rccl_arch_thresholds.cc` for gfx1250, gfx950, and
gfx942. The lookup function `rcclGetArchThresholds(gcn)` maps GCN arch
strings to the appropriate table.
### 2. `tuning_model_11` — gfx1250 Fabric Tuning Model (`tuning.cc`)
A new latency/bandwidth tuning model (`tuning_model_10`) is registered
for gfx1250 and indexed in `rcclGetTuningIndexForArch()`. The model
covers ring/tree correction factors and LL protocol ranges for
ReduceScatter, AllGather, and AllReduce. gfx1250 is now also added to
the AINIC tuning index map.
### 3. `rcclSelectAlltoAll()` — Unified AlltoAll Dispatch
(`rccl_wrap.cc`, `collectives.cc`)
AlltoAll algorithm selection is consolidated into a single
`rcclSelectAlltoAll()` function, mirroring the existing
`rcclSelectAllReduce/AllGather/ReduceScatter` pattern. The priority
chain is:
1. Pivot — large, cache-line-aligned messages on pivot-capable comms
2. GDA (RocSHMEM) — when enabled and within threshold
3. DDA LL / LL128 / VMM — from arch table thresholds (gfx1250 fabric)
4. DDA IPC — for gfx942/gfx950
5. CE registered — single-node, CTA_POLICY_ZERO, symmetric registered
buffers
6. Hierarchical CE — multi-node non-LSA
7. CE scratch (GIN SDMA) — AlltoAll via GIN LSA/SDMA path
8. CE scratch — RCCL_FORCE_CE with unregistered buffers via DDA scratch
9. Direct (p2p) fallback — default; taskAppend decomposes to per-peer
Send/Recv
`rcclGetCollImplInfo()` now returns the concrete backend for AlltoAll
(previously not wired). `rcclGetAlgoInfo()` reports
`RCCL_DIRECT_ALLTOALL` for generic AlltoAll queries (it has no ring/tree
kernel).
### 4. New Threshold Resolver Functions (`rccl_wrap.cc`)
Four public per-tier threshold resolvers replace the old
`rcclDdaEnabled()` per-arch default args pattern:
- `rcclDdaLLThreshold(comm, func)` — env var wins; else arch table
`ddaLLMax[func]`; else `kDdaLLBaseDefault` (64 KiB)
- `rcclDdaLL128Threshold(comm, func)` — same chain for LL128; else
`kDdaLL128BaseDefault` (0)
- `rcclDdaVmmThreshold(comm, func)` — same chain for VMM; else
`kDdaVmmBaseDefault` (128 MiB)
- `rcclDdaEntryThreshold(comm, func)` — context-aware combined cap:
applies R2 (`ddaVmmMaxR2`) and graph-mode (`ddaVmmMaxGraph`) overrides;
this is what callers pass to `rcclDdaEnabled()`
Additional new resolver/helper functions:
- `rcclDdaScratchPayloadCap(comm)` — computes the max across all
DDA/CE-Scratch table entries; used at init to size the DDA fabric
scratch allocation
- `rcclCeAr2ShotMax(comm)` — resolves the CE 2-shot staging cap:
`RCCL_CE_AR_MAX_MSG_BYTES` env var wins; else `ceNonRegMax[AR]` from
arch table; else 256 MiB pre-table default
- `rcclCeArGraphSafe(comm)` — gate for CE AllReduce during graph capture
- `rcclCeAllReduceEnabled(comm)` — env-var gate: `RCCL_CE_ALLREDUCE`
- `rcclForceCeAllReduceEnabled(comm)` — env-var gate:
`RCCL_FORCE_CE_ALLREDUCE`
- `rcclAllGatherCeRegisteredWindow(comm, totalBytes, ...)` — eligibility
check for AllGather CE-registered path; returns true when recv is
registered and `totalBytes <= ceRegMax[AG]`
### 5. New `rcclAddonAlgos_t` Values (`rccl_common.h`)
- `RCCL_CE_SCRATCH` — CE via DDA scratch buffer (force-CE +
unregistered; previously misreported as `RCCL_CE_REGISTERED`)
- `RCCL_A2A_PIVOT` — AlltoAll pivot kernel
- `RCCL_A2A_GDA` — AlltoAll via RocSHMEM/GDA
- `RCCL_A2A_GIN_SDMA` — AlltoAll via GIN LSA/SDMA
- `RCCL_DIRECT_ALLTOALL` — AlltoAll as per-peer Send/Recv (no collective
kernel)
These get `rcclGetAlgoName()` labels: `"A2A-Pivot"`, `"A2A-GDA"`,
`"A2A-GIN-SDMA"`, `"CE-Scratch"`, `"Direct"`.
### 6. New `rcclSymkExtract` Enum (`rccl_common.h`)
A new enum propagates the selector's symk decision into `taskAppend`,
eliminating decision/dispatch disagreement between `rcclSelect*` and the
enqueue path:
```cpp
enum rcclSymkExtract : int8_t {
RCCL_SYMK_EXTRACT_DENY = -1, // selector chose non-symk; do not extract
RCCL_SYMK_EXTRACT_NONE = 0, // no opinion; extractor uses normal logic
RCCL_SYMK_EXTRACT_ALLOW = 1, // rcclSelect* chose RCCL_SYMMETRIC; may extract
};
```
This value is written into the task by `collTaskAppend` and read by the
symmetric scheduler.
### 7. Adjustments to Default Behavior
RCCL_DDA_THRESHOLD default
Before: 128 MiB (hardcoded)
After: -1 (unset); resolved from arch table at runtime
RCCL_DDA_LL_THRESHOLD default
Before: 64 KiB (hardcoded)
After: -1 (unset); resolved from arch table
RCCL_DDA_LL128_THRESHOLD default
Before: 64 MiB (hardcoded)
After: -1 (unset); resolved from arch table
RCCL_DDA_LL128 (enable flag)
Before: 0 (disabled)
After: -1 (auto); enabled iff arch table has non-zero ddaLL128Max for
the collective
RCCL_CE_ALLREDUCE
Before: 0 (disabled by default)
After: -1 (auto/enabled by default)
CE AllReduce 2-shot size cap (RCCL_CE_AR_MAX_MSG_BYTES — new env var)
Before: NCCL_CE_AR_MAX_MSG_BYTES compile-time constant; no runtime
override
After: -1 (unset) = use ceNonRegMax[AR] from arch table; else
NCCL_CE_AR_TMPBUF_DEFAULT_BYTES (256 MiB)
CE AllReduce staging slot size (RCCL_CE_AR_STAGING_BYTES — new env var)
Before: NCCL_CE_AR_STAGING_BYTES compile-time constant (16 MiB); no
runtime override
After: -1 (unset) = use 16 MiB compile-time default; set to override
without recompiling
CE AllReduce registered path cap (RCCL_CE_AR_REG_MAX_MSG_BYTES — new env
var)
Before: no concept; registered CE AR had the same cap as 2-shot
After: -1 (unset) = use ceRegMax[AR] from arch table; 0 = no cap
Arch-table bypass (RCCL_IGNORE_ARCH_TABLE — new env var)
Before: not present
After: 1 (default); when set, arch table is skipped for all arches
without a finalized entry
Protocol threshold logic for gfx1250 single-node
Before: Not applied
After: Applied (ring/tree proto thresholds now apply to gfx1250
regardless of comm->nNodes)
Key behavioral change: DDA LL128 is now auto-enabled for gfx1250 (was
unconditionally off). For
architectures with a non-zero ddaLL128Max in their arch table, the
mid-size fast path becomes
LL -> LL128 -> VMM instead of LL -> VMM.
Key behavioral change: RCCL_FORCE_CE_ALLREDUCE=1 no longer overrides the
ceAr2ShotMax staging buffer
cap. The force flag only bypasses the CTA_POLICY_ZERO check; it cannot
cause a write beyond the
allocated staging buffer.
## Issue Tracking
JIRA ID: AICOMRCCL-1758
JIRA ID: AICOMRCCL-1589
## Test Plan
gfx 1250 single-node sweep
gfx942 1,2,4,8,16 node sweep
gfx950 1,2,4,8,16 node sweep
The following test cases were removed because the behavior they asserted
no longer exists in the new
dispatch model. Each deletion is intentional and documented below.
## Test Result
gfx950 and gfx1250 tests passed.
run-to-run variation observed on gfx942 that is not attributed to this
PR.
## Submission Checklist
- [ ] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
---------
Co-authored-by: isaki001 <Ioannis.Saskiotis@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
## Motivation key same-domain NET path rewrite on the physical device ## Technical Details rcclRewriteSameDomainNetPaths() promoted PHB to PXB on the GPU nodes, but every gfx1250 GPU is an MLOPart partition and ncclTopoGdrDistance() reads the parent DEV entry for those. The verdict therefore stayed on the original PHB, which is above the default netGdrLevel, so ncclTopoCheckGdr() refused GDR and the no-GDR diversion in ncclTopoComputePaths() rewrote the partition entry back to PHB, erasing the promotion that had just been made. Walk the DEV nodes instead and mirror the result onto the partitions attached to each one, so the GDR decision and the graph search read the same distance. The partition index is overlaid on the DEV id, so mask it off before reading the PCI domain. ## Issue Tracking JIRA ID : ROCM-30906 ## Test Plan Add TopoTests coverage for the promotion and for its gfx1250-only scope; the first fails against the previous partition-keyed rewrite. ## Test Result On a 4-NIC gfx1250 node this takes the GPU->NET matrix from uniformly PHB to PXB on each GPU's own domain, which enables GDR and restores NIC locality. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Praneeth Vallem <pvallem@amd.com>
## Motivation This is a follow-up to [PR 10938](#10938) which introduced new tuning infra. This PR updates the changelog and documentation and turns the gfx1250 tuning by default by setting the `RCCL_IGNORE_ARCH_TABLE` env. variable to `0` by default. It also addresses some of the review-comments of PR 10938 that were unaddressed. #10938 (comment) #10938 (comment) #10938 (comment) #10938 (comment) #10938 (comment) ## Technical Details `IGNORE_ARCH_TABLE` is set to `0` by default, which enables the new gfx1250 tuning thresholds. This change doesnt impact gfx942 and gfx950 which were utilizing the table regardless, since their entries reflected the default thresholds as they were before PR 10938. Review-comments addressed from PR are related to documentation, and adjusting tests for correctness. ## Issue Tracking JIRA ID: AICOMRCCL-2422 ## Test Plan gfx942, gfx950 1,2,4,8,16 tests gfx1250 4 GPU tests ## Test Result ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: amd-argus[bot] <321830313+amd-argus[bot]@users.noreply.github.com>
github-actions
Bot
added
documentation
wenkaidu
approved these changes
Sep 25, 2026
Contributor
|
@cppccppccppc @sunguobao PTAL |
isaki001
approved these changes
Sep 25, 2026
astyrrian1
approved these changes
Sep 28, 2026
ScottTodd
approved these changes
Sep 28, 2026
ScottTodd
merged commit Sep 28, 2026
4fbbb31
into
release/bkc/therock-10.2-20260922
67 of 71 checks passed
ScottTodd
added a commit
to ROCm/TheRock
that referenced
this pull request
Sep 29, 2026
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.
Cherry-picks:
JIRA ID: ROCM-31755