Repository navigation
[RCCL] Pick the gfx1250 tuning table by compute-partition mode (RCCL_TUNING_PARTITION_MODE) - #13416
Merged
Merged
Conversation
…RCCL_TUNING_PARTITION_MODE) The gfx1250 arch table was tuned on SPX. On DPX several of its crossovers are wrong in the other direction, so one table cannot serve both modes. This adds the plumbing for a per-mode table; the DPX values come in a follow-up. - rcclGetArchThresholds() takes an rcclTuningPartition (SPX by default). gfx1250 gets a full DPX table, for now an exact copy of the SPX table, so tuning is unchanged. gfx942 and gfx950 keep one table for every mode. - RCCL_TUNING_PARTITION_MODE = auto (default) | SPX | DPX. auto follows the compute-partition mode the GPU reports in sysfs; any mode other than DPX (CPX, or none reported) uses SPX. An unknown value is reported once at ATTN and treated as auto. It selects tuning thresholds only, and has no effect under RCCL_IGNORE_ARCH_TABLE=1. - fillInfo reads current_compute_partition once (hoisted out of the MLOPart probe, which reuses it) and records each rank's request in a new ncclPeerInfo field. After AllGather1 the comm uses DPX only when every rank asks for it, SPX otherwise, so all ranks read the same table; rank 0 reports a mixed comm at ATTN and the chosen table at INFO. Tests: init microtests for the per-rank request (GPU mode, env override, auto/empty, invalid value reported once), the per-comm rule, the table installed in comm->archThresholds and its logging, and the initTransportsRank call site. The rcclGetArchThresholds host fake can now hand out a table per partition so tests see which one a comm got. Every branch is mutation-gated. GPU test TuningPartitionTablePerArch checks the DPX table is a separate copy of SPX and that gfx942/gfx950 ignore the mode. CHANGELOG, env-variables.rst and the microtest fakes map updated. JIRA: AICOMRCCL-1756
github-actions
Bot
added
documentation
✅ All Policy Checks Passed
📖 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
approved these changes
Oct 13, 2026
Contributor
|
No undocumented RCCL or NCCL environment-variable changes. |
This reverts commit 8b9302e.
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
|
|
||
| // Calls to rcclGetArchThresholds and the partition the last one asked for. | ||
| extern int g_archThresholdsCalls; | ||
| extern rcclTuningPartition g_lastArchThresholdsPartition; |
Contributor
There was a problem hiding this comment.
Low
g_lastArchThresholdsPartition has no reader anywhere in the tree: rccl_arch_thresholds_fakes.cc:20 writes it, :26 resets it, and nothing asserts on it. The gap it was added for is real, since the provisional SPX install at init.cc:3157 is not asserted in any host test. Could we add EXPECT_EQ(RCCL_TUNING_PARTITION_SPX, g_lastArchThresholdsPartition) beside the sibling ll128 assertions at init-test.cc:7910, which already drives that rung and stops before rcclSetTuningPartition runs?
mustafabar
approved these changes
Oct 14, 2026
mustafabar
left a comment
mustafabar
left a comment
Contributor
There was a problem hiding this comment.
The PR uses existing infra, but heavily uses AI to get there, and is a manifestation of overfitting using C++ code as it relies on the data collected at that time on the environment. My stance is conservative with respect to what thresholds to add and whether it really pays off to introduce a new arch thresholds knobs. This is for developer's sanity, maintenance and code hygiene. Also, we understand there are dips and performance bugs for some combos, and typically, we don't hide these and go fix perf instead of working around them. In an upcoming PR, please refine this approach so that we don't end up with a massive tech debt. I am approving with the notes above.
Contributor
|
pls check CI issues and watch out for perf regressions CI flags before merging |
sdarisala
approved these changes
Oct 14, 2026
sdarisala
left a comment
sdarisala
left a comment
Contributor
There was a problem hiding this comment.
LGTM. If we can avoid the env variable, that will be better IMO.
sdarisala
reviewed
Oct 14, 2026
sdarisala
left a comment
sdarisala
left a comment
Contributor
There was a problem hiding this comment.
Make tuning table an external config, not a static table embedded into the binary.
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: CHANGELOG.md:27 says the DPX table is a copy of SPX and tuning is unchanged, but the table now carries 14 re-tuned values that auto applies on DPX nodes by default.
Comment @argus review to re-run, or @argus re-review after pushing.
| * Hierarchical symmetric `ncclAllGather` for multi-node using GIN and LSA (`AllGather_RailRing_LsaST`): a GIN ring over the rails team moves data between nodes and LSA stores broadcast it to the other GPUs in each node. | ||
| * nccl4py: per-call collective configuration (`NCCLCollConfig`, `VendorOption`) including the launch completion event, communicator properties (`NCCLCommProperties`), and the `GIN_ONLY` window flag. On ROCm, the HIP `Event` shim provides only an event handle; `record()`, `sync()` and `query` are not implemented. | ||
| * `RCCL_TUNER_EMBEDDED_CONFIG` (default `1`): the built-in CSV tuner configurations in `tuner/*.csv` are now compiled into `librccl.so` and used as the stock tuning defaults, so a cluster no longer has to stage the CSV on every node or on NFS. Set `0` to ignore the embedded configurations. `NCCL_TUNER_CONFIG_FILE` and the on-disk tuner directories still take precedence, and `NCCL_TUNER_PLUGIN=none` still disables the tuner entirely. | ||
| * `RCCL_TUNING_PARTITION_MODE` (`auto`, `SPX` or `DPX`; default `auto`): selects the gfx1250 tuning table for the GPU's compute-partition mode. `auto` uses the DPX table when the GPU reports DPX, and the SPX table otherwise. A communicator uses the DPX table only when every rank selects it. The DPX table is currently a copy of the SPX table, so tuning is unchanged. |
Contributor
There was a problem hiding this comment.
Medium
The RCCL_TUNING_PARTITION_MODE entry says the DPX table "is currently a copy of the SPX table, so tuning is unchanged", but rcclArchThresholds_gfx1250_dpx differs in 14 fields (e.g. ddaLLMax[AG] 512K vs 128K, ddaVmmMax[RS] 16M vs 0, agRingLL128MaxPerRank 16M vs 64M), and RcclWrapTests.cpp:3114 asserts they differ. With the default auto, DPX nodes get the new tuning without opting in. Could we update the entry (and the PR description)?
Suggested change
| * `RCCL_TUNING_PARTITION_MODE` (`auto`, `SPX` or `DPX`; default `auto`): selects the gfx1250 tuning table for the GPU's compute-partition mode. `auto` uses the DPX table when the GPU reports DPX, and the SPX table otherwise. A communicator uses the DPX table only when every rank selects it. The DPX table is currently a copy of the SPX table, so tuning is unchanged. | |
| * `RCCL_TUNING_PARTITION_MODE` (`auto`, `SPX` or `DPX`; default `auto`): selects the gfx1250 tuning table for the GPU's compute-partition mode. `auto` uses the DPX table when the GPU reports DPX, and the SPX table otherwise. A communicator uses the DPX table only when every rank selects it. The DPX table carries crossovers re-measured on H13 DPX nodes, so algorithm and protocol choices on DPX nodes can differ from earlier releases; set `SPX` to keep the previous tuning. |
Re-measured PR 13416's DPX table (921eb0e) out of the box on DPX nodes at 8, 16 and 32 ranks (rccl-tests -n 10 -w 20, median of 3 rounds; 16/32 ranks across the scale-up fabric) and compared each size against the fastest path in the 2026-10-11 DPX sweep. Only entries where a different path won change; the SPX table is untouched. AllGather - ddaLLMax 128K -> 512K total (ddaLLMaxPerRank 16K -> 0): LL beats LL128 1.4-1.6x at 256K-512K on 8 ranks. The crossover falls as ranks grow (1M/512K/~128K at 8/16/32), so a per-rank cap cannot follow it. - ddaLL128Max 128M -> 64M total (ddaLL128MaxPerRank 16M -> 0): Ring beats DDA LL128 1.1-1.4x at 128M-256M on 8/16/32 ranks. - agRingLL128MaxPerRank 64M -> 16M: Ring Simple beats Ring LL128 1.2-1.4x at 256M-512M (8 ranks), 1G (16) and 2G (32). AllReduce - arRingLL128MaxPerRank 32M -> 16M: Ring Simple beats Ring LL128 1.1-1.3x at 256M (8), 512M (16) and 1G (32). - arTreeLLMaxPerRank 32K -> 256K: Tree LL beats Tree LL128 1.3-2.1x to 2M/4M/8M at 8/16/32 (fp8). - ddaArVmmFirstMaxPerRank 512K -> 256K: LL two-shot beats VMM 1.2-1.3x at 16M on 32 ranks. AlltoAll - ddaLLMax 64K -> 2M: LL beats Direct 1.2-4.3x at 128K-2M on 8/16/32 ranks. ddaLL128Max is left as is; with LL above it the LL128 tier is no longer reached for AlltoAll. ReduceScatter - ddaLLMaxPerRank 256K -> 1M: LL beats DDA Simple / Ring 1.1-1.5x to 8M/16M/32M at 8/16/32 (float/bf16, DDA LL128 off). - ddaVmmMaxPerRank 2M -> 0, ddaVmmMax 0 -> 16M total: VMM still takes 16M on 8 ranks, where it wins, while Ring LL128 beats it 1.3-1.5x at 32M/64M on 16/32 ranks. Below 8 ranks the 16M LL total covers the same range. - rsRingLL128MinPerRank 512K -> 1M: Ring LL beats Ring LL128 1.1-1.2x at 16M/32M on 16/32 ranks (fp8). - rsRingLL128MaxPerRank1B 32M -> 16M: Ring Simple beats Ring LL128 1.2-1.3x at 256M/512M/1G on 8/16/32 (fp8). No DDA LL128 threshold is raised for AllReduce, AlltoAll or ReduceScatter. rankScaledMaxBytes stays within the largest total (AllReduce LL128, 256M). TuningPartitionTablePerArch now expects the DPX table to differ from SPX and spot-checks three DPX-only entries. Not covered by this change (needs code or more data): - fp8 AllReduce on 16/32 ranks stays on Tree LL128 from 8M-16M up; Ring is 1.1-1.9x faster. Only a Ring -> Tree switch exists (arTreeMaxPerRank), not the reverse. - AllReduce 64M on 16 ranks and 128M-256M on 32 ranks: Ring LL128 beats DDA LL128 1.1-1.2x, but lowering the 256M DDA LL128 total regresses 128M on 16 ranks and would put rankScaledMaxBytes above the largest total. - AllGather 1M on 8 ranks (LL 1.24x) would need a cap that regresses 1M on 16 ranks. - 64-rank effects are inferred from the sweep only (no 64-rank nodes): the AllGather 64M LL128 cap trades 128M for 256M there. Co-Authored-By: Claude <noreply@anthropic.com>
Conflict in projects/rccl/test/host/MICROTEST_README.md (fakes table): keep develop's reworded src/ras/ras_param.cc row and this branch's src/rccl_arch_thresholds.cc row. Co-Authored-By: Claude <noreply@anthropic.com>
pvallem
force-pushed
the
users/saverine/rccl_spx_dpx_flag
branch
from
October 15, 2026 15:29
6bf78f2 to
d57e7b2
Compare
| const rcclArchThresholds* rcclGetArchThresholds(const char* gcn); | ||
| // Compute-partition mode a gfx1250 comm is tuned for (RCCL_TUNING_PARTITION_MODE). It picks | ||
| // the table only; other archs have one table for every mode. | ||
| enum rcclTuningPartition { |
Contributor
There was a problem hiding this comment.
enum struct would have been better, but ok.
Comment on lines
+1355
to
+1360
| // from sysfs entirely, so ask function 0 rather than our own BDF. | ||
| char physBusIdStr[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE]; | ||
| char partition[MAX_STR_LEN]; | ||
| partition[0] = '\0'; | ||
| if (int64ToBusId(info->busId & ~0xfLL, physBusIdStr) == ncclSuccess) { | ||
| (void)ncclOsGetPciDeviceComputePartitionByBusId(physBusIdStr, partition, sizeof(partition)); |
Contributor
There was a problem hiding this comment.
This code is gross but inherited I guess.
Contributor
TheRock Submodule Bump ActivityNewest first
|
astyrrian1
pushed a commit
that referenced
this pull request
Oct 19, 2026
…TUNING_PARTITION_MODE) (#13416) The gfx1250 arch table (`rcclArchThresholds`) was tuned in SPX mode. The DPX sweep disagrees with it in several places: crossovers that are right for SPX are wrong in the other direction for DPX, and the reverse. One table cannot serve both modes, so this PR adds a way to pick the table by compute-partition mode. This PR is plumbing only. The DPX table is an exact copy of the SPX table, so tuning is unchanged. The DPX values come in a follow-up PR. **New env var `RCCL_TUNING_PARTITION_MODE`** (`auto` default | `SPX` | `DPX`, case-insensitive) - `auto` follows the compute-partition mode the GPU reports (sysfs `current_compute_partition` on the physical function). DPX picks the DPX table; SPX, CPX or no reported mode picks SPX. - `SPX` / `DPX` force a table. - An unknown value is reported once at ATTN and treated as `auto`. - It only selects tuning thresholds. It has no effect on other archs or under `RCCL_IGNORE_ARCH_TABLE=1`. **Table selection** - `rcclGetArchThresholds(gcn, partition)` takes an `rcclTuningPartition`, defaulting to SPX, so existing callers are unchanged. - gfx1250 has a full second table, `rcclArchThresholds_gfx1250_dpx`. gfx942 and gfx950 return the same table for both modes. **Every rank uses the same table** - Ranks choose algorithms independently, so a comm where ranks read different tables could pick different algorithms for one call. - `fillInfo` records each rank's request in a new `ncclPeerInfo::tuningPartition`, appended at the tail. - After AllGather1, `initTransportsRank` re-picks `comm->archThresholds`: DPX only when every rank asks for DPX, otherwise SPX on every rank. - Rank 0 logs a mixed comm at ATTN and the chosen table at INFO (`gfx1250 tuning table: SPX|DPX`). - The sysfs read is hoisted out of the MLOPart probe. One read now serves both uses, and MLOPart behaviour is unchanged. **Host microtests (`rccl-UnitTestsMicroInit`)** - Per-rank request: GPU mode SPX/DPX/CPX/none, env override, `auto` and empty value, invalid value reported once (isolated test), and a function above the MLOPart bound. - Per-comm rule: DPX only when every rank asks for it. - The table installed in `comm->archThresholds`, on rank 0 and other ranks, for gfx1250 and other archs, and its logging. - The `initTransportsRank` call site, including a comm where only one peer asks for SPX. - The `rcclGetArchThresholds` host fake can now return a separate table per partition, so tests check which table the comm got rather than only which partition was asked for. - Mutation-gated on `init.cc`. **GPU unit test** - `RcclDdaTierThresholds.TuningPartitionTablePerArch` checks: - the gfx1250 DPX table is a separate object with the same contents as SPX; - explicit SPX matches the default; - gfx942/gfx950 return one table for both modes. **On H13 nodes** - SPX node: INFO reports `SPX`, and the paths chosen are identical to develop. - DPX node: INFO reports `DPX`, and the paths chosen are identical to develop (the tables are equal). - `RCCL_TUNING_PARTITION_MODE=DPX` on an SPX node and `SPX` on a DPX node force the table. - An invalid value produces one ATTN. - [x] Host microtests: all 13 binaries pass. - [x] Mutation gate on `init.cc` (23 mutants) - [x] `rccl_check_default_behavior.py`: no undocumented env-var changes. - [x] GPU unit tests on gfx1250 - [x] SPX and DPX node checks on H13 - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md - [x] CHANGELOG.md updated (Added) - [x] docs/api-reference/env-variables.rst updated --------- Co-authored-by: Praneeth Vallem <Praneeth.Vallem@amd.com> Co-authored-by: Claude <noreply@anthropic.com> (cherry picked from commit f710804)
pmoutsias-amd
added a commit
that referenced
this pull request
Nov 3, 2026
…the RCCL 2.32.3 for ROCm 10.3.0 (Unreleased) section (#14929) This PR does two things to `projects/rccl/CHANGELOG.md`. 1. Removes the duplicate `## RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` heading. #13544 removed it, and #13128 brought it back. The section is not a pure copy: 48 of its entries have no identical copy elsewhere. 47 of them are folded into `## RCCL 2.32.3 for ROCm 10.2.0` instead of being deleted, and the 48th is the older wording of the `RCCL_CE_AR_STAGING_BYTES` entry, which is dropped in favor of the newer one. Where the same entry appeared in both sections, one copy is kept. 2. Moves the entries below, which sit in release sections but whose code was merged to develop after the `rocm-systems` commit that TheRock `release/therock-10.2` pins (`ded49fd`) and is not in that commit. This moves 53 of them into `## RCCL 2.32.3 for ROCm 10.3.0 (Unreleased)`, under the matching subsections. Not moved: 10 entries that were attributed to #13544 (the CHANGELOG-only PR that last touched them), so the code PR could not be traced. They stay in the 10.2.0 section and are unverified, so the RCCL owners should confirm them. They begin: `RCCL_TDM_LL128_ENABLE`, `Ll128TdmStoreTest`, the TDM reduction path for SIMPLE, Anvil SDMA strong GIN signals, VA GIN signals on Anvil SDMA, `NCCL_P2P_DISABLE=1` on shared topology nodes, queue evictions at process exit, multi-receive `irecv` with `NCCL_IB_USE_INLINE=1`, Anvil SDMA signal cells across logical contexts, and GIN signal shadow sizing. - #13328 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_BCM_LINKS_PATH`: fallback directory path for Broadcom PEX890xx multi-host switch inter-switch-link topology. On systems where a single physical PEX switch is partitioned across multiple CPU domains (upstream ports), RCCL needs to know which ports share the same switch so it can classify cross-partition PCIe paths as `PXB` instead of `PHB` and enable GDR across them. The `switch_discovery` kernel module exposes this via sysfs (`/sys/kernel/pci_switch_link/virtual_switch_links`); when that path is not available, RCCL reads the same directory layout from `RCCL_BCM_LINKS_PATH`. The companion userspace tool `rccl_bcm_isl_discover.sh` discovers the switch groups by reading the PCIe Device Serial Number (DSN) from extended config space (vendor `1000`, device `c030`) and populates this directory. - #12742 (merged 2026-10-18, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_MLOPART`: allows GIN on MLO/CPX partitions. A partition reaches xGMI and the network over the physical device's paths, so GIN is as available to a partition as to the whole GPU. By default a single-node partitioned communicator keeps GIN, and one spanning several nodes does not, because past one node admitting a partition also reaches the hierarchical Copy Engine path. Set `1` to allow GIN on partitions on any number of nodes, or `0` to restore the previous partition exclusion. - #3894 (merged 2026-10-16, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_TREE_COMPACT` (default `0`): builds the inter-node trees with a half-interleave node remapping, for jobs whose nodes are ordered by network domain. Both trees keep cross-domain hops at the leaves only when the nodes split evenly into two domains; with an uneven split, the first tree does so when the first domain holds at least half of the nodes, and the second tree still crosses domains. All ranks in a communicator must use the same value; initialization fails with `ncclInvalidUsage` otherwise. - #12856 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_FABRIC_TELEMETRY_ENABLE` (default `0`): samples every amd-smi fabric telemetry category for the GPUs this process opens communicators on and logs the counters that moved, plus a closing report over the whole sampled span when a device's last communicator is destroyed. Diagnostic only, and off by default. Also needs `RCCL_USE_AMD_SMI_LIB=1` and `render`-group access to the IFoE config character device; without either, RCCL logs why telemetry was skipped and initialization continues. - #12856 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_FABRIC_TELEMETRY_INTERVAL_MS` (default `1000`): sampling period for `RCCL_FABRIC_TELEMETRY_ENABLE`, clamped to a 100 ms floor and a one-hour ceiling. A non-positive value disables telemetry entirely. - #12573 (merged 2026-10-14, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Hierarchical symmetric `ncclAllGather` for multi-node using GIN and LSA (`AllGather_RailRing_LsaST`): a GIN ring over the rails team moves data between nodes and LSA stores broadcast it to the other GPUs in each node. - #14140 (merged 2026-11-01, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Hierarchical symmetric `ncclAllGather` for partition mode in gfx1250 scale-up domain using LSA (`AllGater_HierLsaST`): each rank stores its data once to its counterpart partition on every other GPU, which forwards it to the other partitions there, so data crosses to each GPU once instead of once per partition. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_ALLREDUCE_GFX1250_ENABLE` (default `0`): opt in to GinScatter AllReduce on gfx1250 for symmetric sum messages from 64 MiB through 4 GiB. Left off, gfx1250 keeps the pre-existing AllReduce selectors. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_ALLREDUCE_SCRATCH_BYTES` (default 512 MiB): GinScatter incoming scratch. `0` skips the allocation and GinScatter is not selected. The buffer is allocated only on the GinScatter path. - #13416 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_TUNING_PARTITION_MODE` (`auto`, `SPX` or `DPX`; default `auto`): selects the gfx1250 tuning table for the GPU's compute-partition mode. `auto` uses the DPX table when the GPU reports DPX, and the SPX table otherwise. A communicator uses the DPX table only when every rank selects it. The DPX table is currently a copy of the SPX table, so tuning is unchanged. - #13727 (merged 2026-11-03, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_IB_UDMA_POLICY` (default `0`): selects how QPs are assigned to the two uDMA pipelines of an AINIC with the `IB-CAST` transport. `0` pins each channel's send QPs, and separately its receive QPs, to one pipeline per device, `1` round-robins P2P QPs, `2` leaves the choice to the driver, and `3` splits each connection across both pipelines (the default of two QPs per connection is enough; splitting is lost with one QP per connection and on RMA comms). Out-of-range values log a warning and fall back to `0`. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_AR_2SHOT_MIN_BYTES` (default `-1`): overrides the minimum message size for the CE 2-shot AllReduce path. When `-1`, resolved from `ceNonRegMin[AllReduce]` in the per-arch table. - #12310 (merged 2026-10-28, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_REDUCESCATTER` (default `0`): opt-in Copy Engine ReduceScatter on a single-node symmetric communicator with `CTA_POLICY_ZERO`. Always stages through `ceARTmpBuf`; user-buffer registration is not required. No message-size cap (large messages are pipelined). - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_COLL_MAX_BYTES` (default `-1`): overrides the CE registered-window and 2-shot size cap for all non-AllReduce collectives (AllGather, AlltoAll, etc.). When `-1`, resolved from `ceRegMax[func]` or `ceNonRegMax[func]` in the per-arch table. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_COLL_MIN_BYTES` (default `-1`): overrides the minimum message size for CE non-AllReduce paths (both registered-window and scratch). When `-1`, resolved from `ceNonRegMin[func]` in the per-arch table. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_SYM_K_MAX_BYTES` (default `-1`): overrides the symmetric-kernel upper bound for all collectives. When `-1`, resolved from `symMaxR2[func]` (eager) or `symMaxR2Graph[func]` (graph-capture) in the per-arch table. `kThreshUnlimited` (SIZE_MAX) in the table means no suppression. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_SYM_K_MIN_BYTES` (default `-1`): overrides the symmetric-kernel lower bound for all collectives. When `-1`, resolved from `symMinR2[func]` in the per-arch table. Below this size, DDA is preferred over symk. - #13995 (merged 2026-10-27, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_IB_GDR_FLUSH_FLAG` (default `0`): replaces the GPUDirect RDMA flush of Simple-protocol receives in the `IB` network transport. Instead of the proxy fencing each receive with an RDMA read and waiting for it, the sender writes a per-QP flag into GPU memory after the payload without relaxed ordering, and the GPU waits for it. A communicator uses it only when every rank enables it and can use it (gfx942 or gfx950, the `IB` transport, and no adaptive routing, out-of-order receive queues, ID-based receive matching (`NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=1`) or port-failover resiliency (`NCCL_IB_RESILIENCY_PORT_FAILOVER`)); otherwise every rank keeps the RDMA read flush. Connections over a fused NIC (NIC Fusion, including a fused vNIC chosen by subnet-aware routing), with more than 8 QPs per request, or that cannot register the flag memory also keep it. On gfx950 with 2 to 8 nodes of 8 GPUs, such a communicator also selects Ring/LL128 from 16 MB per node, Tree/Simple from 32 MB per node and Ring/Simple from 64 MB per node for `ncclAllReduce`, and switches `ncclReduceScatter` from LL128 to Simple above 6 MB per rank instead of 8 MB, unless `NCCL_ALGO` or `NCCL_PROTO` is set. On MI355X it raises large `ncclAllReduce` and `ncclReduceScatter` bandwidth by 10 to 22% and 1 to 8 MB `ncclAlltoAll` by 4 to 12%. - #13995 (merged 2026-10-27, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_IB_GDR_FLUSH_FLAG_TUNE` (default `0`): additional tuning for communicators that use `RCCL_IB_GDR_FLUSH_FLAG` on gfx950 with 8 GPUs per node on 2 to 8 or 16 nodes. Ring and tree connections get a 4x LL128 buffer (send/recv keeps the default, and `NCCL_LL128_BUFFSIZE` disables it); `ncclAllReduce` selects Ring/LL128 from 12 MB per node up to what one ring loop carries, then Tree/Simple up to 64 MB per node and Ring/Simple on 48 channels up to 128 MB per node, and Tree/Simple from 8 to 64 MB per node on 16 nodes; `ncclReduceScatter` keeps LL128 up to 7 MB per rank, then Simple, on 16 nodes too. All ranks must use the same value. - #13757 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_PCIE_ORDERED_RING` (default `1`, also accepted as `NCCL_PCIE_ORDERED_RING`): pins the ring graph search to the order GPUs appear in the topology array, so every channel uses the same ring. Without it the search picks a different permutation per channel, so a device exchanges with many different peers; pinning the order reduces fan-out to one peer per direction per device across all rings, which cuts the connections, FIFO buffers and setup work each device carries. On gfx1250 partitioned (DPX) GPUs the unpinned search also mixes scale-up and scale-in hops between channels; pinning keeps the two partitions of a die adjacent, though ring throughput is bounded by its slowest hop either way. This constrains GPU visitation order only: path validation, bandwidth constraints and backtracking are unchanged. It applies to gfx1250 ring graphs on a single host where every rank is a GPU of that topology, so trees, CollNet, other architectures and multi-host communicators are unaffected, and `NCCL_GRAPH_FILE` and `NCCL_RINGS` still take precedence. Set `0` to restore the unconstrained search, or `2` to fail communicator initialization instead of falling back when no ordered ring passes validation. - #10885 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_MULTI_SEGMENT_REGISTER` default changed `1` -> `0` on HIP and stays `1` on other builds. On HIP, multi-segment registration stays off until a job sets the variable. Set `1` to register one DMA-BUF memory region per physical segment where the transport supports it. - #10885 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * NET registration counts physical segments over the whole registered buffer, not only the range a collective touches. With `NCCL_MULTI_SEGMENT_REGISTER=0`, a collective on any part of a multi-segment registration skips NET registration and stages, instead of retrying a single-region registration on every collective. - #10891 (merged 2026-10-25, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * GIN window registration on HIP fails with `ncclInvalidUsage` when the window spans several device-memory segments, instead of registering a memory region that covers only the first segment, and when it contains host VMM memory and `NCCL_DMABUF_ENABLE=0`. - #12947 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `install.sh --all_unrolls` (`-DBUILD_ALL_UNROLLS=ON`) now adds the unroll factors the multi-arch default skips (8 and 16) on top of each architecture's existing set. It no longer replaces that set or clears the gfx1250 pin on unroll 32. The flag stays off by default, so a normal build does not compile unroll 8 or 16. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * The GIO topology model is applied only to systems with exactly 16 GPUs. It was previously applied to any system with 16 or more GPUs, where the 16-GPU model does not describe the topology; such systems now fall through to the other models. - #13581 (merged 2026-10-19, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Enabled the single-node one-slice optimization for gfx1250, which was previously limited to gfx942 and gfx950. Ring Simple AllGather, AllReduce and ReduceScatter now use one slice per chunk instead of two, halving the per-chunk peer handshakes. The gain is largest below about 128 MB and tapers to a few percent at multi-GB sizes, where the ring is already near the link bandwidth. LL, LL128 and non-ring algorithms are unaffected. - #14123 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * On gfx1250, the DDA LL128 `ncclAllGather`, `ncclAllReduce` and `ncclReduceScatter` receivers keep each line's payload from the cooperative load that saw its flag instead of reading the slice a second time. The DDA scratch bank is now floored to 128 bytes instead of 16, so an LL128 wire line never straddles two 128-byte lines when `RCCL_DDA_FABRIC_BUFFER_SIZE` sets an unaligned size. Results are unchanged. Other architectures are unaffected. - #13415 (merged 2026-10-19, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `fp8_e4m3` and `fp8_e5m2` sum reductions on gfx1250 convert eight elements per instruction using the packed conversion instructions, instead of two. Results are unchanged. Other architectures and reduction operators are unaffected. - #13406 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * The rocSHMEM GDA GIN backend (`NCCL_GIN_TYPE=6`) now sets the RoCE traffic class of its queue pairs from `NCCL_GIN_IB_TC`, then `NCCL_IB_TC`, then the communicator traffic class, matching the IB proxy backend. It previously read `ROCSHMEM_GDA_TRAFFIC_CLASS`, which no longer affects these queue pairs. The default remains traffic class 0 when none of these is set. - #13357 (merged 2026-10-23, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_SYM_TMA_ENABLE` now defaults to `1`, matching NCCL 2.31.2, so on gfx1250 (MI450) the Tensor Data Mover symmetric kernels (`AllGather_TmaST`, `AllReduce_RSxTmaLD_AGxTmaST`, `ReduceScatter_TmaLD`) are offered to the symmetric cost model whenever an application uses symmetric registered memory. The cost model selects them only where their deep loop engages, and launches them at the widest allowed block count. They require 16-byte aligned buffers and a build with gfx1250 in `GPU_TARGETS` and a ROCm SDK that ships the TDM descriptor header. For in-place AllGather and ReduceScatter the per-rank slice must also be a multiple of 16 bytes, so that every rank selects the same kernel. Other architectures are unaffected. Set `NCCL_SYM_TMA_ENABLE=0` to restore the previous behavior. - #13357 (merged 2026-10-23, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * rccl-tests `-A 1` (and `rcclGetCollImplInfo` / `rcclSymKGetInfo`) now report a symmetric collective that selects a Tensor Data Mover kernel as `SymTDM` instead of `SYM`. The symmetric reporting query also accounts for buffer alignment, so it no longer omits the TDM kernels that the live dispatch would pick. - #12291 (merged 2026-10-21, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * The AMD Pensando 'ionic' RDMA provider for the rocSHMEM GDA GIN backend (`NCCL_GIN_TYPE=6`) now properly enforces the required support for creation of 'collapsed' single-entry completion queues. If network driver and firmware support for this feature is unavailable, an initialization failure will occur rather than a silent fallback to plain completion queues. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_DDA_LL128_RS_THREADS` removed (was default `1024`): reduce-scatter LL128 now launches a fixed 512-thread block, one warp per slice, so the thread count is no longer a tuning knob. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed topology GPU indices that followed the HIP device order instead of the PCI order. This reordered rings and NIC selection, causing an `ncclGather` performance regression on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed LL-protocol latency on gfx950 when RCCL is built against a ROCm older than 7.0.2 without the 128-bit global load/store builtins (for example ROCm 7.0.1, or `-DDWORDX4_INTRINSICS=OFF`). Every store to the user output buffer was followed by a release fence and an L2 writeback, which made 4 KB to 64 KB `ncclBroadcast` about 30% slower on MI355X. Builds against ROCm 7.0.2 or later were not affected and are unchanged. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed a 15 to 30% `ncclScatter` bandwidth loss from 16 MB upward on multi-node AINIC jobs, caused by the move to two network channels per peer (for example on 4 nodes of 8 GPUs). Scatter now uses and connects the most channels per peer whose flows still fit the P2P channel pool. Point-to-point operations and the other collectives keep their channel count. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed `ncclCommAbort` and `ncclCommRevoke` hanging when a DDA collective (IPC on gfx942/gfx950, fabric LL/LL128/VMM on gfx1250) was waiting on a peer that never joined. DDA kernels now observe the communicator abort flag while waiting. A parent communicator shrunk with `NCCL_SHRINK_ABORT` no longer uses DDA for later collectives. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * DDA LL128 eligibility functions (`ddaLL128ArOneShotEligible`, `ddaLL128ArTwoShotEligible`) no longer check `RCCL_DDA_LL`. LL and LL128 are now independently gated: `RCCL_DDA_LL=0` disables only the LL tier; `RCCL_DDA_LL128=0` disables only the LL128 tier. Both checks are enforced via `rcclDdaLLThresholdTab` / `rcclDdaLL128ThresholdTab` at the selector level, which return 0 when the respective param is disabled. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * DDA LL128 is disabled for ReduceScatter and AlltoAll on gfx1250: `ddaLL128Max[RS]` and `ddaLL128Max[A2A]` are 0 in the gfx1250 arch table. - #13005 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * With `NCCL_CTA_POLICY=2` (`NCCL_CTA_POLICY_ZERO`), AllReduce on registered symmetric windows now takes the registered CE path whenever it is available, even when the symmetric kernel is eligible, matching AllGather and upstream NCCL. A non-zero arch-table `ceRegMax[AllReduce]` no longer caps it under this policy; a zero entry and an explicit `RCCL_CE_AR_REG_MAX_MSG_BYTES` still apply. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_CE` (default `1`): acts as a prerequisite gate for CE-scratch selection. When `RCCL_FORCE_CE=0`, CE-scratch is disabled entirely for all collectives regardless of other flags. When non-zero, CE-scratch fires if `CTA_POLICY_ZERO` (`-x 2`) is set (bypassing the size window) or if the message size falls within the arch-table `ceNonRegMax` window. AlltoAll CE-scratch now has a size window (`ceNonRegMax[AlltoAll]` in the per-arch table, currently 0 on all arches). - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Registered-window CE AllReduce is now eligible without requiring ``NCCL_CTA_POLICY=2`` > (``CTA_POLICY_ZERO``); it is gated only by the per-arch table window and ``RCCL_CE_ALLREDUCE``. > ``RCCL_FORCE_CE_ALLREDUCE`` now applies exclusively to the non-registered (2-shot/staged) CE > AllReduce path. - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_SYMMETRIC` (default `0`): when set to a non-zero value, biases the `ncclAllReduce` > backend selector to prefer the symmetric kernel path over all other candidates. The selector > still requires the operand buffers to be registered symmetric windows; if the symmetric > path is ineligible the selector falls back normally. - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_DDA` (default `0`): when set to a non-zero value, biases the `ncclAllReduce` > backend selector to prefer DDA (Direct Data Access) paths over all other candidates, > including symmetric and Copy Engine. If no DDA path is eligible for the given message > size and communicator the selector falls back normally. - #14254 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed an intermittent hang, and possible wrong results, in the first DDA fabric LL all-reduces after communicator initialization on gfx1250 (`NCCL_MNNVL_ENABLE=1`). Initialization zeroed the LL scratch buffer and epoch counters with an asynchronous null-stream memset that could land after a rank's first LL kernel, resetting its epoch so that rank used a different flag and scratch bank than its peers. Initialization now waits for the zero fills, and the ranks agree that every rank zeroed its buffers before any rank can launch an LL kernel. If any rank fails to set up its LL buffers or GPU barrier, every rank disables the fabric DDA path together. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed memory corruption, uninitialized reads and a division by zero when parsing malformed or oversized `NCCL_RINGS` values. A ring list with more channels than `MAXCHANNELS`, a ring naming more GPUs or more NET endpoints than supported, or a GPU or NET number too large for an `int` wrote past the channel and parser arrays or wrapped around; NET slots a ring did not name were copied into the graph uninitialized; and input from which no channel could be parsed divided the system bandwidth by zero channels. Such input is now rejected with `ncclInvalidUsage`, and bandwidth is zero when no channel is parsed. Parsing a `NULL` ring string returns `ncclInvalidArgument` instead of dereferencing it. A ring that names its NETs between its GPUs rather than at its ends, for example `"0 N0 1 2 N1 3"`, no longer stores a negative NET id in the graph, which failed multi-node initialization; such a channel is ignored. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed a crash and silent misconfiguration when parsing malformed `RCCL_TREES` values. A tree not terminated by `)`, for example `"(1)(2)(3)(4)(5)(6"`, searched past the end of the string and segfaulted; a GPU number larger than an `int` wrapped around and built a tree over an unintended GPU, as `"(4294967299)"` did over GPU 3; and a stray non-digit character resolved to a GPU number left over from the previous tree, so `"(5)(6)|(0)(x)"` parsed as `"(0)(6)"`. The number of trees, the GPUs per tree, the transcribed tree length and the GPU index used for remapping are now bounded. Invalid input is rejected with `ncclInvalidUsage`, and a `NULL` tree string returns `ncclInvalidArgument`. - #13815 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed wrong results and hangs in application kernels that use the GIN device API, such as the `07_kernel_fusion` GIN and hybrid RMSNorm examples. `nccl_device/coop.h` selected `__syncwarp` from `ROCM_VERSION`, which only RCCL's own build defines, so in application code the thread and warp cooperative groups synchronized the whole block with `__syncthreads()`. The GIN barrier calls them from a subset of threads, which left the block's waves out of step. The headers now fall back to the HIP version when `ROCM_VERSION` is not defined. - #12743 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed the GIN plugin never being finalized for a communicator that had connected a GIN backend, leaking the plugin's per-communicator context and leaving its reference count one too high after every `ncclCommDestroy`. `ncclGinHostFinalize` closed the GIN connections and then zeroed the whole `ncclGinState`, so the `ncclGinFinalize` that follows it in `commFree` saw no active backend and skipped the plugin `finalize()` call. Host finalize now clears only the per-connection state and hands the backend records to `ncclGinFinalize`, which finalizes each one on the last shared-resource reference. `commFree` returns the first of a `closeColl` or plugin-finalize error after those frees; if both fail, the `closeColl` error is the one reported. An `ncclRmaProxyFinalize` error still aborts the rest of teardown immediately. On the Anvil SDMA backend (`NCCL_GIN_TYPE=7`) this was one leaked init context per connected communicator. - #12679 (merged 2026-10-16, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed a crash (`HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION`) on gfx110x and gfx120x when the Tree algorithm runs over shared memory with more channels than the Tree graph search found, for example with `NCCL_ALGO=Tree` and `NCCL_P2P_DISABLE=1`. Channels added after `ncclTopoPreset` had no intra-node tree links; `ncclTopoPostset` now sets the tree and CollNet chain links on every channel, as it already did for rings. - #11133 (merged 2026-10-14, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed the host one-sided RMA API being unavailable on multi-node communicators whose ranks had fused NICs. NIC fusion merges devices for the regular transport only; GIN and RMA select from their own topology node types and never receive a fused device. - #12750 (merged 2026-10-18, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed Anvil SDMA GIN small-message AllToAll latency after `ncclGinFenceLevel` support. A sub-threshold `gin.put`, a `gin.putValue`, or a standalone `gin.signal` emitted two system-scope fences per peer: one in `fenceBeforeSignal` and one in `signalPeer`. `signalPeer` now emits a single system-scope release immediately before the remote atomic, rather than an agent-scope release before it and a system-scope fence after it, which also stops payload stores from passing the signal cell. A clean-queue IPC `SignalInc` skips `fenceBeforeSignal` entirely, since it would only repeat that release. The SDMA `quiet()` before a signal now runs only when the call itself issued SDMA; previously any call carrying a signal or counter quieted the peer queue, including an IPC put after an earlier SDMA put to the same peer, which now relies on `Flush` to drain that queue. The `NCCL_GIN_ANVIL_SDMA_IPC_AGENT_FENCE` debug knob is removed. Each `Flush` thread drains the peers it owns and then clears exactly the bits it drained, over a mask disjoint from every other thread's. Previously `Flush` loaded a snapshot, quieted it, and cleared the whole word on the coop's leading thread, on every rank, which could drop a concurrent mark whose doorbell had not yet fired and then skip the quiet that mark needed. Clearing after the drain also keeps the invariant that a clear bit means a drained queue, so a second coop flushing the same context cannot return from `Flush` while the first is still inside `quiet()`. - #5708 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.30.7 for ROCm 10.1.0` section: > * `hipcc` is no longer accepted as the CMake CXX compiler for RCCL builds. Use `amdclang++` directly (`-DCMAKE_CXX_COMPILER=amdclang++` or set `ROCM_PATH`). `hipcc` was a wrapper around `amdclang++` that is being phased out across the ROCm stack. Only the CHANGELOG is changed. Version macros are left to the component owners. JIRA ID : AIROCDOC-4467 --------- Co-authored-by: Corey Derochie <161367113+corey-derochie-amd@users.noreply.github.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
The gfx1250 arch table (
rcclArchThresholds) was tuned in SPX mode. The DPX sweep disagrees with it in several places: crossovers that are right for SPX are wrong in the other direction for DPX, and the reverse. One table cannot serve both modes, so this PR adds a way to pick the table by compute-partition mode.This PR is plumbing only. The DPX table is an exact copy of the SPX table, so tuning is unchanged. The DPX values come in a follow-up PR.
Technical Details
New env var
RCCL_TUNING_PARTITION_MODE(autodefault |SPX|DPX, case-insensitive)autofollows the compute-partition mode the GPU reports (sysfscurrent_compute_partitionon the physical function). DPX picks the DPX table; SPX, CPX or no reported mode picks SPX.SPX/DPXforce a table.auto.RCCL_IGNORE_ARCH_TABLE=1.Table selection
rcclGetArchThresholds(gcn, partition)takes anrcclTuningPartition, defaulting to SPX, so existing callers are unchanged.rcclArchThresholds_gfx1250_dpx. gfx942 and gfx950 return the same table for both modes.Every rank uses the same table
fillInforecords each rank's request in a newncclPeerInfo::tuningPartition, appended at the tail.initTransportsRankre-pickscomm->archThresholds: DPX only when every rank asks for DPX, otherwise SPX on every rank.gfx1250 tuning table: SPX|DPX).Test Plan
Host microtests (
rccl-UnitTestsMicroInit)autoand empty value, invalid value reported once (isolated test), and a function above the MLOPart bound.comm->archThresholds, on rank 0 and other ranks, for gfx1250 and other archs, and its logging.initTransportsRankcall site, including a comm where only one peer asks for SPX.rcclGetArchThresholdshost fake can now return a separate table per partition, so tests check which table the comm got rather than only which partition was asked for.init.cc.GPU unit test
RcclDdaTierThresholds.TuningPartitionTablePerArchchecks:On H13 nodes
SPX, and the paths chosen are identical to develop.DPX, and the paths chosen are identical to develop (the tables are equal).RCCL_TUNING_PARTITION_MODE=DPXon an SPX node andSPXon a DPX node force the table.Test Result
init.cc(23 mutants)rccl_check_default_behavior.py: no undocumented env-var changes.Submission Checklist
JIRA ID: AICOMRCCL-3001