feat(rccl/gin-sdma): GIN SDMA AllGather Implementation (-D3 hybrid xGMI-SDMA) by dlamd1dai · Pull Request #10930 · ROCm/rocm-systems

dlamd1dai · 2026-08-28T23:57:42Z

Summary

Implements the GIN device-kernel AllGather collective (JIRA AICOMRCCL-1798): each rank issues gin.put() of its chunk to all peers.

Size-adaptive hybrid (-D 3):

LSA-direct tier (≤ 32 KiB/rank): all-peers xGMI store, 16 CTAs
GIN-put / Anvil-SDMA tier (> 32 KiB/rank): copy-engine puts with signals, 4 CTAs

Device-time via shared rccl-tests CLI (--device_timing, -L/-P), matching the A2A path. Stacks on merged upstream #10658 (A2A devtime), #10672 (GIN unit tests), and #10675 (HIP scope fence guard).

JIRA ID : AICOMRCCL-1798

Stack

Base: develop (64135e20f4)
Tip: 12975a5842 — 12 AllGather commits, 13 files (+1379/−17)

Review status

All seven @pvallem review threads (Aug 18) addressed and resolved on fork PR dlamd1dai/rocm-systems#4. Re-review pending.

Test plan (MI355X / gfx950, NP=8)

Host policy unit tests: 19/19 pass (gin_sdma_allgather_policy_test.cpp)
On-GPU tier/addressing tests: 4/4 pass (gin_sdma_allgather_gpu_test.cpp)
Functional 128 B–4 GiB: #wrong = 0 in LSA and SDMA tiers (all_gather_perf -D 3, NCCL_GIN_TYPE=5)
Device-time: adaptive CTA self-select (16 LSA / 4 SDMA), tier boundary at 32 KiB/rank
Multi-segment hang guard: opt-in RCCL_TESTS_GIN_SDMA_AG (256 MiB–2 GiB/rank)
Docker validation (ROCm 7.13, 8× MI355X): ~390–427 GB/s bus BW, BUILD SANITY OK

therock-pr-bot · 2026-08-28T23:58:46Z

✅ 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?

dlamd1dai · 2026-08-29T00:11:52Z

MI355X validation (Docker, 8× gfx950)

GIN-SDMA hybrid AllGather (all_gather_perf -D 3, NCCL_GIN_TYPE=5) validated on 8× MI355X using a ROCm 7.13 Docker image built from the NCCL 2.30.7 validation line (equivalent AG kernel/design to this PR branch).

Configuration

Parameter	Value
NP / GPUs	8 / 1 GPU per rank
Kernel	`GinHybridAllGatherKernel` (`-D 3`)
GIN backend	Anvil-SDMA (`NCCL_GIN_TYPE=5`)
LSA↔SDMA threshold	32 KiB/rank (`NCCL_GIN_ANVIL_SDMA_THRESHOLD_ALLGATHER=32768`)
Dtype / validation	`int8`, `-A 1` (datacheck on)
cuMem	`NCCL_CUMEM_ENABLE=1`
Warmup / iters	5 / 20

Build gate

BUILD SANITY OK — GIN host plugins present in librccl.so; Anvil SDMA factory linked into all_gather_perf.

Functional results (`#wrong = 0` throughout)

128 B – 128 MiB (powers-of-two sweep):

Total size	Per-rank chunk	Tier	Bus BW (out-of-place)
128 B – 32 KiB	≤ 32 KiB/rank	LSA-direct	0.01 – 2.7 GB/s
64 KiB – 8 MiB	> 32 KiB/rank	GIN/SDMA	5.1 – 170 GB/s
16 MiB – 128 MiB	SDMA (multi-seg)	GIN/SDMA	243 – 391 GB/s

128 B – 4 GiB (extended sweep):

Total size	Bus BW (out-of-place)	#wrong
128 MiB	391 GB/s	0
256 MiB – 1 GiB	407 – 421 GB/s	0
2 GiB	424 GB/s	0
4 GiB	427 GB/s	0

Summary

