rccl(feat): gfx1250 default mode tuning by isaki001 · Pull Request #10938 · ROCm/rocm-systems

isaki001 · 2026-08-29T16:47:16Z

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:

Pivot — large, cache-line-aligned messages on pivot-capable comms
GDA (RocSHMEM) — when enabled and within threshold
DDA LL / LL128 / VMM — from arch table thresholds (gfx1250 fabric)
DDA IPC — for gfx942/gfx950
CE registered — single-node, CTA_POLICY_ZERO, symmetric registered buffers
Hierarchical CE — multi-node non-LSA
CE scratch (GIN SDMA) — AlltoAll via GIN LSA/SDMA path
CE scratch — RCCL_FORCE_CE with unregistered buffers via DDA scratch
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:

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.

therock-pr-bot · 2026-08-29T16:47:45Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

Copilot

Pull request overview

This PR introduces a table-driven tuning/dispatch infrastructure for RCCL collectives, with new gfx1250 defaults and consolidated decision logic (notably for AlltoAll) so that DDA/CE/symmetric selection is consistently resolved from per-architecture thresholds rather than scattered hard-coded constants.

Changes:

Added rcclArchThresholds per-architecture threshold tables (gfx942/gfx950/gfx1250) and wired them into comm initialization.
Consolidated AlltoAll backend selection into rcclSelectAlltoAll() and extended reporting APIs so rccl-tests can label the executed backend.
Updated DDA and CE AllReduce sizing/threshold resolution to be env-overridable but table-driven by default; expanded/adjusted tests accordingly.

Reviewed changes

Copilot reviewed 30 out of 30 changed files in this pull request and generated 8 comments.

Show a summary per file

File	Description
projects/rccl/test/RcclWrapTests.cpp	Adds unit tests for tier threshold resolution and env overrides.
projects/rccl/test/DdaIpcEligibilityTests.cpp	Updates tests to source thresholds from arch tables.
projects/rccl/test/DdaAlltoAllThresholdTests.cpp	Refactors DDA enable tests to use new threshold resolvers/table values.
projects/rccl/test/common/DdaAlltoAllTestHelpers.hpp	Updates helpers to use the new VMM threshold resolver.
projects/rccl/test/common/CeAllReduceTestHelpers.hpp	Introduces a default CE AR max constant and initializes comm mock staging cap.
projects/rccl/test/CeAllReduceEligibilityTests.cpp	Updates CE AllReduce tests to use the new staging-cap model.
projects/rccl/src/rccl_wrap.cc	Adds threshold resolver APIs, unroll selection helper, and unified AlltoAll selection/reporting changes.
projects/rccl/src/init.cc	Attaches `archThresholds` to the comm during init.
projects/rccl/src/include/rccl_common.h	Extends addon algo enums, declares new APIs/params, and updates public comments/contracts.
projects/rccl/src/include/graph.h	Defines `rcclArchThresholds` and exposes `rcclGetArchThresholds()`.
projects/rccl/src/include/dda_init_detail.h	Updates documentation to reflect resolved (table/env) DDA threshold semantics.
projects/rccl/src/include/dda_alltoall.h	Adds block-count helpers for AlltoAll DDA launch reporting.
projects/rccl/src/include/comm.h	Adds `archThresholds` pointer to `ncclComm`.
projects/rccl/src/include/ce_coll.h	Removes fixed CE max macro usage and adds `ceArMaxBytes` in CE state.
projects/rccl/src/include/algorithms/CollCommon.h	Fixes index arithmetic types to avoid overflow/sign issues.
projects/rccl/src/include/algorithms/alltoall/alltoall_dda.h	Adjusts AlltoAll DDA indexing types.
projects/rccl/src/include/algorithms/alltoall/alltoall_dda_fabric.h	Adjusts fabric AlltoAll DDA indexing types.
projects/rccl/src/graph/tuning.cc	Adds gfx1250 tuning model, arch threshold tables, and arch→tuning-index mapping.
projects/rccl/src/fabric_init.cu	Sizes DDA scratch based on resolved (table/env) thresholds.
projects/rccl/src/enqueue.cc	Hardens tuner-cost-table guard for addon funcs; aligns CE AllReduce size gating with runtime cap; honors AllGather precomputed decision.
projects/rccl/src/device/generate.py	Expands gfx1250 supported unroll options.
projects/rccl/src/dda_alltoall_ipc.cu	Refactors launch-geometry computation and adds block-count helper.
projects/rccl/src/dda_alltoall_fabric.cu	Refactors launch-geometry computation and adds block-count helper.
projects/rccl/src/dda_alltoall_fabric_ll128.cu	Refactors LL128 launch geometry and adds block-count helper.
projects/rccl/src/dda_alltoall_fabric_ll.cu	Refactors LL launch geometry and adds block-count helper.
projects/rccl/src/dda_all_reduce_fabric_ll128.cu	Uses resolved LL128 threshold rather than raw env value.
projects/rccl/src/collectives.cc	Migrates AlltoAll dispatch to `rcclSelectAlltoAll()` and adjusts DDA gating helper to table-driven threshold.
projects/rccl/src/ce_coll.cc	Resolves CE AllReduce staging cap at init and uses it for layout computations.
projects/rccl-tests/src/common.h	Updates rccl-tests `ncclFunc_t` numbering/comments to match librccl ABI expectations.
projects/rccl-tests/src/alltoall.cu	Adds algo/proto/channel reporting hooks for AlltoAll, including backend-aware impl query.

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

