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

dlamd1dai · 2026-08-18T07:18:56Z

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)

Upstream PR

Tracking PR: ROCm#10930

Review status

All seven @pvallem review threads (Aug 18) addressed and resolved. 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

LargestSysBufferTest, LargestVramBufferTest, and BigSysBufferStressTest allocate buffers sized near total system RAM. On APUs the GPU shares system RAM, so these allocations OOM-kill kfdtest or hang the system. The existing `!hsakmt_is_dgpu()` guard misclassifies APUs as dGPUs on platforms where KFD exposes CPU and GPU as separate topology nodes, so the guard is skipped and the tests run on hardware they shouldn't. Also check HsaNodeProperties::Integrated (set from the kernel AMD_IS_APU flag) as the authoritative APU indicator. The legacy hsakmt_is_dgpu() check is kept as a fallback so behavior on existing dGPU and older APU platforms is unchanged. ROCM-21587 Built and ran kfdtest on an APU target before and after the fix. - Before fix: kfdtest is OOM-killed mid-run. - After fix: 175/175 PASS, no OOM, no regressions. - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Signed-off-by: Priya Hosur <priya.hosur@amd.com>

Port of the AllGather PR #4 review fixes onto the NCCL 2.30.7 test WIP branch (same set as the AG PR branch commit), so validation runs on the SUT directly. - rccl-tests: clamp NCCL_GIN_ANVIL_AG_CTAS to the launched barrier/signal pool (allGatherPoolCtas = max(-V, ladder peak)) to fix the blockIdx.x-indexed OOB. - rccl: fold Anvil-SDMA conn-check infra failures into the collective decision (localFail -> localMissing=nRanks through the shared barrier/allgather) so all ranks abort together instead of hanging peers. Coexists with the existing per-context signalSlot deadlock fix on this branch. - rccl: scope the gfx950 cuMem free-skip to the peer-aperture unmap in ncclCuMemFreeAddr only; ncclCuMemFree and ncclCudaFree free normally. - rccl: printf + __trap() instead of __builtin_unreachable() in the ANVIL_SDMA/ROCSHMEM_GDA get() specializations. - rccl-tests: share allGatherRecvSliceOffset() between AllGatherLsaDirect and the GPU addressing test; scope the test comment. - rccl: rocshmem-gda template test stub adds atomic_add_single for parity with the queue_pair_device.h declarations (atomic_add already present here). - rccl-tests: behavioral "grid never exceeds pool" regression for the CTA clamp. Co-authored-by: Cursor <cursoragent@cursor.com>

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>

The local heap VaMgr backs every kLocal allocation, including the pure VA reservations of the VMM API. Its size was fixed at 2x VRAM, but CLR's VmHeapArray alone reserves 2x VRAM of VA, so hipMemAddressReserve failed with HSA_STATUS_ERROR_OUT_OF_RESOURCES with large vmm reservation. Reserve 8x VRAM instead, falling back to 4x and 2x when the CPU/GPU VA ranges cannot be reserved. Signed-off-by: Flora Cui <flora.cui@amd.com>

__builtin_amdgcn_cvt_scalef32_pk_fp8_f16 and its bf8 counterpart take a tied input as their first parameter, but hadd, hadd_b, hadd2 and hadd2_b pass the f16 source vector v1 there instead of the shortx2_t destination. Older compilers declared the parameter _Vector and tolerated this; newer ones declare it _ExtVector and reject it, so the gfx950 device path fails to compile. Pass u.i16_vec (u1.i16_vec in hadd_b) as the tied input. No functional change on compilers that accepted the old form: the tied operand is the destination the result is written to in either case.

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>

Register gfx1250-strict so A0-targeted code objects (EF_AMDGPU_MACH 0x0eb) load and identify correctly, keeping them separate from the gfx1250 (B0) target the compiler now emits for B0 silicon.