0 wrong elements across the full 128 B – 4 GiB range.
LSA-direct tier behaves correctly at and below the 32 KiB/rank crossover.
SDMA tier reaches ~390 GB/s bus BW at 128 MiB and ~427 GB/s at 4 GiB with no hang.
Multi-segment SDMA path (128 MiB segments) exercised cleanly through 4 GiB total.

Host policy unit tests (19/19) and on-GPU addressing tests (4/4) pass on the same stack; details in the PR test plan.

wenkaidu · 2026-08-29T00:52:28Z

+}
+
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
+testResult_t AllGatherGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs, ncclCommProperties_t* commProperties) {


This doesn't match common.h:
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
testResult_t (getDevCommRequirements)(int deviceImpl, ncclDevCommRequirements reqs, ncclComm_t comm);

Good catch — thanks for pointing this out. You're right that the callback didn't match common.h.

Fixed in cbdc463: AllGatherGetDevCommRequirements now takes ncclComm_t comm as the third argument and calls ncclCommQueryProperties() internally, following the same pattern as AlltoAllGetDevCommRequirements and AllReduceGetDevCommRequirements. The GIN-not-enabled path now returns testInvalidUsage for consistency with alltoall.

Please let me know if anything else looks off.

wenkaidu · 2026-08-29T00:53:50Z

  static std::once_flag once;
  std::call_once(once, []() { atexit(rcclShutdownHandler); });
 }
+}


Thanks — that was a stray brace introduced when adding rcclSkipCuMemFree(). Removed in cbdc463; the shutdown helpers and the gfx950 workaround are back at file scope as intended.

This was also what was breaking the RCCL host unit test build (extraneous closing brace at line 59). Appreciate the quick spot.

… wip Cherry-pick of cbdc463 from users/dondai/gin-stage3b-sdma-ag. On this branch the signature/alloc.h fixes were already present (c18dd4c); this commit only switches the GIN-not-enabled path to testInvalidUsage, matching alltoall and the upstream PR review. Co-authored-by: Cursor <cursoragent@cursor.com>

Port the GIN Anvil-SDMA hybrid AllGather kernel into stage3b, at parity with the AllReduce path: per-rank chunk size gates LSA-direct peer stores (<= sdmaThreshold) vs. all-peers GIN copy-engine puts (> threshold). - all_gather.cu: GinHybridAllGatherKernel (NCCL_GIN_TYPE=5) + host launcher; sizing/threshold/tier/bandwidth logic factored into a testable policy header. - gin_sdma_allgather_policy.h: pure host/device policy helpers (single source of truth for the LSA<->SDMA crossover) shared by kernel and unit tests. - rccl/test: host policy unit test + on-GPU tier-predicate/addressing test. - gin_plugin_anvil_sdma.cc: honor explicit NCCL_GIN_ANVIL_SDMA_THRESHOLD=0. - common.cu: don't report host cputime column for device-impl kernels. Validated on 8x MI355X across 128 B-128 MB: 0 wrong in both tiers, crossover lands exactly at 2 MiB/rank, SDMA tier ~390 GB/s bus bw. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit ffe0d60)