amd-argus

ARGUS

This PR changes 2,397 lines across 37 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: the rcclUseCeAllReduce to rcclUseCeAr2Shot rename left three call sites in the two touched test files, so the unit-test target does not compile.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

This PR changes 2,399 lines across 36 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: the rcclUseCeAr2Shot rename still leaves three call sites on the old name, so the unit-test target does not compile.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Findings

1 finding(s) that fall outside the diff and could not be anchored to a line

projects/rccl/test/RcclWrapTests.cpp:2597 (High) Gfx1250_CeEligible_StillTakesDda expects DDA at 64 MiB but the new gfx1250 AllReduce entry cap is 32 MiB

This PR changes 2,403 lines across 37 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: the DDA entry gate never maxes in ddaVmmMaxGraph, so gfx1250's 256 MiB graph AllReduce tier is unreachable while fabric scratch still pays for it. The rcclUseCeAr2Shot rename is fixed; 26 earlier findings are still open.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

mch

Here's a few comments so far. Haven't got to the good part (collectives.cc, enqueue.cc) yet.

amd-argus

ARGUS

This PR changes 2,399 lines across 37 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: two gfx1250 AlltoAll threshold cases assert TRUE against a zero table entry and cannot pass, and symMinR2 withdraws symk with no guaranteed replacement.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Findings

1 finding(s) that fall outside the diff and could not be anchored to a line

projects/rccl/test/RcclWrapTests.cpp:2596 (High) Gfx1250_CeEligible_StillTakesDda asserts a 64 MiB AllReduce takes DDA, but the new entry cap is 32 MiB on gfx1250

This PR changes 2,394 lines across 37 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: Gfx1250_CeEligible_StillTakesDda picks 64 MiB against the new 32 MiB DDA entry cap, so it now fails.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

isaki001 · 2026-09-16T15:00:51Z

@argus re-review

amd-argus

ARGUS

This PR changes 3,750 lines across 46 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: rcclAllGatherCeRegisteredWindowTab is declared in no header, so rccl-UnitTestsFixturesDebug does not compile.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

This PR changes 3,908 lines across 48 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: two shipped tests fail as written (RcclWrapTests.cpp:2597 and the null comm->topo in RcclAlltoAllDecision), and three new suites are missing from test_categories_fixtures_debug.yaml.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Findings

1 finding(s) that could not be posted as inline comments

projects/rccl/test/DdaFabricEligibilityTests.cpp:761 (High) Two base-fixture AllGather at-threshold cases still use the pre-PR env defaults and now assert on sizes the resolvers reject In projects/rccl/test/DdaFabricEligibilityTests.cpp, totalCap at line 761 is still the pre-PR 64 MiB RCCL_DDA_LL128_THRESHOLD default and kDdaLLThresholdBytes at line 476 is still the pre-PR 64 KiB, while their twins under DdaFabricLL128EligibilityTest were re-pinned at line 1493. The base fixture leaves comm.archName NULL, so the resolvers take the null-table branch and return kDdaLL128BaseDefault = 32 MiB and kDdaLLBaseDefault = 32 KiB (projects/rccl/src/rccl_wrap.cc:864 and :874). ''' AllGatherLL128_AtThresholdEligible line 759 total 64 MiB vs cap 32 MiB -> rejected at dda_all_gather_fabric_ll128.cu:163 AllGatherLL_AtThresholdEligible line 558 total 64 KiB vs cap 32 KiB -> rejected at dda_all_gather_fabric_ll.cu:135 ''' Both cases assert TRUE, and DdaFabricEligibilityTest.* is registered at test_categories_fixtures_debug.yaml:28. Can lines 476 and 761 move to the kDda*BaseDefault constants the way line 1493 did?