Add Llama-3.1-70B training workload (via MADEngine) to the RCCL CI pipeline on the Ruby cluster. The mechanism is a Docker overlay image that swaps the CI-built librccl.so into a rocm/primus base container, enabling end-to-end validation of RCCL builds against real training workloads without rebuilding the full PyTorch/ROCm stack. Key components: - test_madengine.py: orchestration script — overlay build, manifest generation, SLURM dispatch, live log metric extraction, JSONL perf datastore with per-precision rolling regression detection - therock-rccl-test-madengine.yml: workflow for MADEngine workloads - CI schedule changed from weekly to nightly (06:17 UTC) - CI scripts moved to projects/rccl/ci/scripts/ (from .github/scripts/) Clones pinned: madengine at ec4de0b58c49, MAD at 688828bd9d4a on the mad-rccl branch. HF token passed via process environment only (not written to manifest artifact). Validated single-node (8x MI325X): BF16: 773.9 TFLOP/s/GPU, 1722 tokens/s/GPU (50/50 iterations) FP8: 1249.4 TFLOP/s/GPU, 2780 tokens/s/GPU (50/50 iterations)

Fixes discovered during manual validation on Ruby cluster: - Add --network=none to Docker overlay build (bridge not available on all compute nodes) - Pin SLURM job to the overlay build node when no registry is configured (image only exists locally) - Disable madengine node health check that overrides nodelist - Strip multi-NIC config for single-node runs (Gloo requires all listed NICs to exist, unlike NCCL) - Fix MAD scripts path (primus_megatron-lm, not primus/megatron-lm) - Use full SHA for MAD pin - Smart MAD clone: verify HEAD before fetch+checkout - Handle PermissionError when saving run artifacts to shared dir Validated: 1N/8GPU llama-3.1-70b-training on Ruby BF16: 763.2 TFLOP/s/GPU, 1698.4 tok/s/GPU FP8: 1254.7 TFLOP/s/GPU, 2792.2 tok/s/GPU

Fix critical issues from mkuznet1 and i-kosarev review: Manifest structure (generate_manifest): - Move slurm config into deployment_config.slurm (was root-level, ignored by madengine) - Move env vars into context.docker_env_vars and deployment_config.env_vars (root-level env was dead) - Move container mounts into context.docker_mounts (root-level container_mounts was dead) - Add multiple_results to built_models card so madengine can produce per-precision structured output (root cause of empty metrics) - Set n_gpus to -1 (madengine resolves per-node) - Leave training_precision empty (card runs both BF16 and FP8) - Add GLOO_SOCKET_IFNAME to docker_env_vars - Remove MAD_MULTI_NODE_RUNNER (noise for primus launcher) - Add qos to deployment_config.slurm (was silently dropped) - Add docker_run_options from reference template Simplify run_madengine: - Remove redundant additional_context (slurm, distributed, env_vars) since manifest now has deployment_config in the right place and madengine merges it automatically RCCL commit provenance (get_rccl_commit): - Remove git rev-parse HEAD fallback — in CI this returns TheRock's pinned commit (constant), causing stale image cache hits - Use RCCL_COMMIT_HASH env, then GITHUB_RUN_ID, then sha256 of librccl.so for unique image tags Push tag format (build_rccl_overlay_image): - Replace both / and : in base image name to avoid invalid Docker reference with two colons Results parsing: - Replace parse_perf_csv with parse_perf_results reading perf_entry_super.json (31 fixed columns, per-precision rows) - Fall back to perf.csv with all rows (was collapsing to last row) - Use per-run status in datastore (was stamping one status on all) - Copy perf_entry_super files to run artifacts - Require metric_value for exit code override (was only checking iteration count) Workflow YAML: - Add secrets: inherit to therock-test-madengine job (HF_TOKEN was silently empty) - Add packages: write permission (needed for GHCR push) - Upload perf_entry_super files as CI artifacts

…ne metrics Drive regression checks and datastore writes from structured output (perf_entry_super.json) instead of live-log scraping. Addresses three review items from mkuznet1 (Aug 12): - Filter structured rows by metric_key and key by precision instead of overwriting metric_value on every row (last-row-wins bug) - Invert priority: structured results are primary for regression and datastore, live-log scraping is fallback only when structured data is missing - Build a single precision_results list consumed by all downstream stages (exit-code override, regression, datastore, report) Also fixes: - Preserve header-only runs (precision detected, zero iterations) as incomplete instead of silently discarding them - Record status=fail with metric_value=None when no results exist (was recording pass with no metric)

- Only pin SLURM job to build node for single-node runs; multi-node needs SLURM to allocate across nodes freely - Remove --network=none from docker build so overlay push to GHCR works - Drop vip_prio QoS (DenyQos on meta64 partition rejects it)

