rccl(feat): gfx1250 default mode tuning - #10938
Open
isaki001 wants to merge 28 commits into
Open
Conversation
…ding entry in tuning table
…ernel at particular thresholds
Stop getAlgoInfo from reading past latencies[]/bandwidths[] for funcs beyond NCCL_NUM_FUNCTIONS. Report Direct over p2p channels from rcclGetAlgoInfo, and the real backend from rcclGetCollImplInfo. Co-authored-by: Cursor <cursoragent@cursor.com>
…ric is suppressed from thresholding
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Contributor
There was a problem hiding this comment.
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
rcclArchThresholdsper-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.
Comment on lines
53
to
56
| // For alltoall, we use count for grid calculation (data per rank pair) | ||
| auto gridBlock = meta::comms::getGridAndBlockDims(count, sizeof(T), nBlocksMax); | ||
| auto gridBlock = ddaAllToAllIpcGeom(count); | ||
| const auto& grid = gridBlock.first; | ||
| const auto& block = gridBlock.second; |
Comment on lines
+49
to
52
| // Use the block cap chosen at init (barrier flag buffer is sized for it); | ||
| // count is already the byte count (kernel instantiated for int8_t). | ||
| auto gridBlock = ddaAllToAllFabricGeom(comm, count); | ||
| const auto& grid = gridBlock.first; |
Comment on lines
+1622
to
+1626
| // (6) CE scratch: unregistered buffers, RCCL_FORCE_CE, recv fits in DDA scratch. | ||
| if (rcclParamForceCe() && comm->ddaScratch != nullptr && totalBytes <= comm->ddaScratchBytes && | ||
| ncclCeScratchAvailable(comm, ncclFuncAlltoAll, ncclDevSum, datatype, ncclSymSendNonregRecvNonreg)) { | ||
| decision->algo = RCCL_CE_REGISTERED; | ||
| return ncclSuccess; |
Comment on lines
+765
to
+766
| // Per-arch DDA/CE dispatch thresholds â populated at comm init from rcclGetArchThresholds(). | ||
| // NULL on architectures without a dedicated threshold table (falls back to env-var params). |
Comment on lines
+252
to
+255
| // R2 variant: DDA VMM cap when recv buffer is registered (winRegType != | ||
| // ncclSymSendNonregRecvNonreg). On gfx1250, RS fires DDA even when | ||
| // symEligible=true; lowering this cap for R2 lets CE win at smaller sizes. | ||
| // 0 means "use ddaVmmMax" (same as R0 behavior) -- only RS needs an override. |
Comment on lines
+1806
to
+1809
| // ddaVmmMaxR2: DDA VMM cap when recv buffer is registered (R2 mode). | ||
| // RS fires DDA even when symEligible=true on gfx1250 (!symEligible || ddaFabricArch). | ||
| // Lowering this for RS lets CE win at smaller sizes for registered buffers. | ||
| // 0 means use ddaVmmMax (no R2-specific override for that collective). |
Comment on lines
+1818
to
+1820
| 0, // [7] Recv -- not used | ||
| 0, // [8] AlltoAll -- no R2-specific override (use 4 MiB) | ||
| }, |
Comment on lines
+744
to
+774
| // Context-aware VMM threshold resolver. Callers pass the full winRegType so | ||
| // the function can apply the correct override without the caller having to | ||
| // interpret registration semantics. Policy: only recv registration shifts the | ||
| // DDA VMM cap (send-only registration does not change DDA/CE dispatch). | ||
| // When the recv buffer is registered (ncclSymSendNonregRecvReg or | ||
| // ncclSymSendRegRecvReg) and the arch table has a non-zero R2 override for | ||
| // this collective, that cap wins over the default ddaVmmMax. | ||
| // When inside a graph capture (graphMode=true) and the table has a non-zero | ||
| // graph-mode override, that cap wins. Graph-mode is checked first. | ||
| // Env var (RCCL_DDA_THRESHOLD) always wins over all context variants. | ||
| size_t rcclDdaVmmThresholdCtx(const ncclComm* comm, ncclFunc_t func, | ||
| ncclSymRegType_t winRegType, bool graphMode) { | ||
| // Env var wins unconditionally -- same as rcclDdaVmmThreshold(). | ||
| size_t threshold; | ||
| if (ddaThresholdFromEnv(rcclParamDdaThreshold(), &threshold)) return threshold; | ||
| const rcclArchThresholds* table = ddaArchTable(comm); | ||
| if (table == nullptr) return 0; | ||
| // Graph-mode override: CE AR is blocked during captures; let DDA extend further. | ||
| if (graphMode) { | ||
| size_t graphCap = ddaThresholdFromTable(table->ddaVmmMaxGraph, func); | ||
| if (graphCap != 0) return graphCap; | ||
| } | ||
| // R2 override: recv buffer registered; only recv registration shifts the cap. | ||
| const bool recvReg = (winRegType == ncclSymSendNonregRecvReg || | ||
| winRegType == ncclSymSendRegRecvReg); | ||
| if (recvReg) { | ||
| size_t r2Cap = ddaThresholdFromTable(table->ddaVmmMaxR2, func); | ||
| if (r2Cap != 0) return r2Cap; | ||
| } | ||
| return ddaThresholdFromTable(table->ddaVmmMax, func); | ||
| } |
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
This branch introduces tuning infra. 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 governs DDA protocol tier selection, CE AllReduce sizing, symmetric kernel suppression, and per-size unroll factors.AlltoAll is migrated to a unified
rcclSelectAlltoAll()decision function.All existing architectures (gfx942, gfx950) retain their existing behavior through their own table entries; gfx1250 gets a fully specified set of defaults tuned from benchmark data (placeholder values flagged for sweep validation under AICOMRCCL-1756).
This branch also enables rccl-tests to report algo/protocol for alltoall.
TUNING THRESHOLDS HAVE NOT BEEN FINALIZED AND PLACE-HOLDERS STILL EXIST.
Technical Details
1.
rcclArchThresholds— Per-Architecture Dispatch Table (graph.h,tuning.cc)A new
rcclArchThresholdsstruct 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 recv buffer is registered (R2 mode); only ReduceScatter needs this on gfx1250ddaVmmMaxGraph[func]— alternative VMM cap during HIP graph capture (CE AllReduce is blocked during capture; DDA can absorb the window)ceArMin/ceArMax— CE AllReduce 2-shot staging window bounds;ceArMaxis both an allocation limit and a tuning capceArRegMax— CE AllReduce tuning cap for registered (zero-copy) path; independent ofceArMaxsymMaxR2[func]— size above which symmetric kernel is suppressed in favor of CE when recv is registered (AllReduce: 512 KiB on gfx1250)unrollMapAR/AG/RS/A2A— per-size unroll factor breakpoint tables (size → unroll index pairs); first entry withmaxBytes >= msgByteswinsTables are defined in
tuning.ccfor gfx1250, gfx950, and gfx942. The lookup functionrcclGetArchThresholds(gcn)maps GCN arch strings to the appropriate table.2.
tuning_model_10— gfx1250 Fabric Tuning Model (tuning.cc)A new latency/bandwidth tuning model (
tuning_model_10) is registered for gfx1250 and indexed inrcclGetTuningIndexForArch(). Values are based on gfx950's model (tuning_model_6) as a placeholder, pending sweep validation (AICOMRCCL-1756). 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 existingrcclSelectAllReduce/AllGather/ReduceScatterpattern. The priority chain is:rcclGetCollImplInfo()now returns the concrete backend for AlltoAll (previously not wired).rcclGetAlgoInfo()reportsRCCL_DIRECT_ALLTOALLfor generic AlltoAll queries (it has no ring/tree kernel).4. New Threshold Resolver Functions (
rccl_wrap.cc)Three public per-tier threshold resolvers replace the old
rcclDdaEnabled()per-arch default args pattern:rcclDdaLLThreshold(comm, func)— env var wins; else arch tableddaLLMax[func]; else 0rcclDdaLL128Threshold(comm, func)— same chain for LL128rcclDdaVmmThreshold(comm, func)— same chain for VMM; context-aware variantrcclDdaVmmThresholdCtx()applies R2 and graph-mode overridesA new
rcclCeArRegisteredMax(comm)resolves the registered CE cap:RCCL_CE_AR_REG_MAX_MSG_BYTESenv var wins; elseceArRegMaxfrom arch table; 0 = no cap.rcclApplyUnrollForSize(comm, func, msgBytes)applies per-size unroll breakpoints at the start of each collective's_impl()function.5. New
rcclAddonAlgos_tValues (rccl_common.h)RCCL_CE_SCRATCH— CE via DDA scratch buffer (force-CE + unregistered; previously misreported asRCCL_CE_REGISTERED)RCCL_A2A_PIVOT— AlltoAll pivot kernelRCCL_A2A_GDA— AlltoAll via RocSHMEM/GDARCCL_DIRECT_ALLTOALL— AlltoAll as per-peer Send/Recv (no collective kernel)These get
rcclGetAlgoName()labels:"A2A-Pivot","A2A-GDA","CE-Scratch","Direct".6. 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: 32 KiB (hardcoded)
After: -1 (unset); resolved from arch table
RCCL_DDA_LL128_THRESHOLD default
Before: 32 MiB (hardcoded)
After: -1 (unset); resolved from arch table
RCCL_DDA_LL128 (enable flag)
Before: 0 (disabled)
After: 1 (enabled by default)
CE AllReduce size cap
Before: NCCL_CE_AR_MAX_MSG_BYTES (compile-time constant)
After: comm->ceColl.ceArMaxBytes (set at init from arch table's ceArMax)
Channel tuning override for gfx1250
Before: Applied
After: Skipped (pending sweep data, AICOMRCCL-1756)
Protocol threshold logic for gfx1250 single-node
Before: Not applied
After: Applied (ring/tree proto thresholds now apply to gfx1250 regardless of comm->nNodes)
Channel override for gfx1250
Before: Could apply
After: Explicitly blocked
Key behavioral change: DDA LL128 is now ON by default (was off). This changes the default fast path
for mid-size messages on gfx1250 from LL -> VMM to LL -> LL128 -> VMM.
Key behavioral change: RCCL_FORCE_CE_ALLREDUCE=1 no longer overrides the ceArMaxBytes staging buffer
cap. The force flag only bypasses the CTA_POLICY_ZERO check; it cannot cause a write beyond the
allocated staging buffer.
7. Important Differences from
developDDA threshold resolution is now table-driven, not hardcoded.
The old rcclDdaEnabled(comm, totalBytes, gfx942Default, gfx950Default, gfx1250Default) signature
with inline per-call defaults is replaced by rcclDdaEnabled(comm, totalBytes, threshold) where
threshold comes from rcclDdaVmmThreshold(comm, func). Callers no longer embed arch-specific constants.
ReduceScatter DDA on gfx1250 no longer fires when symEligible=true.
Previously, the fabric arch flag ddaFabricArch unconditionally bypassed the !symEligible guard for RS.
Now RS follows the same !symEligible gate as other arches — DDA only enters when symmetric kernel is
unavailable. The ddaVmmMaxR2 mechanism handles the registered-buffer case separately (lowering the
DDA VMM cap so CE can compete).
Symmetric kernel suppression for AllReduce on gfx1250.
The new symMaxR2 field (512 KiB default) causes symEligible to be set false when recv is registered
and the message exceeds this threshold. This allows CE 2-shot or CE-registered to win over symk for
large registered AllReduces, which benchmark data shows is faster on gfx1250.
CE AllReduce for AllGather honors rcclSelectAllGather's decision in taskAppend().
A new decisionValid branch in taskAppend() routes AllGather CE dispatch based on the pre-computed
decision, eliminating decision/dispatch disagreement. RCCL_CE_SCRATCH is now a distinct enum value
(previously misreported as RCCL_CE_REGISTERED), so algo reporting accurately reflects whether CE
is using registered windows or DDA scratch.
updateCollCostTable guard updated.
The condition checking for specific AlltoAll func variants is replaced by (int)info->func >= NCCL_NUM_FUNCTIONS.
This is safer (covers future addon funcs) and matches the actual reason: comm->latencies[]/bandwidths[]
are indexed by kernel collective funcs only.
Graph-mode DDA for AllReduce is extended to 256 MiB on gfx1250 (vs. 16 MiB in non-graph mode).
During graph capture CE AllReduce is blocked, so DDA must cover the window CE would normally handle.
Issue Tracking
JIRA ID: AICOMRCCL-1758
JIRA ID: AICOMRCCL-1589
Test Plan
Test Result
Submission Checklist