This PR changes 3,981 lines across 54 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: rcclGetTuningIndexForArch still maps gfx1250 to tuning index 10, which is gfx110x's model, so the new tuning_model_11 is never selected and gfx110x drops to model 0.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

…nabled

…g in rcclSelect routine

…atchAvailable for checking eligiblity CE-scratch algo, fix assertion to validation rejection of target algorithm instead of assuming ring-fallback

…called

mch · 2026-09-21T18:24:23Z

Changelog check — 🔴 Missing changelog entry

This PR introduces user-facing changes but does not include a CHANGELOG.md update.

Why: This PR adds new GPU architecture support (gfx1250), introduces new per-architecture tuning infrastructure with user-visible behavior changes to algorithm/protocol selection for AllReduce, AllGather, ReduceScatter, and AlltoAll, and adds new env var controls (RCCL_CE_AR_MAX_MSG_BYTES, RCCL_CE_AR_STAGING_BYTES) with changed semantics for existing thresholds.

AI assessment: This PR adds gfx1250 tuning support to RCCL by introducing a new rcclArchThresholds per-architecture dispatch table (in rccl_arch_thresholds.h, graph.h, and tuning.cc) that governs DDA protocol tier selection, CE AllReduce sizing, symmetric kernel suppression, and per-size unroll factors. It adds a new tuning_model_10 latency/bandwidth model for gfx1250 and consolidates AlltoAll algorithm selection into a new rcclSelectAlltoAll() function. The ncclFunc_t enum in rccl-tests is renumbered/extended to match the internal librccl ABI, and AlltoAll in rccl-tests gains algo/proto/channel reporting.

Suggested entry (draft — review before using):

### Added
- Added `rcclArchThresholds` per-architecture dispatch table (`rccl_arch_thresholds.h`, `graph.h`) governing DDA protocol tier selection (`ddaLLMax`, `ddaLL128Max`, `ddaVmmMax`, `ddaVmmMaxR2`, `ddaVmmMaxGraph`), CE AllReduce sizing (`ceArMin`, `ceArMax`, `ceArRegMax`), symmetric kernel suppression (`symMaxR2`), and per-size unroll factor maps (`unrollMapAR`, `unrollMapAG`, `unrollMapRS`, `unrollMapA2A`) for gfx942, gfx950, and gfx1250.
- Added `rcclGetArchThresholds()` lookup function mapping GCN arch strings to the appropriate `rcclArchThresholds` table entry.
- Added `tuning_model_10` latency/bandwidth tuning model for gfx1250 in `tuning.cc`; gfx1250 is now included in the AINIC tuning index map via `rcclGetTuningIndexForArch()`.
- Added `rcclSelectAlltoAll()` unified AlltoAll algorithm-selection function in `collectives.cc` / `rccl_wrap.cc`, mirroring `rcclSelectAllRed()`.
- Added `ncclAllToAllDdaIpcBlocks()`, `ncclAllToAllDdaFabricBlocks()`, `ncclAllToAllDdaFabricLLBlocks()`, and `ncclAllToAllDdaFabricLL128Blocks()` query functions reporting real launch-grid occupancy for AlltoAll DDA paths.
- Added `AlltoAllGetAlgoProtoChannels()` and `AlltoAllGetCollImplInfo()` hooks in rccl-tests so AlltoAll reports algo, protocol, and channel count alongside benchmark results.