…only have fenic0 Gloo fails with "Unable to find address for: enp49s0f0np0" because that interface doesn't exist on Ruby compute nodes. Only fenic0 is present across all nodes.

github-actions · 2026-08-29T05:52:10Z

✅ MINOR ABI BREAKAGE resolved - ABI compliance check is now passing!

…#10000) ## Motivation Fix `reduce_wave` and `reduce_wg` operations that produce incorrect results when running with more than 2 PEs and message sizes ≥512B. Data validation errors indicate partial reductions, suggesting not all PEs are contributing to the reduction. ## Technical Details Needs additional fencing to ensure that all threads have access to the correct data in order to proceed with the reduction. Added appropriate fencing, on IPC and GDA. ## Issue Tracking JIRA ID: AIROCSHMEM-485 ## Test Plan Verified using the rocshmem functional test suite with the team-based reduction test (test ID 33) across multiple PE counts and message sizes: `mpirun -n 4 -mca pml ucx -mca osc ucx -x ROCSHMEM_MAX_NUM_CONTEXTS=1 -x UCX_ROCM_IPC_SIGPOOL_MAX_ELEMS=16384 -x ROCSHMEM_HEAP_SIZE=6442450944 --timeout 300 --map-by numa ./tests/functional_tests/rocshmem_functional_tests -a 33 -w 1 -z 64 -localbuftype heap -s 32768` ## Test Result - Tests with power of 2 PEs (2,4,8) Pass on IPC and GDA. - Tests with non-power of 2 PEs pass on GDA ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

) ## Motivation A couple of memory tests hang when synchronizing after the SDMA engine returns an error from `memory_async_batch_copy` for a batch copy with more than one operation (i.e. `hipMemcpyBatchAsync()` with more than one source). When a batch copy fails, there is one completion signal per operation created via `ActiveSignal()`, but upon failure we only call `ResetCurrentSignal()` once. This leaves N-1 signals stuck since the copy was never submitted and we hang at `hipStreamSynchronize()`. ## Technical Details Reset one completion signal per operation on batch copy failure. ## Issue Tracking AIRUNTIME-2712 ## Test Plan Run Unit_hipMemcpyBatchAsync_Broadcast_MultipleGroups and Unit_hipMemcpyBatchAsync_Broadcast_MixedWithLinear. ## Test Result Both tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

…Cm#10396) ## Summary - Adds MADEngine workload integration into RCCL CI, validated end-to-end on 2N/16GPU (MI355X, Ruby cluster) - Fixes multi-node blockers: SLURM node-pinning, QoS rejection, `NCCL_SOCKET_IFNAME`, Docker `--network=none` - Resolves rdma-core ABI mismatch (container rdma-core 50 vs host rdma-core 61) by bind-mounting the host's libibverbs and IB provider directory at the compiled-in search path - Fixes `perf_entry_super.json` metric name parsing so results report correctly (tok/s/GPU, TFLOPS/GPU) JIRA ID : AICOMRCCL-1332 ## Validation (Ruby cluster, job 26928) - **Llama 3.1 70B BF16 pretrain (2N/16GPU):** 1,719 tok/s/GPU, 773 TFLOPS/GPU - **Llama 3.1 70B FP8 pretrain (2N/16GPU):** 2,848 tok/s/GPU, 1,280 TFLOPS/GPU - Bare-metal rccl-tests 2N all_reduce_perf: 105.3 GB/s avg bus bandwidth, 0 errors ## Key changes - `test_madengine.py`: multi-node manifest generation, host IB lib bind-mounts, metric parsing fixes - `therock-rccl-test-madengine.yml`: new workflow for MADEngine workloads - `rccl_ci_utils.py`: shared CI utilities (GitHub summary, outputs) - `test_jax_collective.py`, `test_pytorch_c10d.py`: additional CI test scripts ## Test plan - [x] 2N MADEngine Llama 3.1 70B training passes on Ruby (BF16 + FP8) - [x] Bare-metal 2N rccl-tests all_reduce_perf passes - [x] `perf_entry_super.json` parsing produces correct metrics - [x] `python3 -m py_compile` passes on all new scripts 🤖 Generated with [Claude Code](https://claude.com/claude-code)

