[dlamd1dai]

Summary

Fixes a hang in GIN-SDMA AllToAll for transfers >1 GiB on MI355X + ROCm 7.13 (NCCL_GIN_TYPE=5/SDMA). A single fused COPY_LINEAR_WAIT_SIGNAL descriptor at/above 256 MiB/peer stalls the SDMA engine, so the copy never lands and its SignalInc never fires, leaving every rank spinning in waitSignal.

The size cap now lives inside the Anvil-SDMA backend Put API, not in the shared rccl-tests A2A kernels. GinAlltoAllKernel/HybridAlltoAllKernel are the entry points for all GIN backends (proxy, rocSHMEM GDA, Anvil-SDMA), so capping there wrongly throttled proxy/GDA and exposed an app-visible ginPutChunked() helper with no NCCL equivalent. Moving the cap into the backend scopes the fix to SDMA and keeps chunking out of user code.

What changed

• ncclGinApi_Put<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> (gin_anvil_sdma.h): a large put is split into <= kGinPutSegBytes (128 MiB) SDMA copies. The fused signal (or the explicit SignalInc) is carried only on the final segment. The backend runs a single in-order SDMA queue per peer (numChannels forced to 1, see 852d20d4b9), so the trailing signal still means the whole message has landed. A <=128 MiB put stays a single segment, so the common path is unchanged.
• gin_anvil_sdma_put_policy.h (new device header): ships the two put-size limits — kGinPutMaxBytes (1 GiB, the 30-bit descriptor count bound) and kGinSdmaSafeCopyBytes (128 MiB, the empirical single-copy limit) — plus the pure segment math (ginPutSegmentCount(), ginPutSegmentAt()). Host-compilable (GIN_SDMA_HOST_ONLY) so the unit test links without a GPU. Added to the rccl install manifest.
• Host unit test relocated to projects/rccl/test/gin/gin_sdma_policy_test.cpp and wired into the rccl test suite (rccl-UnitTestsGinSdmaPolicy): asserts segment counts and contiguous, single-final-signal tiling over {0, 128MiB, 128MiB+1, 256MiB, 1GiB, 2GiB}.
• rccl-tests reverted to plain gin.put(): dropped ginPutChunked, the rccl-tests gin_sdma_collective_policy.h, and the BUILD_TESTS wiring.

Why 128 MiB and not 1 GiB

1 GiB is the hard correctness bound (30-bit descriptor count). 128 MiB is the empirical reliability bound: on MI355X + ROCm 7.13 a fused copy+signal at/above 256 MiB stalls the engine. 128 MiB is the largest tier that never stalls, with no measured bandwidth loss vs. the 1 GiB case.

Test plan

• Host unit test (rccl-UnitTestsGinSdmaPolicy): segment math correct over the boundary set, exactly one final signal — 5/5 pass, warning-clean.
• alltoall_perf (SDMA) on 8x MI355X (gfx950), 256 MiB -> 2 GiB total: no hang at any size (incl. 256M/512M/1G/2G per-peer), #wrong = 0, ~426 GB/s busbw at 2 GiB. Segmentation exercised (2 GiB total -> 256 MiB/peer -> 2 x 128 MiB segments).
• NCCL_GIN_TYPE=2 (proxy) / 4 (GDA) unaffected — now guaranteed by construction (kernels use plain gin.put()); clean datapoint pending a host with a working RDMA stack.

Made with Cursor

[pvallem]

There's no regression test for the chunking here. Could we add the following test cases?

1. A host unit test of the segment math (pure integer arithmetic, no GPU needed) over bytes = {0, 128MiB, 128MiB+1, 256MiB, 1GiB, 2GiB}, asserting exactly one signal is carried, no segment exceeds 128MiB, and the segments sum back to bytes. (gin_sdma_collective_policy.h already advertises a GIN_SDMA_HOST_ONLY GoogleTest path that is defined but currently unused.)

2. An AllToAll correctness test above 128MiB/peer (exercises the new multi-segment loop) and above 1GiB/peer (crosses the 1GiB 30-bit truncation boundary), using exact-integer datatypes so a truncated or stale tail can't be masked by float tolerance.

3. The original 2GiB-total AllToAll hang repro as a guard that it now completes (MI355X + ROCm 7.13, NCCL_GIN_TYPE=5).

Today there's no AllToAll test in rccl-tests and the pytest byte_range caps at 128M total, so nothing reaches the new multi-segment path.

[dlamd1dai]

Sure. I will add my test code into the "rccl-tests/" directory.

[dlamd1dai]