### Changed
- Extended `ncclFunc_t` enum in rccl-tests `common.h` to match the librccl internal ABI: renumbered existing values and added `ncclFuncAlltoAll` (8), `ncclFuncScatter` (9), `ncclFuncGather` (10), `ncclFuncAllToAllPivot` (11), `ncclFuncAlltoAllGda` (12), `ncclFuncAlltoAllvGda` (13), `ncclFuncAllGatherV` (14), `ncclFuncPutSignal` (15), `ncclFuncSignal` (16), `ncclFuncWaitSignal` (17), `ncclFuncAlltoAllv` (18); `ncclNumFuncs` updated to 19. **ABI-sensitive: mismatched values silently mislabel results.**
- DDA LL and LL128 eligibility checks for AllGather, AllReduce, and ReduceScatter now use per-architecture threshold helpers `rcclDdaLLThreshold(comm, func)` and `rcclDdaLL128Threshold(comm, func)` instead of the flat `rcclParamDdaLLThreshold()` / `rcclParamDdaLL128Threshold()` env-var reads.
- AlltoAll index variables in `alltoall_dda.h` and `alltoall_dda_fabric.h` widened from `int` to `size_t` to prevent overflow on large message sizes.
- DDA fabric and IPC AlltoAll launchers refactored to derive launch geometry from a single inline helper (`ddaAllToAllFabricGeom`, `ddaAllToAllFabricLLGeom`, `ddaAllToAllFabricLL128Geom`, `ddaAllToAllIpcGeom`) shared between the dispatch path and the new block-count query functions.
- `ddaFabricScratchSizing()` sizing comment and logic updated to reference `rcclDdaScratchPayloadCap()` instead of the hard-coded `DDA_THRESHOLD` default.

Confidence: medium

Rationale: The diff is large and partially truncated, but the public-facing changes are clearly identifiable from headers (rccl_arch_thresholds.h, rccl_decision.h, graph.h, dda_alltoall.h), the rccl-tests ABI enum renumbering, and new exported functions. The tuning thresholds are explicitly flagged as placeholder/not finalized, which is noted in the description but does not affect the structural API additions that warrant changelog coverage. Confidence is medium rather than high because the diff is truncated and some files (collectives.cc, rccl_wrap.cc, tuning.cc, rccl_arch_thresholds.h, rccl_decision.h) are referenced but not shown in full.

To resolve: add an entry to CHANGELOG.md before merging. If this PR genuinely does not need a changelog entry, add the label no-changelog and the check will pass.

…selection

isaki001 · 2026-09-21T18:58:02Z

Changelog check — 🔴 Missing changelog entry

This PR introduces user-facing changes but does not include a CHANGELOG.md update.

Why: This PR adds new GPU architecture support (gfx1250), introduces new per-architecture tuning infrastructure with user-visible behavior changes to algorithm/protocol selection for AllReduce, AllGather, ReduceScatter, and AlltoAll, and adds new env var controls (RCCL_CE_AR_MAX_MSG_BYTES, RCCL_CE_AR_STAGING_BYTES) with changed semantics for existing thresholds.

AI assessment: This PR adds gfx1250 tuning support to RCCL by introducing a new rcclArchThresholds per-architecture dispatch table (in rccl_arch_thresholds.h, graph.h, and tuning.cc) that governs DDA protocol tier selection, CE AllReduce sizing, symmetric kernel suppression, and per-size unroll factors. It adds a new tuning_model_10 latency/bandwidth model for gfx1250 and consolidates AlltoAll algorithm selection into a new rcclSelectAlltoAll() function. The ncclFunc_t enum in rccl-tests is renumbered/extended to match the internal librccl ABI, and AlltoAll in rccl-tests gains algo/proto/channel reporting.

Suggested entry (draft — review before using):
### Added
- Added `rcclArchThresholds` per-architecture dispatch table (`rccl_arch_thresholds.h`, `graph.h`) governing DDA protocol tier selection (`ddaLLMax`, `ddaLL128Max`, `ddaVmmMax`, `ddaVmmMaxR2`, `ddaVmmMaxGraph`), CE AllReduce sizing (`ceArMin`, `ceArMax`, `ceArRegMax`), symmetric kernel suppression (`symMaxR2`), and per-size unroll factor maps (`unrollMapAR`, `unrollMapAG`, `unrollMapRS`, `unrollMapA2A`) for gfx942, gfx950, and gfx1250.
- Added `rcclGetArchThresholds()` lookup function mapping GCN arch strings to the appropriate `rcclArchThresholds` table entry.
- Added `tuning_model_10` latency/bandwidth tuning model for gfx1250 in `tuning.cc`; gfx1250 is now included in the AINIC tuning index map via `rcclGetTuningIndexForArch()`.
- Added `rcclSelectAlltoAll()` unified AlltoAll algorithm-selection function in `collectives.cc` / `rccl_wrap.cc`, mirroring `rcclSelectAllRed()`.
- Added `ncclAllToAllDdaIpcBlocks()`, `ncclAllToAllDdaFabricBlocks()`, `ncclAllToAllDdaFabricLLBlocks()`, and `ncclAllToAllDdaFabricLL128Blocks()` query functions reporting real launch-grid occupancy for AlltoAll DDA paths.
- Added `AlltoAllGetAlgoProtoChannels()` and `AlltoAllGetCollImplInfo()` hooks in rccl-tests so AlltoAll reports algo, protocol, and channel count alongside benchmark results.