## Summary - Add `rocgdb-corefile` as a test target alongside `rocgdb-cpu` and `rocgdb-gpu` in `therock_matrix.py` for the `debug_tools-dbgapi`, `runtimes`, `all`, and `nightly` groups - Update `valid_test_targets` in the unit test to keep it in sync ## Test plan - [x] Verify CI picks up `rocgdb-corefile` in the affected build groups - [x] Confirm unit test in `therock_configure_ci_test.py` passes 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude <noreply@anthropic.com>

## Motivation  Fix a `rocprof-sys-sample` crash seen when profiling TransferBench through a Python launcher. The direct TransferBench binary path works, but the Python parent/child path can crash while child threads are setting up sampling. ## Technical Details  `sampling::shutdown()` used the same child-process path for two different cases: normal sampled thread exit and post-fork cleanup. In a forked child, an exiting thread could release every sampler slot, including slots owned by other live threads still inside `sampling::configure()`. This PR separates the two cleanup paths: - normal thread-exit shutdown releases only the current thread's sampler slot - `postfork_child()` calls `postfork_child_release_samplers()` to release all inherited sampler slots after fork - `sampling::shutdown()` derives the sampler index from the persistent `thread_info` record, avoiding a thread-local destruction-order dependency during teardown - the child thread-exit path validates and bounds-checks the index before accessing sampler storage - `sampling::shutdown()` keeps its existing signature ## Issue Tracking    JIRA ID - AIPROFSYST-437 ## Test Plan  Validated the reported TransferBench Python parent/child reproduction on an MI300X system using the latest PR commit. Also validated a fork-without-exec workload to cover the `postfork_child()` cleanup path. Repeated both affected paths in a local 2-GPU container. ## Test Result  MI300X TransferBench Python launcher: - 5/5 runs completed successfully - exit code: 0 - crash count: 0 - aggregate GPU bandwidth: 2688.511–2694.895 GB/s MI300X fork-without-exec validation: - 8 threads: 3/3 runs completed, crash=0, child rc=0 - 64 threads: 3/3 runs completed, crash=0, child rc=0 Local 2-GPU validation: - build completed: 29/29 targets - thread-exit path: 3/3 runs completed without crashes - aggregate GPU bandwidth: 184.730–185.610 GB/s - postfork path with 8 and 32 threads completed without crashes - child exit code: 0 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

rocjitsu emulates a GPU at the KFD ioctl boundary, so nothing below that line is ever exercised: there is no configuration space, no BARs, no MMIO, no guest DMA and no interrupts anywhere in the tree, and the real amdgpu and amdkfd kernel drivers therefore cannot run against it at all. Validating what the driver actually sees needs a lower boundary, one that is the silicon ABI rather than a userspace call. Add that boundary. simdojo gains PciDevice, a Component that declares its bus shape and reacts to guest-driven access while staying independent of any VMM, reaching the guest only through an injected IrqSink and DmaEngine. rocjitsu gains the libvfio-user transport that backs it, a scratch device and a protocol-level test client that together exercise the transport without a VMM at all, and the two diagnostics the driver bring-up loop depends on: a burndown of registers the device does not model, and a watchdog for a guest spinning on a status bit that will never change, both symbolized by register name. Serving it is a mode of the existing CLI rather than a second binary, since it is another front end onto the same machine as --daemon. The split between the two halves is in the layout: the device model lives in vm/amdgpu/pci and is portable, shared by every future backend and built everywhere, while vmm/ holds the backends that face a VMM, of which the libvfio-user one is the first. That half is Linux-only, needs uapi headers carrying the VFIO definitions the protocol reuses, and is reached only through ROCJITSU_ENABLE_VFIO, which is off by default. Its dependency is fetched into third_party alongside the existing ones and compiled here, so no meson, ninja or system json-c is required. One constraint is worth knowing before use: a device reaches only the guest memory its client shares mmap-ably, so the VMM must be configured for that, which with QEMU means backing guest RAM with memory-backend-memfd,share=on. Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

…U is An AMD GPU does not tell the driver what it is through its PCI IDs. The driver binds on class, then reads a table out of the device listing every IP block, its version, and where its registers live; without one it cannot decide which block drivers to instantiate, and it refuses the device. Generate that table and write it where the driver looks, 64 KiB below the top of the megabyte-rounded capacity the device reports, since the driver computes that address from what it read rather than being told and never learns the true byte count. Each block names its registers as the list of segments the driver indexes by _BASE_IDX rather than as one address, so publishing a single base leaves most of a block resolving against nothing. Validate all of it here rather than by reading a guest's dmesg: the validator walks the structures the way the driver walks them, from the start of the binary rather than from the table holding them, because an offset measured from the wrong place still lands inside the table and the driver rejects it with a bare -EINVAL. Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