Mirror the AllToAll GIN-SDMA >1 GiB hang-fix tests (PR ROCm#9927) for AllGather so the multi-channel/SDMA ordering concern is covered by an automated test, not just manual runs. Drives GinHybridAllGatherKernel (-D 3, NCCL_GIN_TYPE=5) with the SDMA tier forced (threshold=0) at per-rank chunks that cross the <=128 MiB segmentation boundary (256 MiB -> 2 seg, 2 GiB -> 16 seg), plus a 2 GiB-total completion guard with a subprocess timeout that turns a reintroduced SDMA hang into a failure. Exact-integer dtypes + built-in data check (-c 1) so a truncated /stale tail can't hide under float tolerance. Opt-in via RCCL_TESTS_GIN_SDMA_AG (needs an 8x MI355X-class node + MPI + a GIN-SDMA build); config mirrors test_AllToAll.py (NP/launcher/xenv/exe/timeout). Includes the harness's intermittent gfx950 cuMem-VMM connectivity-gate retry so the transient abort doesn't flake the test; a genuine data-check mismatch is never retried. Verified on 8x MI355X: 2 GiB-total hang guard passes, 0 wrong, ~425 GB/s (SDMA/SIMPLE tier, 2x128 MiB segments). Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 38b741a)

Give the GIN Anvil-SDMA hybrid AllGather its own LSA<->SDMA crossover knob, NCCL_GIN_ANVIL_SDMA_THRESHOLD_ALLGATHER, decoupled from the backend gin.put inline-vs-copy-engine threshold (the global NCCL_GIN_ANVIL_SDMA_THRESHOLD / ctx->sdmaThreshold). The device API is untouched: the threshold is host-resolved and passed to GinHybridAllGatherKernel as a kernel argument, consumed only by the host-testable tier predicate (chunkUsesLsaTier); gin.put's signature is unchanged. - gin_sdma_allgather_policy.h: replace the context-based resolveSdmaThreshold() with a pure, precedence-based pickSdmaThreshold() (per-collective env > global env > compiled default) plus the 32 KiB/rank compiled default constant. - all_gather.cu: host AllGatherResolveSdmaThreshold() reads the per-collective env, falls back to the global env then the compiled default, parsed as 64-bit (explicit 0 honored -> all-SDMA tier; >2 GiB values do not wrap). Thread the resolved value through a new AllGather-specific launch helper into the kernel; guard the RunColl -D 3 case with ENABLE_DEVICE_API so a device-API-off build cleanly falls through to testNotImplemented (matches the shared launcher's prior behavior). - gin_sdma_allgather_policy_test.cpp: cover pickSdmaThreshold precedence, explicit-0, and the 64-bit no-wrap path. Validated: 13/13 host policy unit tests; docker device compile + A2A and AllGather GIN smoke gates pass on 8x MI355X (0 wrong). Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 212d3fc)

AllGather launched a fixed deviceCtaCount (-V) grid for all sizes, which over- subscribes the GIN-put/SDMA tier: only nRanks threads issue the puts (the copy engines move the bytes), so extra CTAs are pure barrier/signal overhead. Mirror the ReduceScatter design and self-select a size-adaptive CTA count decoupled from -V, keyed off the same LSA<->SDMA tier predicate the kernel evaluates: * LSA-direct tier (chunk <= threshold): 16 CTAs (grid-stride store scales with threads; peaks ~16 on 8x MI355X, tiny sizes are latency-bound/CTA-indifferent). * GIN-put/SDMA tier (chunk > threshold): 4 CTAs (stable near-peak 1 MiB-128 MiB). Measured gains vs the old fixed -V 32 launch: 512 KiB +21%, 2 MiB +18%, 8 MiB +13%, 32 MiB +5%; LSA tier unchanged. Datacheck bit-exact 128 B-256 MiB; A2A/RS unaffected. Add allGatherCtas()/allGatherMaxCtas() to the policy header (host- unit-testable), NCCL_GIN_ANVIL_AG_CTAS to pin a fixed count (diagnostic), an explicit-grid launcher, and size the devComm barrier/signal pool to max(deviceCtaCount, allGatherMaxCtas()) so the self-selected grid never exceeds it. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit e5ce620)

…tage3e Bring the in-kernel wall_clock64 device-timing hook (AllGatherDeviceTime, opt-in via NCCL_GIN_ANVIL_DEVICE_TIMING) to the AG branch, reconciled with the size- adaptive CTA selection just cherry-picked here. To share one implementation between the production and timed launches, factor the LSA/SDMA logic out of GinHybridAllGatherKernel into ginAllGatherBody<T>() and add GinHybridAllGather- TimedKernel that loops the body under one persistent launch. The devtime path self-selects the same allGatherCtas() count as the perf path (verified: the resolved per-collective sdmaThreshold so the timed tier matches the launched cfg. common.h already carries the testColl.deviceTime field and gin_sdma_devtime.h exists on this branch (used by AllToAll), so only all_gather.cu changes: include the scaffold, add the body/timed kernel, add AllGatherDeviceTime, and wire it as the testColl.deviceTime member. Builds clean; datacheck bit-exact; devtime prints. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 4435b34)

