From 7d35eb163c4447f938afa07b95cbed0887319b54 Mon Sep 17 00:00:00 2001 From: Anatolii Rozanov Date: Tue, 6 Oct 2026 12:20:23 +0200 Subject: [PATCH 1/6] test(rccl): cover dense all-to-all ncclPutSignal on the CE path Every rank puts twice to every other rank in one group and waits for all of them, once with uniform sizes and once with sizes around the CE chunk size. With four or more ranks the uneven profile is the only case that makes the copy engine split a batch into round-robin chunk waves, which no test ran before. The CI batch config runs the test in the 2-rank Host API suites and in a new single-node 8-GPU suite, since chunking needs more than one remote peer per round. Co-authored-by: Cursor --- projects/rccl/test/HostApiMPITests.cpp | 122 ++++++++++++++++++ .../configs/mi300x_mellanox_ib.json | 24 ++++ 2 files changed, 146 insertions(+) diff --git a/projects/rccl/test/HostApiMPITests.cpp b/projects/rccl/test/HostApiMPITests.cpp index cc686795496a..b3df6b40a14a 100644 --- a/projects/rccl/test/HostApiMPITests.cpp +++ b/projects/rccl/test/HostApiMPITests.cpp @@ -24,6 +24,8 @@ * M1 - TwoCommunicatorsIndependentWindows: two comms, independent windows * M2 - StressManySmallPuts: 100×64-byte PUTs, opCnt=100 wait * P6 - PutToSelf: PUT to own window (peer=self loopback) + * P7 - DenseAllToAllPutSignal: every rank puts to every peer in one group, + * uniform and uneven sizes * W3 - DoubleDeregister: ncclCommWindowDeregister twice on same handle * W4 - DestroyCommWithoutDeregister: window auto-freed during commFree * R1 - RmaDisableIntraNodePutUsesCe: NCCL_RMA_DISABLE=1 leaves LSA puts working @@ -69,6 +71,7 @@ #include #include #include +#include using namespace MPITestConstants; using namespace RCCLTestGuards; @@ -1429,6 +1432,125 @@ TEST_F(HostApiTest, AllToAllPut) TEST_INFO("P5 rank %d: AllToAllPut passed (recv from rank %d).", myRank, recvFrom); } +// ============================================================================ +// P7 — DenseAllToAllPutSignal +// ============================================================================ + +namespace +{ +constexpr int kA2aPutsPerPeer = 2; +// Put k reads its source k * kA2aSrcShift bytes in, so each put carries its own pattern. +constexpr size_t kA2aSrcShift = 64; +constexpr size_t kA2aGuard = 4096; +constexpr uint8_t kA2aSentinel = 0xA5; +constexpr size_t kA2aUniformBytes = 4u << 20; +// Around the default NCCL_CE_CHUNK_SIZE (8 MiB), none a multiple of it. +constexpr size_t kA2aUnevenBytes[] = {(3u << 20) + 4096, (9u << 20) + 12288, (12u << 20) + 262144}; +constexpr size_t kA2aSlot = 13u << 20; +static_assert(kA2aUnevenBytes[2] + kA2aSrcShift * kA2aPutsPerPeer + kA2aGuard <= kA2aSlot, + "a slot must hold the largest put plus its guard"); + +size_t a2aPutBytes(bool uneven, int src, int dst, int k) +{ + return uneven ? kA2aUnevenBytes[(src + dst + k) % 3] : kA2aUniformBytes; +} + +// Slot 0 is the send region; put k from rank src lands in its own slot. +size_t a2aRecvOffset(int src, int k) +{ + return kA2aSlot * (1 + src * kA2aPutsPerPeer + k); +} +} // namespace + +/** + * @test HostApiTest.DenseAllToAllPutSignal + * @brief Every rank puts twice to every other rank in one group, then waits for + * all of them (AICOMRCCL-1964). + * + * Runs once with uniform sizes and once with uneven sizes around the copy-engine + * chunk size; with 4+ ranks on one node the uneven profile is what makes the + * copy engine split its batches into chunks. Each put lands in its own slot with + * its own pattern, and the guard after it must stay untouched. + * + * A signal that overtook its data is caught only if the copy is still in flight + * when the receiver reads the slot back, so that ordering is checked + * probabilistically. + */ +TEST_F(HostApiTest, DenseAllToAllPutSignal) +{ + if(!validateTestPrerequisites(/*min=*/2, /*max=*/0, kNoPowerOfTwoRequired, + /*min_nodes=*/1, kRequireSingleNode)) + { + GTEST_SKIP() << "Need at least 2 MPI processes on a single node"; + } + + const int myRank = rank(); + const int nRanks_ = nRanks(); + ncclComm_t comm = getActiveCommunicator(); + hipStream_t stream = getActiveStream(); + const size_t winSize = kA2aSlot * (1 + nRanks_ * kA2aPutsPerPeer); + + void* winBuf = nullptr; + ASSERT_MPI_EQ(ncclSuccess, allocFineGrainBuffer(&winBuf, winSize)); + auto winBufGuard = makeScopeGuard([&]() { freeFineGrainBuffer(winBuf); }); + + ncclWindow_t win = nullptr; + NcclWindowGuard wg(comm, winBuf, winSize, &win, winMode()); + ASSERT_MPI_NE(win, nullptr); + ASSERT_MPI_EQ(ncclSuccess, wg.initResult()); + + uint8_t* base = static_cast(winBuf); + + std::vector descs; + for(int peer = 0; peer < nRanks_; ++peer) + { + if(peer != myRank) + descs.push_back({/*opCnt=*/kA2aPutsPerPeer, peer, kSigIdx, kCtx}); + } + + for(bool uneven : {false, true}) + { + FillSentinel(base, winSize, kA2aSentinel); + FillBuf(base, kA2aSlot, myRank); + ASSERT_MPI_SUCCESS(MPI_Barrier(MPI_COMM_WORLD)); + + ncclResult_t putRes = ncclGroupStart(); + for(int peer = 0; peer < nRanks_ && putRes == ncclSuccess; ++peer) + { + if(peer == myRank) continue; + for(int k = 0; k < kA2aPutsPerPeer && putRes == ncclSuccess; ++k) + { + putRes = ncclPutSignal(base + k * kA2aSrcShift, a2aPutBytes(uneven, myRank, peer, k), ncclUint8, + peer, win, a2aRecvOffset(myRank, k), + kSigIdx, kCtx, kFlags, comm, stream); + } + } + ncclResult_t endRes = ncclGroupEnd(); + if(putRes == ncclSuccess) putRes = endRes; + ASSERT_MPI_EQ(ncclSuccess, putRes); + + ncclResult_t waitRes = ncclWaitSignal(static_cast(descs.size()), descs.data(), comm, stream); + ASSERT_MPI_EQ(ncclSuccess, waitRes); + ASSERT_MPI_EQ(hipSuccess, hipStreamSynchronize(stream)); + + bool ok = true; + for(int src = 0; src < nRanks_ && ok; ++src) + { + if(src == myRank) continue; + for(int k = 0; k < kA2aPutsPerPeer && ok; ++k) + { + const size_t bytes = a2aPutBytes(uneven, src, myRank, k); + const uint8_t* slot = base + a2aRecvOffset(src, k); + ok = VerifyBuf(slot, bytes, static_cast(src + k * kA2aSrcShift)) && + AllSentinel(slot + bytes, kA2aGuard, kA2aSentinel); + } + } + ASSERT_MPI_TRUE(ok); + + TEST_INFO("P7 rank %d: DenseAllToAllPutSignal passed (%s sizes).", myRank, uneven ? "uneven" : "uniform"); + } +} + // ============================================================================ // O2 — SignalImpliesPriorPutsDelivered // ============================================================================ diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json index 0e7febb39672..b758ed6d2674 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json @@ -992,6 +992,7 @@ {"name": "HostApi_WaitSignalMultipleDescriptors", "description": "Host API: wait signal multiple descriptors", "test_filter": "HostApiTest.WaitSignalMultipleDescriptors"}, {"name": "HostApi_LargePut", "description": "Host API: large put", "test_filter": "HostApiTest.LargePut"}, {"name": "HostApi_AllToAllPut", "description": "Host API: all-to-all put", "test_filter": "HostApiTest.AllToAllPut"}, + {"name": "HostApi_DenseAllToAllPutSignal", "description": "Host API: dense all-to-all put/signal, uniform and uneven sizes", "test_filter": "HostApiTest.DenseAllToAllPutSignal"}, {"name": "HostApi_SignalImpliesPriorPutsDelivered", "description": "Host API: signal implies prior puts delivered", "test_filter": "HostApiTest.SignalImpliesPriorPutsDelivered"}, {"name": "HostApi_PutSignalInvalidCtx", "description": "Host API: put signal invalid ctx", "test_filter": "HostApiTest.PutSignalInvalidCtx"}, {"name": "HostApi_WaitSignalNullDescs", "description": "Host API: wait signal null descs", "test_filter": "HostApiTest.WaitSignalNullDescs"}, @@ -1056,6 +1057,20 @@ {"name": "HostApi_RmaDisableCrossNodePutFails_NonSym", "description": "Host API: NCCL_RMA_DISABLE=1, non-symmetric window, cross-node put/wait fail with ncclInvalidUsage instead of using the proxy", "test_filter": "HostApiTest.RmaDisableCrossNodePutFails", "env_variables": {"NCCL_CUMEM_ENABLE": "0"}} ] }, + "host_api_dense_a2a_8rank": { + "extends": "net_ib_base", + "num_ranks": 8, + "num_nodes": 1, + "num_gpus": 8, + "description": "Dense all-to-all ncclPutSignal on all 8 GPUs of one node. Only with 4+ ranks does the uneven profile make the copy engine split its batches into chunks, which the 2-rank Host API suites never reach.", + "env_variables": { + "NCCL_CUMEM_ENABLE": "0", + "RCCL_RMA_USE_DMABUF": "1" + }, + "tests": [ + {"name": "HostApi_DenseAllToAllPutSignal_8rank", "description": "Host API: dense all-to-all put/signal on 8 GPUs (copy-engine chunking)", "test_filter": "HostApiTest.DenseAllToAllPutSignal"} + ] + }, "rma_external_plugin_put_signal": { "extends": "default", @@ -5159,6 +5174,15 @@ "config": "host_api_tests_default", "enabled": true }, + { + "name": "Host API Tests - Dense All-to-All (8 GPUs)", + "description": "Dense all-to-all put/signal over the copy engine on one node, covering chunked batches", + "config": "host_api_dense_a2a_8rank", + "num_ranks": 8, + "num_nodes": 1, + "num_gpus": 8, + "enabled": true + }, { "name": "Host API Graph Capture - Proxy", "description": "Two-node proxy graph capture and repeated replay for PutSignal/WaitSignal", From e015947e030300015464879142078a233e52c86d Mon Sep 17 00:00:00 2001 From: Anatolii Rozanov Date: Tue, 6 Oct 2026 15:27:40 +0200 Subject: [PATCH 2/6] test(rccl): run the 8-rank dense all-to-all in the existing MPI suite Move the 8-rank DenseAllToAllPutSignal run into mpi_8rank_single_node, which already hosts MPI tests that need eight ranks on one node, and drop the dedicated host_api_dense_a2a_8rank config and its suite. Co-authored-by: Cursor --- .../configs/mi300x_mellanox_ib.json | 29 ++++--------------- 1 file changed, 5 insertions(+), 24 deletions(-) diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json index b758ed6d2674..e86a36b50e7c 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json @@ -1057,21 +1057,6 @@ {"name": "HostApi_RmaDisableCrossNodePutFails_NonSym", "description": "Host API: NCCL_RMA_DISABLE=1, non-symmetric window, cross-node put/wait fail with ncclInvalidUsage instead of using the proxy", "test_filter": "HostApiTest.RmaDisableCrossNodePutFails", "env_variables": {"NCCL_CUMEM_ENABLE": "0"}} ] }, - "host_api_dense_a2a_8rank": { - "extends": "net_ib_base", - "num_ranks": 8, - "num_nodes": 1, - "num_gpus": 8, - "description": "Dense all-to-all ncclPutSignal on all 8 GPUs of one node. Only with 4+ ranks does the uneven profile make the copy engine split its batches into chunks, which the 2-rank Host API suites never reach.", - "env_variables": { - "NCCL_CUMEM_ENABLE": "0", - "RCCL_RMA_USE_DMABUF": "1" - }, - "tests": [ - {"name": "HostApi_DenseAllToAllPutSignal_8rank", "description": "Host API: dense all-to-all put/signal on 8 GPUs (copy-engine chunking)", "test_filter": "HostApiTest.DenseAllToAllPutSignal"} - ] - }, - "rma_external_plugin_put_signal": { "extends": "default", "is_gtest": true, @@ -3052,6 +3037,11 @@ "env_variables": { "NCCL_GRAPH_STREAM_ORDERING": "0" } + }, + { + "name": "HostApi_DenseAllToAllPutSignal_8Ranks", + "description": "Dense all-to-all put/signal over the copy engine; with 8 ranks the uneven profile splits batches into chunks, which the 2-rank Host API suites never reach", + "test_filter": "HostApiTest.DenseAllToAllPutSignal" } ] }, @@ -5174,15 +5164,6 @@ "config": "host_api_tests_default", "enabled": true }, - { - "name": "Host API Tests - Dense All-to-All (8 GPUs)", - "description": "Dense all-to-all put/signal over the copy engine on one node, covering chunked batches", - "config": "host_api_dense_a2a_8rank", - "num_ranks": 8, - "num_nodes": 1, - "num_gpus": 8, - "enabled": true - }, { "name": "Host API Graph Capture - Proxy", "description": "Two-node proxy graph capture and repeated replay for PutSignal/WaitSignal", From 4aa66582d2a99ef8030b59ece88f90c9fdbd843b Mon Sep 17 00:00:00 2001 From: Anatolii Rozanov Date: Tue, 6 Oct 2026 15:58:43 +0200 Subject: [PATCH 3/6] test(rccl): give every dense all-to-all put a unique data pattern The puts read a byte pattern that repeated every 256 bytes and did not depend on the destination, so a chunk delivered to the wrong peer or at the wrong chunk offset left the data unchanged. Give each put its own source slot and fill it with 64-bit words that encode the sender, the receiver, the put index and the word index. Chunking already starts at three ranks, so say 3+ in the test header. Co-authored-by: Cursor --- projects/rccl/test/HostApiMPITests.cpp | 70 ++++++++++++++++++++------ 1 file changed, 56 insertions(+), 14 deletions(-) diff --git a/projects/rccl/test/HostApiMPITests.cpp b/projects/rccl/test/HostApiMPITests.cpp index b3df6b40a14a..b38687a22d92 100644 --- a/projects/rccl/test/HostApiMPITests.cpp +++ b/projects/rccl/test/HostApiMPITests.cpp @@ -1439,26 +1439,62 @@ TEST_F(HostApiTest, AllToAllPut) namespace { constexpr int kA2aPutsPerPeer = 2; -// Put k reads its source k * kA2aSrcShift bytes in, so each put carries its own pattern. -constexpr size_t kA2aSrcShift = 64; constexpr size_t kA2aGuard = 4096; constexpr uint8_t kA2aSentinel = 0xA5; constexpr size_t kA2aUniformBytes = 4u << 20; // Around the default NCCL_CE_CHUNK_SIZE (8 MiB), none a multiple of it. constexpr size_t kA2aUnevenBytes[] = {(3u << 20) + 4096, (9u << 20) + 12288, (12u << 20) + 262144}; constexpr size_t kA2aSlot = 13u << 20; -static_assert(kA2aUnevenBytes[2] + kA2aSrcShift * kA2aPutsPerPeer + kA2aGuard <= kA2aSlot, - "a slot must hold the largest put plus its guard"); +static_assert(kA2aUnevenBytes[2] + kA2aGuard <= kA2aSlot, "a slot must hold the largest put plus its guard"); size_t a2aPutBytes(bool uneven, int src, int dst, int k) { return uneven ? kA2aUnevenBytes[(src + dst + k) % 3] : kA2aUniformBytes; } -// Slot 0 is the send region; put k from rank src lands in its own slot. -size_t a2aRecvOffset(int src, int k) +// The first nRanks * kA2aPutsPerPeer slots hold what this rank sends, the rest what it receives. +size_t a2aSendOffset(int dst, int k) { - return kA2aSlot * (1 + src * kA2aPutsPerPeer + k); + return kA2aSlot * (dst * kA2aPutsPerPeer + k); +} + +size_t a2aRecvOffset(int nRanks, int src, int k) +{ + return kA2aSlot * ((nRanks + src) * kA2aPutsPerPeer + k); +} + +// Word i of put k from src to dst; no two words of the test are equal, so a +// chunk that lands at the wrong offset or at the wrong peer changes the data. +uint64_t a2aWord(int nRanks, int src, int dst, int k, size_t i) +{ + const uint64_t op = (static_cast(src) * nRanks + dst) * kA2aPutsPerPeer + k; + return op << 32 | i; +} + +void a2aFill(void* buf, size_t bytes, int nRanks, int src, int dst, int k) +{ + std::vector words(bytes / sizeof(uint64_t)); + for(size_t i = 0; i < words.size(); ++i) + words[i] = a2aWord(nRanks, src, dst, k, i); + ASSERT_EQ(hipMemcpy(buf, words.data(), bytes, hipMemcpyHostToDevice), hipSuccess); +} + +bool a2aVerify(const void* buf, size_t bytes, int nRanks, int src, int dst, int k) +{ + std::vector words(bytes / sizeof(uint64_t)); + if(hipMemcpy(words.data(), buf, bytes, hipMemcpyDeviceToHost) != hipSuccess) + return false; + for(size_t i = 0; i < words.size(); ++i) + { + if(words[i] != a2aWord(nRanks, src, dst, k, i)) + { + fprintf(stderr, "a2aVerify: put %d from rank %d to rank %d, word %zu: expected %#llx got %#llx\n", + k, src, dst, i, static_cast(a2aWord(nRanks, src, dst, k, i)), + static_cast(words[i])); + return false; + } + } + return true; } } // namespace @@ -1468,7 +1504,7 @@ size_t a2aRecvOffset(int src, int k) * all of them (AICOMRCCL-1964). * * Runs once with uniform sizes and once with uneven sizes around the copy-engine - * chunk size; with 4+ ranks on one node the uneven profile is what makes the + * chunk size; with 3+ ranks on one node the uneven profile is what makes the * copy engine split its batches into chunks. Each put lands in its own slot with * its own pattern, and the guard after it must stay untouched. * @@ -1488,7 +1524,7 @@ TEST_F(HostApiTest, DenseAllToAllPutSignal) const int nRanks_ = nRanks(); ncclComm_t comm = getActiveCommunicator(); hipStream_t stream = getActiveStream(); - const size_t winSize = kA2aSlot * (1 + nRanks_ * kA2aPutsPerPeer); + const size_t winSize = kA2aSlot * 2 * nRanks_ * kA2aPutsPerPeer; void* winBuf = nullptr; ASSERT_MPI_EQ(ncclSuccess, allocFineGrainBuffer(&winBuf, winSize)); @@ -1511,7 +1547,13 @@ TEST_F(HostApiTest, DenseAllToAllPutSignal) for(bool uneven : {false, true}) { FillSentinel(base, winSize, kA2aSentinel); - FillBuf(base, kA2aSlot, myRank); + for(int peer = 0; peer < nRanks_; ++peer) + { + if(peer == myRank) continue; + for(int k = 0; k < kA2aPutsPerPeer; ++k) + a2aFill(base + a2aSendOffset(peer, k), a2aPutBytes(uneven, myRank, peer, k), nRanks_, myRank, peer, k); + } + ASSERT_MPI_EQ(hipSuccess, hipDeviceSynchronize()); ASSERT_MPI_SUCCESS(MPI_Barrier(MPI_COMM_WORLD)); ncclResult_t putRes = ncclGroupStart(); @@ -1520,8 +1562,8 @@ TEST_F(HostApiTest, DenseAllToAllPutSignal) if(peer == myRank) continue; for(int k = 0; k < kA2aPutsPerPeer && putRes == ncclSuccess; ++k) { - putRes = ncclPutSignal(base + k * kA2aSrcShift, a2aPutBytes(uneven, myRank, peer, k), ncclUint8, - peer, win, a2aRecvOffset(myRank, k), + putRes = ncclPutSignal(base + a2aSendOffset(peer, k), a2aPutBytes(uneven, myRank, peer, k), ncclUint8, + peer, win, a2aRecvOffset(nRanks_, myRank, k), kSigIdx, kCtx, kFlags, comm, stream); } } @@ -1540,8 +1582,8 @@ TEST_F(HostApiTest, DenseAllToAllPutSignal) for(int k = 0; k < kA2aPutsPerPeer && ok; ++k) { const size_t bytes = a2aPutBytes(uneven, src, myRank, k); - const uint8_t* slot = base + a2aRecvOffset(src, k); - ok = VerifyBuf(slot, bytes, static_cast(src + k * kA2aSrcShift)) && + const uint8_t* slot = base + a2aRecvOffset(nRanks_, src, k); + ok = a2aVerify(slot, bytes, nRanks_, src, myRank, k) && AllSentinel(slot + bytes, kA2aGuard, kA2aSentinel); } } From f56e61c2d76ef62dbad12d380313c9687c698ae5 Mon Sep 17 00:00:00 2001 From: Anatolii Rozanov Date: Wed, 7 Oct 2026 12:02:34 +0200 Subject: [PATCH 4/6] test(rccl): tighten the dense all-to-all put/signal checks Check that a rank's own receive slots stay untouched, since no put targets them, so a chunk landing there no longer goes unnoticed. Derive the size index and the slot bound from the size array, and assert that every put is a whole number of 64-bit words, which the fill and verify helpers rely on. Name the size profile in failure output. Pin NCCL_CUMEM_ENABLE=1 and RCCL_CE_BATCH_ASYNC_ENABLE=1 on the 8-rank CI entry so it runs on a symmetric window and the batch path it is meant to cover, and stop its description from promising chunking. Co-authored-by: Cursor --- projects/rccl/test/HostApiMPITests.cpp | 38 ++++++++++++++----- .../configs/mi300x_mellanox_ib.json | 8 +++- 2 files changed, 35 insertions(+), 11 deletions(-) diff --git a/projects/rccl/test/HostApiMPITests.cpp b/projects/rccl/test/HostApiMPITests.cpp index b38687a22d92..10303aaa8a31 100644 --- a/projects/rccl/test/HostApiMPITests.cpp +++ b/projects/rccl/test/HostApiMPITests.cpp @@ -21,11 +21,11 @@ * E2 - PutSignalNullWindow: null peerWin → error * E3 - PutSignalOffsetOutOfBounds: offset past end → error (or skip) * E4 - PutSignalInvalidSigIdx: sigIdx=1 → error (or skip) + * P7 - DenseAllToAllPutSignal: every rank puts to every peer in one group, + * uniform and uneven sizes * M1 - TwoCommunicatorsIndependentWindows: two comms, independent windows * M2 - StressManySmallPuts: 100×64-byte PUTs, opCnt=100 wait * P6 - PutToSelf: PUT to own window (peer=self loopback) - * P7 - DenseAllToAllPutSignal: every rank puts to every peer in one group, - * uniform and uneven sizes * W3 - DoubleDeregister: ncclCommWindowDeregister twice on same handle * W4 - DestroyCommWithoutDeregister: window auto-freed during commFree * R1 - RmaDisableIntraNodePutUsesCe: NCCL_RMA_DISABLE=1 leaves LSA puts working @@ -71,6 +71,7 @@ #include #include #include +#include #include using namespace MPITestConstants; @@ -1445,11 +1446,23 @@ constexpr size_t kA2aUniformBytes = 4u << 20; // Around the default NCCL_CE_CHUNK_SIZE (8 MiB), none a multiple of it. constexpr size_t kA2aUnevenBytes[] = {(3u << 20) + 4096, (9u << 20) + 12288, (12u << 20) + 262144}; constexpr size_t kA2aSlot = 13u << 20; -static_assert(kA2aUnevenBytes[2] + kA2aGuard <= kA2aSlot, "a slot must hold the largest put plus its guard"); + +constexpr bool a2aPutFitsSlot(size_t bytes) +{ + return bytes % sizeof(uint64_t) == 0 && bytes + kA2aGuard <= kA2aSlot; +} + +constexpr bool a2aAllPutsFitSlot() +{ + for(size_t bytes : kA2aUnevenBytes) + if(!a2aPutFitsSlot(bytes)) return false; + return a2aPutFitsSlot(kA2aUniformBytes); +} +static_assert(a2aAllPutsFitSlot(), "every put must be whole 64-bit words and fit its slot with the guard"); size_t a2aPutBytes(bool uneven, int src, int dst, int k) { - return uneven ? kA2aUnevenBytes[(src + dst + k) % 3] : kA2aUniformBytes; + return uneven ? kA2aUnevenBytes[(src + dst + k) % std::size(kA2aUnevenBytes)] : kA2aUniformBytes; } // The first nRanks * kA2aPutsPerPeer slots hold what this rank sends, the rest what it receives. @@ -1501,7 +1514,8 @@ bool a2aVerify(const void* buf, size_t bytes, int nRanks, int src, int dst, int /** * @test HostApiTest.DenseAllToAllPutSignal * @brief Every rank puts twice to every other rank in one group, then waits for - * all of them (AICOMRCCL-1964). + * all of them (AICOMRCCL-1964). P5 covers the ring-neighbour case; this + * is the dense N-to-N one. * * Runs once with uniform sizes and once with uneven sizes around the copy-engine * chunk size; with 3+ ranks on one node the uneven profile is what makes the @@ -1514,7 +1528,7 @@ bool a2aVerify(const void* buf, size_t bytes, int nRanks, int src, int dst, int */ TEST_F(HostApiTest, DenseAllToAllPutSignal) { - if(!validateTestPrerequisites(/*min=*/2, /*max=*/0, kNoPowerOfTwoRequired, + if(!validateTestPrerequisites(/*min=*/2, kNoProcessLimit, kNoPowerOfTwoRequired, /*min_nodes=*/1, kRequireSingleNode)) { GTEST_SKIP() << "Need at least 2 MPI processes on a single node"; @@ -1546,6 +1560,7 @@ TEST_F(HostApiTest, DenseAllToAllPutSignal) for(bool uneven : {false, true}) { + SCOPED_TRACE(uneven ? "uneven sizes" : "uniform sizes"); FillSentinel(base, winSize, kA2aSentinel); for(int peer = 0; peer < nRanks_; ++peer) { @@ -1578,11 +1593,16 @@ TEST_F(HostApiTest, DenseAllToAllPutSignal) bool ok = true; for(int src = 0; src < nRanks_ && ok; ++src) { - if(src == myRank) continue; for(int k = 0; k < kA2aPutsPerPeer && ok; ++k) { - const size_t bytes = a2aPutBytes(uneven, src, myRank, k); - const uint8_t* slot = base + a2aRecvOffset(nRanks_, src, k); + const uint8_t* slot = base + a2aRecvOffset(nRanks_, src, k); + if(src == myRank) + { + // No put targets this rank's own receive slots. + ok = AllSentinel(slot, kA2aSlot, kA2aSentinel); + continue; + } + const size_t bytes = a2aPutBytes(uneven, src, myRank, k); ok = a2aVerify(slot, bytes, nRanks_, src, myRank, k) && AllSentinel(slot + bytes, kA2aGuard, kA2aSentinel); } diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json index e86a36b50e7c..c530b562316a 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json @@ -3040,8 +3040,12 @@ }, { "name": "HostApi_DenseAllToAllPutSignal_8Ranks", - "description": "Dense all-to-all put/signal over the copy engine; with 8 ranks the uneven profile splits batches into chunks, which the 2-rank Host API suites never reach", - "test_filter": "HostApiTest.DenseAllToAllPutSignal" + "description": "Dense all-to-all put/signal over the copy engine on 8 ranks; the uneven profile can split batches into chunks, which the 2-rank Host API suites never reach", + "test_filter": "HostApiTest.DenseAllToAllPutSignal", + "env_variables": { + "NCCL_CUMEM_ENABLE": "1", + "RCCL_CE_BATCH_ASYNC_ENABLE": "1" + } } ] }, From f16521feb47939bc881f71325d2fa2d32bfefed1 Mon Sep 17 00:00:00 2001 From: Anatolii Rozanov Date: Wed, 7 Oct 2026 17:10:09 +0200 Subject: [PATCH 5/6] test(rccl): check the whole slot tail after each dense all-to-all put The check after each put read only the 4 KiB guard, so a chunk written to the wrong offset further into the 13 MiB slot went unnoticed. Scan the rest of the slot instead, as the own-rank slot check already does. Co-authored-by: Cursor --- projects/rccl/test/HostApiMPITests.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/rccl/test/HostApiMPITests.cpp b/projects/rccl/test/HostApiMPITests.cpp index 10303aaa8a31..0a88c8924000 100644 --- a/projects/rccl/test/HostApiMPITests.cpp +++ b/projects/rccl/test/HostApiMPITests.cpp @@ -1520,7 +1520,7 @@ bool a2aVerify(const void* buf, size_t bytes, int nRanks, int src, int dst, int * Runs once with uniform sizes and once with uneven sizes around the copy-engine * chunk size; with 3+ ranks on one node the uneven profile is what makes the * copy engine split its batches into chunks. Each put lands in its own slot with - * its own pattern, and the guard after it must stay untouched. + * its own pattern, and the rest of the slot after it must stay untouched. * * A signal that overtook its data is caught only if the copy is still in flight * when the receiver reads the slot back, so that ordering is checked @@ -1604,7 +1604,7 @@ TEST_F(HostApiTest, DenseAllToAllPutSignal) } const size_t bytes = a2aPutBytes(uneven, src, myRank, k); ok = a2aVerify(slot, bytes, nRanks_, src, myRank, k) && - AllSentinel(slot + bytes, kA2aGuard, kA2aSentinel); + AllSentinel(slot + bytes, kA2aSlot - bytes, kA2aSentinel); } } ASSERT_MPI_TRUE(ok); From c871c55d2dfb49d6df0222571422a8e0a5b46393 Mon Sep 17 00:00:00 2001 From: Anatolii Rozanov Date: Thu, 8 Oct 2026 10:54:12 +0200 Subject: [PATCH 6/6] test(rccl): fail the dense all-to-all test when filling a send buffer fails ASSERT_EQ inside a void helper returns only from the helper, so a failed hipMemcpy let the test put an unfilled buffer and report a data mismatch later. Return the result instead and assert on it once after the fill loop; the collective check makes every rank stop together. Co-authored-by: Cursor --- projects/rccl/test/HostApiMPITests.cpp | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/projects/rccl/test/HostApiMPITests.cpp b/projects/rccl/test/HostApiMPITests.cpp index 0a88c8924000..2bd15a363a0b 100644 --- a/projects/rccl/test/HostApiMPITests.cpp +++ b/projects/rccl/test/HostApiMPITests.cpp @@ -1484,12 +1484,12 @@ uint64_t a2aWord(int nRanks, int src, int dst, int k, size_t i) return op << 32 | i; } -void a2aFill(void* buf, size_t bytes, int nRanks, int src, int dst, int k) +bool a2aFill(void* buf, size_t bytes, int nRanks, int src, int dst, int k) { std::vector words(bytes / sizeof(uint64_t)); for(size_t i = 0; i < words.size(); ++i) words[i] = a2aWord(nRanks, src, dst, k, i); - ASSERT_EQ(hipMemcpy(buf, words.data(), bytes, hipMemcpyHostToDevice), hipSuccess); + return hipMemcpy(buf, words.data(), bytes, hipMemcpyHostToDevice) == hipSuccess; } bool a2aVerify(const void* buf, size_t bytes, int nRanks, int src, int dst, int k) @@ -1562,12 +1562,14 @@ TEST_F(HostApiTest, DenseAllToAllPutSignal) { SCOPED_TRACE(uneven ? "uneven sizes" : "uniform sizes"); FillSentinel(base, winSize, kA2aSentinel); - for(int peer = 0; peer < nRanks_; ++peer) + bool filled = true; + for(int peer = 0; peer < nRanks_ && filled; ++peer) { if(peer == myRank) continue; - for(int k = 0; k < kA2aPutsPerPeer; ++k) - a2aFill(base + a2aSendOffset(peer, k), a2aPutBytes(uneven, myRank, peer, k), nRanks_, myRank, peer, k); + for(int k = 0; k < kA2aPutsPerPeer && filled; ++k) + filled = a2aFill(base + a2aSendOffset(peer, k), a2aPutBytes(uneven, myRank, peer, k), nRanks_, myRank, peer, k); } + ASSERT_MPI_TRUE(filled); ASSERT_MPI_EQ(hipSuccess, hipDeviceSynchronize()); ASSERT_MPI_SUCCESS(MPI_Barrier(MPI_COMM_WORLD));