The vfio-user front end so far serves a scratch function: enough to prove the transport, but nothing amdgpu would look at twice. Before a real driver can be pointed at rocjitsu it needs a device whose identity it recognizes, apertures it can map, and answers to the handful of registers it reads before it knows what hardware it is talking to. Add that device, driven entirely by configuration. Identity comes from the section that already describes the GPU to KFD, because the values a guest driver reads and the values KFD reports are the same part and must not be able to disagree; a new section adds only what describes the bus rather than the GPU. The class code carries the most weight there, since the driver's PCI table wildcards the device ID and binds on class alone, then asks the device what it is through IP discovery. Video memory is backed by a descriptor and mapped by the guest directly, since an aperture that trapped every access would be unusable, while the register aperture always traps because a register read has to be answered by the model rather than served from a page. A window smaller than memory is the normal case for a real part, so the index and data registers reach whatever the aperture does not, which is where the driver looks for the discovery table. The registers report a memory size that is neither zero nor all-ones, either of which makes the driver abandon discovery before it starts, firmware initialization already complete since an emulated device has nothing to wait for, and scratch registers left at zero so the driver looks for that table at the top of memory. Configurations that no bus could realize are refused up front rather than part-way through building a transport: apertures must be powers of two above the PCI floor and large enough to reach the registers the device claims to answer, and a capacity must survive being stated in the megabytes the driver reads. A device that shares memory by descriptor serves one client, because the mapping outlives the connection and cannot be reclaimed. Anything unmodelled reads as absent hardware and is counted, which is the list the next stage works through. Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

…U is (ROCm#10732) An AMD GPU does not tell the driver what it is through its PCI IDs. The driver binds on class, then reads a table out of the device listing every IP block, its version, and where its registers live; without one it cannot decide which block drivers to instantiate, and it refuses the device. Generate that table and write it where the driver looks, 64 KiB below the top of the megabyte-rounded capacity the device reports, since the driver computes that address from what it read rather than being told and never learns the true byte count. Each block names its registers as the list of segments the driver indexes by _BASE_IDX rather than as one address, so publishing a single base leaves most of a block resolving against nothing. Validate all of it here rather than by reading a guest's dmesg: the validator walks the structures the way the driver walks them, from the start of the binary rather than from the table holding them, because an offset measured from the wrong place still lands inside the table and the driver rejects it with a bare -EINVAL.

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>

github-actions · 2026-08-29T20:29:07Z

✅ MINOR ABI BREAKAGE resolved - ABI compliance check is now passing!

…e and signals Fill the initCommConfig slot in ncclTestEngine (nullptr, matching alltoall.cu) so getDevCommRequirements is not shifted when ENABLE_DEVICE_API is on. Use per-CTA signalIndex (blockIdx.x) in the GIN/SDMA AllGather body so the signal baseline, barrier pool, and waitSignal target stay aligned across CTAs. Co-authored-by: Cursor <cursoragent@cursor.com>

Centralize AllGather env parsing and devtime helpers in the shared policy header, fix LSA rank iteration and CMake ENABLE_DEVICE_API gating, and strengthen unit/pytest coverage for the default 32 KiB tier crossover. Co-authored-by: Cursor <cursoragent@cursor.com>

Match #wrong=N with N>0 instead of the table header so connectivity-gate retries are not misclassified as genuine data failures. Co-authored-by: Cursor <cursoragent@cursor.com>

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

dlamd1dai changed the title ~~[AICOMRCCL-1798] rccl: GIN device kernels — AllGather (-D3 hybrid Anvil-SDMA)~~ [AICOMRCCL-1798] feat(rccl): GIN SDMA AllGather Implementation (-D3 hybrid Anvil-SDMA) Aug 18, 2026

dlamd1dai changed the title ~~[AICOMRCCL-1798] feat(rccl): GIN SDMA AllGather Implementation (-D3 hybrid Anvil-SDMA)~~ [AICOMRCCL-1798] feat(rccl): GIN SDMA AllGather Implementation (-D3 hybrid xGMI-SDMA) Aug 18, 2026