Reword the size-adaptive CTA comments so the AllGather design does not reference the ReduceScatter design (this is the AG branch). No functional change. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 3ae778c)

Cover the new allGatherCtas()/allGatherMaxCtas() policy helpers: LSA vs SDMA tier selection keyed off the crossover predicate, tracking of a moved threshold, env pin + 128 clamp + 0-falls-back-to-ladder, and max-covers-both-tiers. 18/18 host policy tests pass. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit e35adb7)

Resolves the seven review comments on #4 (GIN SDMA AllGather): - rccl-tests: clamp NCCL_GIN_ANVIL_AG_CTAS to the launched barrier/signal pool (allGatherPoolCtas = max(-V, ladder peak)); a pin above the pool previously launched more CTAs than slots -> OOB on the blockIdx.x-indexed pools. - rccl: fold Anvil-SDMA conn-check infra failures (write/verify kernel, D2H copy, slot-reset memset) into the collective decision instead of a per-rank goto, so a single-rank failure aborts all ranks together rather than leaving peers hung on bootstrapBarrier/allgather. - rccl: scope the gfx950 cuMem free-skip to the peer-aperture unmap in ncclCuMemFreeAddr only; ncclCuMemFree (physical handle) and ncclCudaFree (ordinary buffers) now free normally, so long-lived procs no longer leak on every ncclCommDestroy. Leak is confined to the peer VA aperture. - rccl: replace __builtin_unreachable() in the ANVIL_SDMA/ROCSHMEM_GDA get() specializations with __builtin_trap() so a future get() fails loudly. Uses the compiler builtin (no HIP runtime include) so it also compiles in the symmetric gensrc TU that pulls these headers via nccl_device.h. - rccl-tests: share allGatherRecvSliceOffset() between AllGatherLsaDirect and the GPU addressing test so a wrong slice formula is caught, not just self-consistent; scope the test comment to what it does/doesn't cover. - rccl: fix the rocshmem-gda template test stub to define atomic_add / atomic_add_single (was stale atomic_nofetch, broke --rocshmem-gin builds). - rccl-tests: replace tautological policy asserts with a behavioral "grid never exceeds pool" regression guarding the CTA-pin clamp. Validated on 8x MI355X (gfx950): --rocshmem-gin image builds clean; host policy unit tests pass (AllGather 19/19 incl. the new clamp+grid guards, SDMA 5/5); all_gather_perf -D3 (NCCL_GIN_TYPE=5) 128 B-128 MiB is bit-exact (#wrong=0) in both the LSA-direct and SDMA tiers with no hang. Co-authored-by: Cursor <cursoragent@cursor.com>

Replace stale NCCL_GIN_ANVIL_AG_DEVTIME_* env vars with the harness --device_timing/-L/-P tier knobs (matching A2A), gate to -D 3, and return the -1.0 mode-2 sentinel when no measurement is produced. Co-authored-by: Cursor <cursoragent@cursor.com>

…sync check Add LSA signal connectivity self-test (localFail collective abort) and the gin_anvil_sdma_oss7_device launchers so the PR branch matches the wip branch's plugin safety path. Include the AllGather manifest sync script used to keep the six core AG files identical across gin-stage3b-sdma-ag and the NCCL 2.30.7 wip. Co-authored-by: Cursor <cursoragent@cursor.com>

Remove the AllGather manifest and sync-check scripts; that validation tooling belongs on the NCCL 2.30.7 wip branch only. The conn-check plugin fix from the prior commit is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

…ates Reconcile PR #4 with merged upstream ROCm#10658 (A2A devtime host helpers, BenchTime skip/augment/check path) and ROCm#10672 (GDA/SDMA device template tests + GIN CI wiring). AllGather devtime now buffers via devtimeAugmentLine to match the A2A augment pattern; AllGather-specific unit test targets remain. Co-authored-by: Cursor <cursoragent@cursor.com>