### Changed
- Extended `ncclFunc_t` enum in rccl-tests `common.h` to match the librccl internal ABI: renumbered existing values and added `ncclFuncAlltoAll` (8), `ncclFuncScatter` (9), `ncclFuncGather` (10), `ncclFuncAllToAllPivot` (11), `ncclFuncAlltoAllGda` (12), `ncclFuncAlltoAllvGda` (13), `ncclFuncAllGatherV` (14), `ncclFuncPutSignal` (15), `ncclFuncSignal` (16), `ncclFuncWaitSignal` (17), `ncclFuncAlltoAllv` (18); `ncclNumFuncs` updated to 19. **ABI-sensitive: mismatched values silently mislabel results.**
- DDA LL and LL128 eligibility checks for AllGather, AllReduce, and ReduceScatter now use per-architecture threshold helpers `rcclDdaLLThreshold(comm, func)` and `rcclDdaLL128Threshold(comm, func)` instead of the flat `rcclParamDdaLLThreshold()` / `rcclParamDdaLL128Threshold()` env-var reads.
- AlltoAll index variables in `alltoall_dda.h` and `alltoall_dda_fabric.h` widened from `int` to `size_t` to prevent overflow on large message sizes.
- DDA fabric and IPC AlltoAll launchers refactored to derive launch geometry from a single inline helper (`ddaAllToAllFabricGeom`, `ddaAllToAllFabricLLGeom`, `ddaAllToAllFabricLL128Geom`, `ddaAllToAllIpcGeom`) shared between the dispatch path and the new block-count query functions.
- `ddaFabricScratchSizing()` sizing comment and logic updated to reference `rcclDdaScratchPayloadCap()` instead of the hard-coded `DDA_THRESHOLD` default.
Confidence: medium

Rationale: The diff is large and partially truncated, but the public-facing changes are clearly identifiable from headers (rccl_arch_thresholds.h, rccl_decision.h, graph.h, dda_alltoall.h), the rccl-tests ABI enum renumbering, and new exported functions. The tuning thresholds are explicitly flagged as placeholder/not finalized, which is noted in the description but does not affect the structural API additions that warrant changelog coverage. Confidence is medium rather than high because the diff is truncated and some files (collectives.cc, rccl_wrap.cc, tuning.cc, rccl_arch_thresholds.h, rccl_decision.h) are referenced but not shown in full.

To resolve: add an entry to CHANGELOG.md before merging. If this PR genuinely does not need a changelog entry, add the label no-changelog and the check will pass.

Thank you Matt, I will update the changelog.

amd-argus

ARGUS

This PR changes 4,469 lines across 69 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: RCCL_IGNORE_ARCH_TABLE defaults to 1 (rccl_wrap.cc:98), so the gfx1250 table this PR adds is bypassed on a default run and the suites that pin it cannot pass.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

ddurnov

Please create Jira for additional configurations check
Please create Jira for refactoring of the places where it makes sense in the future iterations
In the future please split PR into smaller orthogonal chunks (e.g. pre-requisite with flexibility/configurability as first PR and check that there were no regressions before integrating new paths which use the funcs) for more fine grain review/acceptance

#10938 made AllReduce lose to SYM and a 32 MiB staging cap rejected FORCE messages that still fit the 256 MiB 2-shot window. Pipeline AllGather through staging slots and select registered CE for ZERO/FORCE so CE BEFORE and pipelined AllReduce take CE.

systems-assistant · 2026-09-24T13:51:30Z