Added all three in 73e0a47.

1. Host unit test (no GPU). I factored the segmentation arithmetic out of the __device__ ginPutChunked into two pure functions in gin_sdma_collective_policy.h — ginPutSegmentCount() and ginPutSegmentAt() (plus kGinPutSegBytes = min(reliability, correctness caps)) — and rewrote the device loop to drive off them, so the code the test exercises is exactly what runs on device. test/gin_sdma_policy_test.cpp (GoogleTest, compiled host-only via GIN_SDMA_HOST_ONLY) asserts over bytes = {0, 1, 128MiB, 128MiB+1, 256MiB, 1GiB, 1GiB+1, 2GiB} — plus a dense +/-2 B scan around every cap multiple — that no segment exceeds 128 MiB, the segments tile [0,bytes) contiguously and sum back to bytes, and exactly one final (signal-carrying) segment lands on the tail. Wired opt-in via -DBUILD_TESTS=ON (find/fetch GTest), ctest label unit. (This re-establishes the GIN_SDMA_HOST_ONLY path you noted — it's now actually used.)

2. AllToAll correctness (GPU). test/test_AllToAll.py drives the real GinHybridAlltoAllKernel (-D 3, NCCL_GIN_TYPE=5, SDMA tier forced via NCCL_GIN_ANVIL_SDMA_THRESHOLD*=0) at 256 MiB/peer (>128 MiB → multi-segment loop) and 2 GiB/peer (>1 GiB → crosses the 30-bit truncation boundary), with exact-integer dtypes (int32/int64/uint8) and the built-in data check (-c 1), so a truncated/stale tail can't hide under float tolerance (nonzero exit on any wrong element).

3. Hang guard (GPU). The original 2 GiB-total repro (256 MiB/peer at 8 ranks) is a dedicated completion sentinel: a subprocess timeout converts a reintroduced SDMA hang into a test failure instead of stalling the runner.

The two GPU tests need an 8x MI355X-class node + MPI + a GIN-SDMA build, so the module skips unless RCCL_TESTS_GIN_SDMA_A2A=1 (size/np/timeout configurable via env); the host unit test runs anywhere. The device loop is behavior-identical to the prior do/while, so no functional change to the shipping path.

[dlamd1dai]

Ran both test tiers on hardware:

Host unit test — built googletest 1.14.0 and compiled gin_sdma_policy_test.cpp host-only (-DGIN_SDMA_HOST_ONLY -Wall -Wextra): all 5 cases pass, warning-clean (fixed one -Wdangling-else in 579861c).

GPU AllToAll (8x MI355X, ROCm 7.13, NCCL_GIN_TYPE=5, force-single-channel) — -D 3, SDMA tier forced, int32, built-in data check:

case	per-peer	segments	busbw	#wrong
>128 MiB/peer	256 MiB	2×128 MiB	~424 GB/s	0
>1 GiB/peer	2 GiB	16×128 MiB (crosses 30-bit boundary)	~428 GB/s	0
2 GiB-total hang guard	256 MiB	2×128 MiB	~424 GB/s	0, no hang

Running the suite on-hardware surfaced two gaps I fixed in 51c2af8:

• The test now forwards the env the GIN-SDMA backend actually needs to initialize (always NCCL_GIN_ENABLE=1, plus deployment-specific vars via RCCL_TESTS_A2A_XENV and launcher opts via RCCL_TESTS_MPI_OPTS) — the minimal 3-var version couldn't bring GIN up on the cluster.
• On the hang timeout it now SIGKILLs the whole launcher process group (new session), so a wedged run can't orphan ranks that pin the GPUs / exhaust SDMA queues for later cases.

[dlamd1dai]

Added and relocated. The host unit test is now projects/rccl/test/gin/gin_sdma_policy_test.cpp, wired into the rccl suite as rccl-UnitTestsGinSdmaPolicy. It exercises the pure segment math (ginPutSegmentCount()/ginPutSegmentAt(), no GPU) over {0, 128MiB, 128MiB+1, 256MiB, 1GiB, 2GiB} and asserts contiguous tiling with exactly one final (signal-carrying) segment — 5/5 pass, warning-clean.

[Kapil-Shyam-Pawar]

GinAlltoAllKernel and HybridAlltoAllKernel are the entry points for the other GIN backends as well, not just SDMA, so this will cap puts for proxy, and rocshmem gda too

It would be good to validate with NCCL_GIN_TYPE=2 and 4

Would it make sense to put the chunking in ncclGinApi_Put<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> instead, the way the proxy backend already does?
That would fix the hang for all GIN SDMA callers and not just alltoall_perf

rocm-systems/projects/rccl/src/include/nccl_device/gin/proxy/gin_proxy.h

Lines 214 to 242 in ea5a1ab

	NCCL_DEVICE_INLINE void put(Coop coop, ncclGinProxyGfd_t* gfd, ncclGinProxyGpuCtx_t* proxyCtx, int peer,
	ncclGinWindow_t dstWnd, size_t dstOff, T srcVal, bool hasInline, ncclGinWindow_t srcWnd,
	size_t srcOff, size_t bytes, ncclGinSignalDescriptor signal, ncclGinSignalOp_t signalOp,
	uint64_t signalVal, bool hasCounter, ncclGinCounter_t counterId,
	cuda::thread_scope required, cuda::thread_scope given) {
	// Release the source to system scope before the GFD is published, since the NIC
	// DMA-reads it. hip_compat.h orders thread_scope with system widest (largest), so
	// "narrower than system" is given < system (upstream libcu++ uses the opposite order).
	#if defined(__HIP_PLATFORM_AMD__)
	if ((int)given < (int)cuda::thread_scope_system) {
	#else
	if ((int)given > (int)cuda::thread_scope_system) {
	#endif
	cuda::atomic_thread_fence(cuda::memory_order_release, cuda::thread_scope_system);
	}
	using nccl::gin::proxy::DataChunkSize;
	while (bytes > DataChunkSize) {
	ncclGinProxyOp_t op;
	constructProxyOp(op, /*isGet*/ false, /*isFlush*/ false, /*hasInline*/ false, NCCL_GIN_SIGNAL_TYPE_NONE, signalOp,
	/*hasCounter*/ false);
	nccl::gin::proxy::buildGfd(gfd, op, /*srcVal*/ 0, /*hasInline*/ false, srcOff, srcWnd, dstOff, dstWnd,
	DataChunkSize, /*counterId*/ 0, /*signalId*/ 0,
	/*signalVal*/ 0, nullptr, 0);
	nccl::gin::proxy::postGfd<Coop>(coop, proxyCtx, gfd, peer);
	bytes -= DataChunkSize;
	srcOff += DataChunkSize;
	dstOff += DataChunkSize;
	}

[abouteiller]

The GDA gin has the same limitation (1GB message size), but I agree with Nusrat that this should not surface in user code.

[dlamd1dai]

Good catch — you're right that GinAlltoAllKernel/HybridAlltoAllKernel are shared by all GIN backends, so capping there was not necessary for proxy and GDA. Fixed in 92b25f7: the size cap now lives inside the SDMA backend's Put API (ncclGinApi_Put<NCCL_NET_DEVICE_GIN_ANVIL_SDMA>), and the kernels are back to a plain gin.put(). Proxy (type 2) and GDA (type 4) are no longer capped by construction — they don't touch the SDMA segmentation path at all.

I haven't been able to get a clean type 2/4 datapoint yet (the RDMA host I have, dell300x, has a broken bnxt_re ABI in-container), but since the kernel path for those types is now byte-for-byte the pre-PR gin.put(), there's no behavioral change to validate. SDMA (type 5) is validated on 8x MI355X up to 2 GiB — no hang, 0 wrong, ~426 GB/s.

[dlamd1dai]

Agreed on both points. As of 92b25f7 the chunking is entirely inside the SDMA backend gin.put() and no longer in user/test code. On the GDA 1 GiB limit: this PR scopes its fix to SDMA only (GDA's kernel path is unchanged here). The same internal-segmentation pattern could be applied to the GDA backend's Put as a follow-up if we want gin.put() to be size-agnostic there too — happy to file that separately.

[nusislam]

Will applications call gin.puChunked()? This looks like a new API that has been added in RCCL and not present in NCCL. The ideal solution would be to gin.put() internally calling gin.putChunked() for large messages. Otherwise, if an application calls gin.put() with large sizes, it may still hit the hang issue.

[dlamd1dai]

Agreed, and that's exactly what 92b25f7 does. There's no ginPutChunked() anymore — it's removed from rccl-tests. Segmentation is now internal to gin.put() for the SDMA backend: a large put is split into <=128 MiB SDMA copies inside ncclGinApi_Put<ANVIL_SDMA>, invisible to callers. No new user-facing API, and it matches NCCL's gin.put() surface.

[abouteiller]

This is entirely empirical do we have any understanding why the limit is not 1G as well here?

[dlamd1dai]

It's two different limits. 1 GiB (kGinPutMaxBytes) is the hard correctness bound — the SDMA linear-copy count field is 30 bits. 128 MiB (kGinSdmaSafeCopyBytes) is the empirical reliability bound: on MI355X + ROCm 7.13 a fused COPY_LINEAR_WAIT_SIGNAL at/above 256 MiB/peer stalls the engine (the original hang). 128 MiB is the largest power-of-two tier below that stall threshold, and it shows no measured bandwidth loss vs. the 1 GiB case (~426 GB/s at 2 GiB). The code takes min(kGinSdmaSafeCopyBytes, kGinPutMaxBytes) so if the empirical limit is ever raised it can never exceed the descriptor bound.

[pvallem]

Suggested change

	#include <gtest/gtest.h>
	#include <cstddef>
	#include <vector>
	#include "gin_sdma_collective_policy.h"
	#include <cstddef>
	#include <vector>
	#include <gtest/gtest.h>
	#include "gin_sdma_collective_policy.h"

[dlamd1dai]

Adopted, with the path updated for the relocation: the test now lives under projects/rccl/test/gin/ and includes the shipped nccl_device/gin/anvil_sdma/gin_anvil_sdma_put_policy.h (the old gin_sdma_collective_policy.h is gone).

[dlamd1dai]

Follow-up work is stacked on top of this PR: dlamd1dai#1 — "rccl-tests/gin-sdma: in-device (wall_clock64) A2A timing mechanism". It is based on this PR's head branch (users/dondai/gin-sdma-a2a-meta-repro) so it shows only the single devtime commit; it should merge after this PR.

[dlamd1dai]

@abouteiller @pvallem — this PR is approved and MERGEABLE; the only red check (GIN / Single node gin testing) is a runner-infra failure unrelated to this PR's change, so I believe it's ready to merge.

Root cause of the red check. The build and SLURM submission succeed; the sole failure in the compute-node logs (node mia1-p01-g22) is rocSHMEM-GDA queue-pair creation:

gin_plugin_rocshmem_gda.cc:118 NCCL WARN GIN rocshmem-gda: failed to create QPs
  ← libmlx5.so: undefined symbol: mlx5dv_get_data_direct_sysfs_path, version MLX5_1.25
  ← libmlx5.so: undefined symbol: mlx5dv_reg_dmabuf_mr, version MLX5_1.25

The node's libmlx5/rdma-core is too old to expose the MLX5_1.25 symbols the GDA transport needs, so NCCL_GIN_TYPE=4 (rocSHMEM-GDA/RDMA) can't create QPs. Notes:

• This is in gin_plugin_rocshmem_gda.cc, which this PR does not touch — rccl-tests/gin-sdma: cap A2A GIN put to 128 MiB to fix >1 GiB SDMA hang #9927 only changes the Anvil-SDMA A2A put path (128 MiB segmentation). The SDMA/proxy paths connected fine in the same run.
• The checks that exercise this PR's change are green: TheRock RCCL CI Linux (gfx94X-dcgpu) / Test single-node ✅ and device-API / Single node device-api testing ✅.

Corroborating validation of the SDMA fix:

• 8× MI300X (gfx942): GIN-SDMA A2A -D3 sweep to 2 GiB total (256 MiB/peer → 2×128 MiB segments) — no hang, #wrong=0, Out of bounds values : 0 OK.
• Host regression test rccl-UnitTestsGinSdmaPolicy: 5/5 pass (segment-count/boundary/tiling/final-segment/zero-cap-guard).
• 8× MI355X (gfx950) per the PR test plan: 256 MiB/peer and the 2 GiB-total guard pass with #wrong=0, busbw unchanged.

Request: since the state is UNSTABLE (not BLOCKED) with an approval, this red check is non-required — could you merge as-is, or advise if you'd prefer I re-run the GIN check first? Separately, I'll flag the libmlx5/MLX5_1.25 gap on mia1-p01-g22 to the CI owners so the GDA path is exercised on future runs.

Follow-on work (devtime, GDA unit tests) is stacked on this branch and will rebase once this merges.

[abouteiller]

The merging of the target branch closed the PR automatically, work is not to be discarded just need to change the target branch to develop.

[abouteiller]

In Open MPI -host expects a comma separated list of hosts, not a hostfile, the option is --hostfile.

[abouteiller]

There is nothing specific to SDMA as far as I can tell in that tester and it could be used to validate the same feature in GDA (or other GIN transports for that matter), maybe scrubbing the ANVIL_SDMA from the comments and function names would be good.