Remove a stray closing brace in alloc.h that broke RCCL host builds, and align AllGatherGetDevCommRequirements with common.h by taking ncclComm_t and querying comm properties internally (matching alltoall/all_reduce). Co-authored-by: Cursor <cursoragent@cursor.com>

dlamd1dai · 2026-08-29T19:24:16Z

Hygiene: rebased `users/dondai/gin-stage3b-sdma-ag` onto current `develop` (`ef24cf0`) — 13 commits replayed cleanly, no conflicts. New tip: `b9b0cb4` (same review fixes as former `cbdc463`). CI re-triggered via force-push.

@wenkaidu — could you take another look when you have a moment? Both review threads (`alloc.h` brace, `AllGatherGetDevCommRequirements` signature) are addressed in the rebased tip.

amd-argus

Automated Review Guard for Upfront Scrutiny

Findings

55 finding(s) are posted inline, on the lines they refer to.

projects/rccl-tests/src/all_gather.cu:1 (Low) the PR description links an internal work item, which both PR templates forbid

Bottom line: Request changes. Two build breaks land before any runtime behaviour is reachable: the testEngine initializer at all_gather.cu:424 and the unlinked conn-check symbols in gin_plugin_anvil_sdma.cc.

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

amd-argus · 2026-08-29T23:56:10Z

 NCCL_WEAK struct testEngine ncclTestEngine = {
  /* .getBuffSize = */ AllGatherGetBuffSize,
-  /* .runTest = */ AllGatherRunTest
+  /* .runTest = */ AllGatherRunTest,


common.h:174 declares initCommConfig as the third testEngine member under NCCL_VERSION_CODE >= 2.14.0, so this positional initializer binds AllGatherGetDevCommRequirements to void()(ncclConfig_t). alltoall.cu:624 has the slot, and its guard also carries the >= 2.29.0 arm that common.h:177 uses.

amd-argus · 2026-08-29T23:56:10Z

 static std::mutex pluginMutex;

+// [GIN-CONN-CHECK] Device connectivity self-test launchers (gin_anvil_sdma_oss7_device.cc).
+extern "C" int ginAnvilConnWrite(void* remoteAddrsDev, int nRanks, int selfRank, unsigned long long stamp);


ginAnvilConnWrite and ginAnvilConnCheck are defined only in gin_anvil_sdma_oss7_device.cc, which src/CMakeLists.txt:523 adds to librccl alone. rccl-UnitTestsGinAnvilPlugin (test/CMakeLists.txt:412) compiles this file with gin_anvil_plugin_test_stubs.cc and intentionally does not link rccl (test/CMakeLists.txt:875), so it fails to link once ENABLE_ROCSHMEM_GIN is on.

amd-argus · 2026-08-29T23:56:10Z

+    // 1) Store my stamp into every peer's test slot [rank].
+    if (rank == injRank) {
+      WARN("GIN anvil-sdma: [TEST] injecting connectivity fault on rank %d (skipping signal writes)", rank);
+    } else if (ginAnvilConnWrite(ctx->signal_remote_addrs_dev, nRanks, rank, stamp) != 0) {


With nSignals=1 and nRanks=8 the conn-check writes 8x the space it was given:

bind sizes the slot 1 * 8 B = 8 B (line 648) ginAnvilConnWriteKernel writes base[selfRank] on every peer 8 * 8 B = 64 B hipMemset(lsaSelf, 0, 8 * nRanks) = 64 B

ginSignalCount=1 ships today (alltoall.cu:140), and this gate runs on the first signal bind of any comm.

amd-argus · 2026-08-29T23:56:10Z

+  const int ginContext = 0;
+  const unsigned int signalIndex = 0;
+  ncclGin gin { devComm, ginContext };
+  const uint64_t signalValue = gin.readSignal(signalIndex);


signalIndex is 0 for every CTA, but the barrier at line 209 is indexed by blockIdx.x, so CTAs sync independently. A peer's CTA0 can clear barrier 0 and start putting before our CTA3 reaches this snapshot, so CTA3 counts those puts into its own baseline and then waits for more than is ever sent:

peer CTA0 clears barrier 0, its puts land our signal[0] += m our CTA3 reaches line 207 snapshots base + m our CTA3 waits at line 222 for base + m + nRanks but only nRanks increments are ever sent base + nRanks -> CTA3 hangs

alltoall.cu:269 uses signalIndex = blockIdx.x for exactly this reason. Confining the put and wait to blockIdx.x == 0 would also work, since the stride at line 215 means only CTA0 puts.

amd-argus · 2026-08-29T23:56:10Z

  // Check if process is shutting down to avoid use-after-free in HIP runtime
-  if (rcclShutdownFlag().load(std::memory_order_acquire)) {
-    INFO(NCCL_ALLOC, "ncclCuMemFreeAddr: Skipping free (process shutdown) pointer %p", ptr);
+  if (rcclShutdownFlag().load(std::memory_order_acquire) || rcclSkipCuMemFree()) {


p2p.cc:429 maps the peer handle then immediately cuMemRelease()s it, so the mapping holds the last reference and skipping cuMemUnmap pins the physical allocation, not just the VA. It also skips ncclMemUntrack. p2pDeregisterMemHandle (p2p.cc:865) routes here too, so a register/deregister loop on gfx950 grows without bound, which is the opposite of what the comment promises.

amd-argus · 2026-08-29T23:56:12Z

+// (force off) or =1 (force on).
+inline bool rcclSkipCuMemFree() {
+  static const bool skip = [](){
+    const char* e = getenv("NCCL_CUMEM_SKIP_FREE");


docs/api-reference/env-variables.rst:55 documents NCCL_CUMEM_ENABLE and NCCL_CUMEM_SKIP_FREE appears nowhere in docs/. Could a row go in with its gfx950 auto-enable default?

amd-argus · 2026-08-29T23:56:12Z

@@ -0,0 +1,154 @@
+/*************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.


Sibling forms differ by directory: rccl-tests/src uses "2024, Advanced Micro Devices" with the comma (gin_sdma_devtime.h:2), rccl/test/gin uses "2026 Advanced Micro Devices" without it (gin_sdma_policy_test.cpp:2). So this file wants "2026, " and both new gin tests want "2026 ". The fourth new file, gin_anvil_sdma_oss7_device.cc, is already correct.

amd-argus · 2026-08-29T23:56:12Z

+
+#if defined(__HIPCC__) || defined(__CUDACC__)
+
+#include <hip/hip_runtime.h>


Suggested change

#include <hip/hip_runtime.h>

#include <cstdint>

#include <hip/hip_runtime.h>

amd-argus · 2026-08-29T23:56:12Z

+// Requires a visible GPU at run time; skips cleanly (GTEST_SKIP) otherwise, so it
+// carries a "gpu" label alongside "unit".
+
+#include <gtest/gtest.h>


Suggested change

#include <gtest/gtest.h>

#include <cstdint>

#include <vector>

#include <gtest/gtest.h>

#include <hip/hip_runtime.h>

#include "gin_sdma_allgather_policy.h"

amd-argus · 2026-08-29T23:56:12Z

+// message sizing, LSA<->SDMA tier crossover, context-threshold fallback and
+// bandwidth math that all_gather.cu (GinHybridAllGatherKernel) relies on.
+
+#include <gtest/gtest.h>


Suggested change

#include <gtest/gtest.h>

#include <cstddef>

#include <cstdint>

#include <gtest/gtest.h>

#define GIN_SDMA_HOST_ONLY 1

#include "gin_sdma_allgather_policy.h"

amd-argus · 2026-08-30T00:16:11Z

+  for (size_t i = tid; i < count; i += nthreads) {
+    T value = src[i];
+    for (int lp = 0; lp < nRanks; lp++) {
+      T* dst = (T*)ncclGetLsaPointer(recvwin, recvoffset, lp) + dstOff;


lp runs to nRanks, but ncclGetLsaPointer indexes the LSA team:

ncclGetLsaPointer(w, off, peer) = lsaFlatBase + peer * stride4G (core__funcs.h:114) ^ bounded by lsaSize, not world nRanks

On a multi-node comm nRanks > lsaSize, so lp >= lsaSize addresses outside the window. alltoall.cu:355 iterates lsa.nRanks with startLsa, and alltoall.cu:71 rejects the mismatch in its requirements hook; AllGatherGetDevCommRequirements has no equivalent check.

amd-argus · 2026-08-30T00:17:18Z

+  }
+}
+
+TEST(AllGatherGpu, TierPredicateMatchesHost) {


TierPredicateMatchesHost compares the kernel output against a host golden at line 99, but both call chunkUsesLsaTier from the same header in the same TU, so the assertion holds for any definition including an inverted one. It proves device and host codegen parity, not the semantics. One absolute host assertion at the default threshold would close it.

amd-argus · 2026-08-30T00:17:20Z

+// returning the "unset" sentinel when absent/empty/unparseable so allGatherCtas()
+// falls back to its size-adaptive ladder.
+static inline size_t AllGatherParseCtasEnv() {
+  const char* e = getenv("NCCL_GIN_ANVIL_AG_CTAS");


AllGatherParseCtasEnv and the parseEnv lambda at line 151 are the same getenv plus strtoull, about 12 lines, but only the lambda reports whether the var was set. So NCCL_GIN_ANVIL_AG_CTAS=0 and unset are indistinguishable once allGatherCtas treats 0 as unset, while the threshold path deliberately honours 0. One shared parser removes the asymmetry.

amd-argus · 2026-08-30T00:25:48Z

+        // tier, few (4) for the GIN-put/SDMA tier where only nRanks threads issue
+        // the puts and extra CTAs are pure barrier overhead. NCCL_GIN_ANVIL_AG_CTAS
+        // pins a fixed count (diagnostic).
+        const size_t agSdmaThreshold = AllGatherResolveSdmaThreshold();


AllGatherResolveSdmaThreshold() runs two getenv calls on every AllGatherRunColl, while line 286 caches AllGatherParseCtasEnv() in a function-local static. Could this one be cached the same way?

amd-argus · 2026-08-30T00:31:37Z

+  if (count == 0 || devtimeLoop < 1) return testSuccess;
+
+  const size_t chunkBytes = count * (size_t)wordSize(type);
+  int loop = devtimeLoop;


Lines 324 to 334 are line for line alltoall.cu:516 to 526, including both the 64 MiB and 8 MiB literals and both skip clamps; the only difference is chunkBytes here against perPeerBytes there. Hoisting it into gin_sdma_devtime.h removes about 11 lines and one copy of the tuning constants.

amd-argus · 2026-08-30T00:47:20Z

+  auto kernel = SPECIALIZE_KERNEL(GinHybridAllGatherTimedKernel, type, op);
+  if (kernel == nullptr) return testSuccess;
+
+  const size_t sdmaThreshold = AllGatherResolveSdmaThreshold();


Lines 339 to 345 are the same three statements as lines 284 to 291 in the perf path, including a second independent function-local static for the CTA env. A small AllGatherResolveLaunchCtas helper removes about 8 lines and one of the statics.

amd-argus · 2026-08-30T00:54:09Z

+    int nRanksGlobal = args->nProcs * args->nThreads * args->nGpus;
+    const size_t totalBytes = count * wordSize(type) * (size_t)nRanksGlobal;
+    double sec = devUs * 1.0e-6;
+    double algBw = (double)totalBytes / 1.0e9 / sec;


algBw and busBw are recomputed by hand here, but AllGatherGetBw at line 68 already delegates to gin_sdma_allgather::bandwidthGBps, which computes the identical values (gin_sdma_allgather_policy.h:144). Reusing the helper drops 3 lines and the drift risk.

amd-argus · 2026-08-30T00:55:07Z


 void AllGatherGetBw(size_t count, size_t typesize, double sec, double* algBw, double* busBw, int nranks) {
-  double baseBw = (double)(count * typesize * nranks) / 1.0E9 / sec;
+  gin_sdma_allgather::bandwidthGBps(count, typesize, sec, nranks, algBw, busBw);


AllGatherGetBw takes size_t typesize but bandwidthGBps at gin_sdma_allgather_policy.h:142 takes int typeSize, so this call narrows. Harmless for 1 to 16, but the inline code it replaced had no narrowing; making the helper take size_t keeps it exact.

amd-argus · 2026-08-30T00:55:10Z

+  EXPECT_EQ(pickSdmaThreshold(false, 0, true, 65536, kAllGatherSdmaThresholdDefault), 65536u);
+}
+
+TEST(AllGatherPolicyThreshold, CompiledDefaultWhenNothingSet) {


CompiledDefaultWhenNothingSet passes kAllGatherSdmaThresholdDefault in and expects it back, so it cannot detect a change to the constant's value. With chunkUsesLsaTier exercised only at 0, 128 and 2097152, and every CTA case passing a bare 32768, changing the default fails no test. One absolute pair at the constant would pin it.

dlamd1dai requested review from a team August 28, 2026 23:57

dlamd1dai mentioned this pull request Aug 28, 2026

feat(rccl/gin-sdma): GIN SDMA AllGather Implementation (-D3 hybrid xGMI-SDMA) dlamd1dai/rocm-systems#4

Open

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

dlamd1dai self-assigned this Aug 29, 2026

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

wenkaidu reviewed Aug 29, 2026

View reviewed changes

dlamd1dai and others added 13 commits August 29, 2026 15:21

dlamd1dai force-pushed the users/dondai/gin-stage3b-sdma-ag branch from cbdc463 to b9b0cb4 Compare August 29, 2026 19:23

dlamd1dai requested a review from wenkaidu August 29, 2026 19:29

amd-argus Bot reviewed Aug 29, 2026

View reviewed changes

amd-argus Bot reviewed Aug 30, 2026

View reviewed changes

		@@ -0,0 +1,154 @@
		/*************************************************************************
		* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.


		#if defined(__HIPCC__) \|\| defined(__CUDACC__)

		#include <hip/hip_runtime.h>

-#include <gtest/gtest.h>
+#include <cstddef>
+#include <cstdint>
+#include <gtest/gtest.h>
+#define GIN_SDMA_HOST_ONLY 1
+#include "gin_sdma_allgather_policy.h"

Uh oh!

Conversation

dlamd1dai commented Aug 28, 2026 • edited by atlassian Bot Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Stack

Review status

Test plan (MI355X / gfx950, NP=8)

Uh oh!

therock-pr-bot Bot commented Aug 28, 2026

✅ All Policy Checks Passed

Uh oh!

dlamd1dai commented Aug 29, 2026

MI355X validation (Docker, 8× gfx950)

Configuration

Build gate

Functional results (#wrong = 0 throughout)

Summary

Uh oh!

wenkaidu Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

dlamd1dai Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

wenkaidu Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

dlamd1dai Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

dlamd1dai commented Aug 29, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

amd-argus Bot Aug 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 29, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

dlamd1dai commented Aug 28, 2026 •

edited by atlassian Bot

Loading

Functional results (`#wrong = 0` throughout)

amd-argus Bot Aug 29, 2026 •

edited

Loading

amd-argus Bot Aug 29, 2026 •

edited

Loading

amd-argus Bot Aug 29, 2026 •

edited

Loading

amd-argus Bot Aug 29, 2026 •

edited

Loading

amd-argus Bot Aug 29, 2026 •

edited

Loading

amd-argus Bot Aug 30, 2026 •

edited

Loading

amd-argus Bot Aug 30, 2026 •

edited

Loading

amd-argus Bot Aug 30, 2026 •

edited

Loading

amd-argus Bot Aug 30, 2026 •

edited

Loading

amd-argus Bot Aug 30, 2026 •

edited

Loading