TheRock Submodule Bump Activity

Newest first

2026-09-24 — Included in TheRock via ROCm/TheRock#8449.

## 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>

#10938 made AllReduce lose to SYM and a 32 MiB staging cap rejected FORCE messages that still fit the 256 MiB 2-shot window. Pipeline AllGather through staging slots and select registered CE for ZERO/FORCE so CE BEFORE and pipelined AllReduce take CE.

…12290) Cherry-picks: - #10938 rccl(feat): gfx1250 default mode tuning - #12051 rccl: key same-domain NET path rewrite - #12074 rccl(feat): gfx1250 tuning followup JIRA ID: ROCM-31755 --------- Co-authored-by: isaki001 <36317038+isaki001@users.noreply.github.com> Co-authored-by: isaki001 <Ioannis.Saskiotis@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Co-authored-by: Praneeth Vallem <pvallem@amd.com> 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>

#10938 made AllReduce lose to SYM and a 32 MiB staging cap rejected FORCE messages that still fit the 256 MiB 2-shot window. Pipeline AllGather through staging slots and select registered CE for ZERO/FORCE so CE BEFORE and pipelined AllReduce take CE.

…n symk is eligible (#13005) CeMPI_AllReduce fails on gfx942/gfx950 because rcclSelectAllReduce lets an eligible symk veto registered CE under CTAPolicy=ZERO. Fix: take CE under ZERO whenever it is available, as upstream and rcclSelectAllGather do, and add Max/Min tests covering both CE variants; single-process grouped multi-GPU CE still hangs as before #10938. JIRA ID: AICOMRCCL-2324

isaki001 requested review from a team and a lite review from Copilot August 29, 2026 16:47

github-actions Bot added project: rccl project: rccl-tests labels Aug 29, 2026

Copilot started reviewing on behalf of isaki001 August 29, 2026 16:48

View session

Copilot AI reviewed Aug 29, 2026

View reviewed changes

systems-assistant Bot added the organization: ROCm label Aug 29, 2026

isaki001 changed the title ~~rccl(feat): gfx1250 default mode tuning~~ draft rccl(feat): gfx1250 default mode tuning Sep 1, 2026

isaki001 force-pushed the users/isaki001/gfx1250_default_mode_tuning branch 3 times, most recently from 392f020 to bc761b6 Compare September 10, 2026 00:56

amd-argus Bot reviewed Sep 11, 2026

View reviewed changes

mch reviewed Sep 11, 2026

View reviewed changes

Comment thread projects/rccl-tests/src/common.h Resolved

mch reviewed Sep 11, 2026

View reviewed changes

Comment thread projects/rccl/src/algorithms/dda/alltoall/dda_alltoall.h Resolved

amd-argus Bot reviewed Sep 11, 2026

View reviewed changes

mch reviewed Sep 11, 2026

View reviewed changes

mch reviewed Sep 12, 2026

View reviewed changes

Comment thread projects/rccl/src/include/rccl_common.h Resolved

mch reviewed Sep 12, 2026

View reviewed changes

Comment thread projects/rccl/src/graph/tuning.cc Resolved

amd-argus Bot reviewed Sep 14, 2026

View reviewed changes

mch reviewed Sep 14, 2026

View reviewed changes

Comment thread projects/rccl/src/enqueue/enqueue.cc Resolved

amd-argus Bot reviewed Sep 16, 2026

View reviewed changes

Comment thread projects/rccl/test/RcclWrapTests.cpp Resolved

Comment thread projects/rccl/test/CeAllReduceEligibilityTests.cpp Resolved

Comment thread projects/rccl/test/CeAllReduceEligibilityTests.cpp Resolved

amd-argus Bot reviewed Sep 16, 2026

View reviewed changes

Comment thread projects/rccl/test/RcclWrapTests.cpp Resolved

Comment thread projects/rccl/test/RcclWrapTests.cpp Resolved

Comment thread projects/rccl/test/test_categories_fixtures_debug.yaml Resolved

isaki001 force-pushed the users/isaki001/gfx1250_default_mode_tuning branch from 44a6b8b to 48d3d55 Compare September 17, 2026 15:46

amd-argus Bot reviewed Sep 17, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_arch_thresholds.cc Resolved

Comment thread projects/rccl/src/rccl_arch_thresholds.cc Resolved

isaki001 changed the title ~~draft rccl(feat): gfx1250 default mode tuning~~ rccl(feat): gfx1250 default mode tuning Sep 18, 2026

isaki001 added 5 commits September 21, 2026 17:51

remove incorrect assumption that ring is fallback when symmetric is e…

52b3c55

…nabled

rename algorithm label used to test assertion in tests due to renamin…

cc32ad9

…g in rcclSelect routine

rename algorithm label used to test assertion in tests due to renamin…

5972867

…g in rcclSelect routine

add back direct check on registration instead of relying on ncclCeScr…

1880872

…atchAvailable for checking eligiblity CE-scratch algo, fix assertion to validation rejection of target algorithm instead of assuming ring-fallback

fix assignment of symmetric in algo decision before rcclSymkQuery is …

b8ccfd4

…called

isaki001 added 2 commits September 21, 2026 18:33

remove tests that are testing no-longer valid assumptions about algo …

d038fea

…selection

add missing symk hooks to tests broken by rcclSymkQuery fallthrough fix

e1d1b6a

amd-argus Bot reviewed Sep 21, 2026

View reviewed changes

Comment thread projects/rccl/test/DdaAlltoAllThresholdTests.cpp Resolved

Comment thread projects/rccl/test/RcclWrapTests.cpp Resolved

ddurnov approved these changes Sep 21, 2026

View reviewed changes

Comment thread projects/rccl-tests/src/alltoall.cu Resolved

Comment thread projects/rccl-tests/src/common.h Resolved

Comment thread projects/rccl/src/algorithms/dda/alltoall/dda_alltoall.h Resolved

mustafabar self-requested a review September 21, 2026 20:54

mustafabar approved these changes Sep 21, 2026

View reviewed changes

mch approved these changes Sep 21, 2026

View reviewed changes

isaki001 merged commit 095786f into develop Sep 22, 2026
86 of 91 checks passed

isaki001 deleted the users/isaki001/gfx1250_default_mode_tuning branch September 22, 2026 04:56

isaki001 mentioned this pull request Sep 22, 2026

rccl(feat): gfx1250 tuning pr followup #12074

Merged

1 task

AbandiGa mentioned this pull request Sep 23, 2026

fix(rccl): close ncclHierCeAvailable signature and eligibility-fixture drift #12054

Merged

1 task

thomas-huber mentioned this pull request Sep 24, 2026

fix(rccl): Don't gate hierarchical CE on hierarchical AllGather sub-communicators #12215

Closed

1 task done

sdarisala mentioned this pull request Sep 25, 2026

cherry-pick: PRs #10938, #12051, #12074 into therock-10.2-20260922 #12290

Merged

pvallem mentioned this pull request Oct 9, 2026

fix(rccl): take registered CE AllReduce under CTAPolicy=ZERO even when symk is eligible #13005

Merged

Conversation

isaki001 commented Aug 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

1. rcclArchThresholds — Per-Architecture Dispatch Table (rccl_arch_thresholds.h, rccl_arch_thresholds.cc)

2. tuning_model_11 — gfx1250 Fabric Tuning Model (tuning.cc)

3. rcclSelectAlltoAll() — Unified AlltoAll Dispatch (rccl_wrap.cc, collectives.cc)

4. New Threshold Resolver Functions (rccl_wrap.cc)

5. New rcclAddonAlgos_t Values (rccl_common.h)

6. New rcclSymkExtract Enum (rccl_common.h)

7. Adjustments to Default Behavior

Issue Tracking

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Aug 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Policy Checks Passed

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

mch left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

isaki001 commented Aug 29, 2026 •

edited

Loading

1. `rcclArchThresholds` — Per-Architecture Dispatch Table (`rccl_arch_thresholds.h`, `rccl_arch_thresholds.cc`)

2. `tuning_model_11` — gfx1250 Fabric Tuning Model (`tuning.cc`)

3. `rcclSelectAlltoAll()` — Unified AlltoAll Dispatch (`rccl_wrap.cc`, `collectives.cc`)

4. New Threshold Resolver Functions (`rccl_wrap.cc`)

5. New `rcclAddonAlgos_t` Values (`rccl_common.h`)

6. New `rcclSymkExtract` Enum (`rccl_common.h`)

therock-pr-bot Bot commented Aug 29, 2026 •

edited

Loading