From 8a826446d61e5b6102561dda1c0010e2b8755afb Mon Sep 17 00:00:00 2001 From: Praneeth Vallem Date: Sat, 10 Oct 2026 14:07:20 -0700 Subject: [PATCH 1/7] fix(rccl): gate DDA AllGather eligibility on the gathered total ncclAllGatherDdaIpcEligible and ncclAllGatherDdaFabricEligible compared the per-rank payload with ddaScratchBytes while their launchers reject when the gathered total exceeds it. With RCCL_DDA_THRESHOLD above the scratch size the selector picked DDA and ncclAllGather returned ncclInvalidArgument instead of falling back. Check sendcount * nRanks, as the RS and A2A predicates do. Add IPC and fabric eligibility cases at and past the scratch boundary, and an 8-rank MPI test that raises RCCL_DDA_THRESHOLD above scratch and checks the call falls back with correct data, while a total that exactly fills scratch still picks DDA IPC. JIRA ID: AICOMRCCL-2788, AICOMRCCL-2775 --- .../dda/all_gather/dda_all_gather_fabric.cu | 3 +- .../dda/all_gather/dda_all_gather_ipc.cu | 3 +- projects/rccl/test/CMakeLists.txt | 1 + .../test/DdaAllGatherThresholdMPITests.cpp | 154 ++++++++++++++++++ .../rccl/test/DdaFabricEligibilityTests.cpp | 12 ++ projects/rccl/test/DdaIpcEligibilityTests.cpp | 19 +++ .../configs/mi300x_mellanox_ib.json | 8 + .../configs/mi355x_thor2_roce.json | 8 + 8 files changed, 206 insertions(+), 2 deletions(-) create mode 100644 projects/rccl/test/DdaAllGatherThresholdMPITests.cpp diff --git a/projects/rccl/src/algorithms/dda/all_gather/dda_all_gather_fabric.cu b/projects/rccl/src/algorithms/dda/all_gather/dda_all_gather_fabric.cu index 704b4e65ae91..f744f4708008 100644 --- a/projects/rccl/src/algorithms/dda/all_gather/dda_all_gather_fabric.cu +++ b/projects/rccl/src/algorithms/dda/all_gather/dda_all_gather_fabric.cu @@ -115,7 +115,8 @@ bool ncclAllGatherDdaFabricEligible(ncclComm* comm, const void* sendbuff, void* return false; } - size_t need = sendcount * ncclTypeSize(datatype); + // Gate on the gathered total: the launcher rejects totals above scratch. + size_t need = sendcount * comm->nRanks * ncclTypeSize(datatype); if (need > comm->ddaScratchBytes) { return false; } diff --git a/projects/rccl/src/algorithms/dda/all_gather/dda_all_gather_ipc.cu b/projects/rccl/src/algorithms/dda/all_gather/dda_all_gather_ipc.cu index 0a7a3449c3ed..e6ae2f542592 100644 --- a/projects/rccl/src/algorithms/dda/all_gather/dda_all_gather_ipc.cu +++ b/projects/rccl/src/algorithms/dda/all_gather/dda_all_gather_ipc.cu @@ -94,7 +94,8 @@ bool ncclAllGatherDdaIpcEligible(ncclComm* comm, const void* sendbuff, void* rec return false; } - size_t need = sendcount * ncclTypeSize(datatype); + // Gate on the gathered total: the launcher rejects totals above scratch. + size_t need = sendcount * comm->nRanks * ncclTypeSize(datatype); if (need > comm->ddaScratchBytes) { return false; } diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index 1cb43e7b9115..acac7f959602 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -1023,6 +1023,7 @@ if(BUILD_TESTS) DeviceApiMPITests.cpp CommNetMismatchMPITests.cpp test_sparse_send_recv.cpp + DdaAllGatherThresholdMPITests.cpp DdaAllReduceMPITests.cpp DdaFabricSimpleMPITests.cpp SymmetricAbortCheckModeMPITests.cpp diff --git a/projects/rccl/test/DdaAllGatherThresholdMPITests.cpp b/projects/rccl/test/DdaAllGatherThresholdMPITests.cpp new file mode 100644 index 000000000000..3be2744b3e2f --- /dev/null +++ b/projects/rccl/test/DdaAllGatherThresholdMPITests.cpp @@ -0,0 +1,154 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +/** + * @file DdaAllGatherThresholdMPITests.cpp + * @brief DDA IPC AllGather with RCCL_DDA_THRESHOLD above the IPC scratch. + * + * Run (example): + * mpirun -np 8 -x RCCL_DDA_THRESHOLD=1073741824 ./rccl-UnitTestsMPI + * --gtest_filter=DdaMPI_AllGatherThreshold.* + */ + +#ifdef MPI_TESTS_ENABLED + +#include "DeviceBufferHelpers.hpp" +#include "MPIHelpers.hpp" +#include "MPITestBase.hpp" +#include "ResourceGuards.hpp" +#include "TestChecks.hpp" +#include "algorithms/dda/dda_init_detail.h" +#include "comm.h" +#include "rccl_common.h" + +#include +#include + +#include +#include + +using namespace RCCLTestGuards; +using namespace RCCLTestHelpers; + +namespace +{ +constexpr int kDdaIpcRanks = nccl_dda_detail::kDdaNranks; + +bool isDdaIpcArch() +{ + hipDeviceProp_t props{}; + if(hipGetDeviceProperties(&props, 0) != hipSuccess) + { + return false; + } + const std::string arch(props.gcnArchName); + return arch.find("gfx942") != std::string::npos + || arch.find("gfx950") != std::string::npos; +} + +// Unique per rank and exact in fp32 (max 4095 * 8 + 7 < 2^24). +float rankPattern(int rank, size_t i) +{ + return static_cast((i % 4096) * 8 + static_cast(rank)); +} +} // namespace + +class DdaMPI_AllGatherThreshold : public MPITestBase +{ +protected: + int selectedAlgo(const void* sendBuf, void* recvBuf, size_t count) + { + int algo = -1; + int protocol = -1; + int maxChannels = -1; + EXPECT_EQ(ncclSuccess, + rcclGetCollImplInfo(getActiveCommunicator(), + ncclFuncAllGather, count, ncclFloat32, + ncclSum, sendBuf, recvBuf, + /*graphCapturing=*/0, &algo, &protocol, + &maxChannels)); + return algo; + } +}; + +// A total above scratch but under the threshold must fall back, not fail. +TEST_F(DdaMPI_AllGatherThreshold, AboveScratchFallsBack) +{ + const size_t threshold = 4 * static_cast(DDA_IPC_BUFFER_SIZE); + // RCCL caches the param on first read, so set it before comm init. + MPIHelpers::MpiEnvGuard thresholdGuard( + "RCCL_DDA_THRESHOLD", std::to_string(threshold).c_str()); + if(!validateTestPrerequisites(kDdaIpcRanks, kDdaIpcRanks, false, 1, 1)) + { + GTEST_SKIP() << "DDA IPC AllGather needs 8 ranks on one node"; + } + const bool optedOut = rcclParamDdaEnable() == 0; + std::string skip = mpiCoordinatedSkipReason( + !isDdaIpcArch() || optedOut, + optedOut ? "RCCL_DDA_ENABLE=0" : "DDA IPC needs gfx942 or gfx950"); + if(!skip.empty()) + { + GTEST_SKIP() << skip; + } + skip = mpiCoordinatedSkipReason( + static_cast(rcclParamDdaThreshold()) != threshold, + "RCCL_DDA_THRESHOLD was read earlier; run alone or set it at launch"); + if(!skip.empty()) + { + GTEST_SKIP() << skip; + } + + ASSERT_MPI_EQ(ncclSuccess, createTestCommunicator()); + ncclComm* comm = getActiveCommunicator(); + int rank = -1; + int nRanks = 0; + ASSERT_MPI_EQ(ncclSuccess, ncclCommUserRank(comm, &rank)); + ASSERT_MPI_EQ(ncclSuccess, ncclCommCount(comm, &nRanks)); + skip = mpiCoordinatedSkipReason(comm->config.launchOrderImplicit == 1, + "NCCL_LAUNCH_ORDER_IMPLICIT=1"); + if(!skip.empty()) + { + GTEST_SKIP() << skip; + } + + // 2x the per-rank share that fills scratch: 64 MiB each, 512 MiB total. + const size_t ranks = static_cast(nRanks); + const size_t fullCount = comm->ddaScratchBytes / (ranks * sizeof(float)); + const size_t count = 2 * fullCount; + const size_t totalCount = count * ranks; + ASSERT_MPI_TRUE(count * sizeof(float) <= comm->ddaScratchBytes); + ASSERT_MPI_TRUE(totalCount * sizeof(float) > comm->ddaScratchBytes); + ASSERT_MPI_TRUE(totalCount * sizeof(float) <= threshold); + + void* sendBuf = nullptr; + void* recvBuf = nullptr; + ASSERT_MPI_EQ(hipSuccess, hipMalloc(&sendBuf, count * sizeof(float))); + DeviceBufferAutoGuard sendGuard(sendBuf); + ASSERT_MPI_EQ(hipSuccess, hipMalloc(&recvBuf, totalCount * sizeof(float))); + DeviceBufferAutoGuard recvGuard(recvBuf); + + // Control: a total that exactly fills scratch still takes DDA IPC. + ASSERT_MPI_EQ(RCCL_DDA_IPC, selectedAlgo(sendBuf, recvBuf, fullCount)); + EXPECT_NE(RCCL_DDA_IPC, selectedAlgo(sendBuf, recvBuf, count)) + << "Rank " << rank << ": DDA IPC picked for a total above scratch"; + + ASSERT_MPI_EQ(hipSuccess, + initializeBufferWithPattern( + sendBuf, count, + [rank](size_t i) { return rankPattern(rank, i); })); + ASSERT_MPI_EQ(hipSuccess, + hipMemset(recvBuf, 0, totalCount * sizeof(float))); + ASSERT_MPI_EQ(ncclSuccess, + ncclAllGather(sendBuf, recvBuf, count, ncclFloat32, comm, + getActiveStream())); + ASSERT_MPI_EQ(hipSuccess, hipStreamSynchronize(getActiveStream())); + ASSERT_MPI_TRUE(verifyBufferData( + recvBuf, totalCount, [count](size_t j) { + return rankPattern(static_cast(j / count), j % count); + })); +} + +#endif // MPI_TESTS_ENABLED diff --git a/projects/rccl/test/DdaFabricEligibilityTests.cpp b/projects/rccl/test/DdaFabricEligibilityTests.cpp index 1dc48b878ffc..3e68983976b4 100644 --- a/projects/rccl/test/DdaFabricEligibilityTests.cpp +++ b/projects/rccl/test/DdaFabricEligibilityTests.cpp @@ -379,6 +379,18 @@ TEST_F(DdaFabricEligibilityTest, AllGather_ScratchTooSmall) mockComm_.get(), sendbuff_, recvbuff_, 4, ncclFloat32)); } +// The launcher checks nRanks * sendcount against scratch, so gate on the total. +TEST_F(DdaFabricEligibilityTest, AllGather_PerRankFitsButTotalExceedsScratch) +{ + ASSERT_EQ(mockComm_.comm.nRanks, 8); + mockComm_.comm.ddaScratchBytes = 1024; + // 8 ranks: 128 B each totals 1024 B and fits; 256 B each does not. + EXPECT_TRUE(ncclAllGatherDdaFabricEligible( + mockComm_.get(), sendbuff_, recvbuff_, 32, ncclFloat32)); + EXPECT_FALSE(ncclAllGatherDdaFabricEligible( + mockComm_.get(), sendbuff_, recvbuff_, 64, ncclFloat32)); +} + TEST_F(DdaFabricEligibilityTest, ReduceScatter_DerivedScratchAcceptsSmallMessage) { // Use derived scratch sizing with default params (8 ranks, 128 MiB threshold, LL+LL128 enabled). diff --git a/projects/rccl/test/DdaIpcEligibilityTests.cpp b/projects/rccl/test/DdaIpcEligibilityTests.cpp index 62887464e7eb..19aaf065d52f 100644 --- a/projects/rccl/test/DdaIpcEligibilityTests.cpp +++ b/projects/rccl/test/DdaIpcEligibilityTests.cpp @@ -77,6 +77,25 @@ TEST_F(DdaIpcEligibilityTest, AllGather_ScratchTooSmall) mockComm_.get(), sendbuff_, recvbuff_, 4, ncclFloat32)); } +// The launcher checks nRanks * sendcount against scratch, so gate on the total. +TEST_F(DdaIpcEligibilityTest, AllGather_PerRankFitsButTotalExceedsScratch) +{ + const size_t nRanks = static_cast(mockComm_.comm.nRanks); + const size_t maxSendCount = DDA_IPC_BUFFER_SIZE / (nRanks * sizeof(float)); + ASSERT_LE((maxSendCount + 4) * sizeof(float), mockComm_.comm.ddaScratchBytes); + EXPECT_FALSE(ncclAllGatherDdaIpcEligible( + mockComm_.get(), sendbuff_, recvbuff_, maxSendCount + 4, ncclFloat32)); +} + +TEST_F(DdaIpcEligibilityTest, AllGather_TotalAtScratchEligible) +{ + const size_t nRanks = static_cast(mockComm_.comm.nRanks); + const size_t maxSendCount = DDA_IPC_BUFFER_SIZE / (nRanks * sizeof(float)); + ASSERT_EQ(maxSendCount * nRanks * sizeof(float), mockComm_.comm.ddaScratchBytes); + EXPECT_TRUE(ncclAllGatherDdaIpcEligible( + mockComm_.get(), sendbuff_, recvbuff_, maxSendCount, ncclFloat32)); +} + TEST_F(DdaIpcEligibilityTest, AllGather_UnalignedCount) { EXPECT_FALSE(ncclAllGatherDdaIpcEligible( 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 4ef992f5b54f..0b97c71d6ae3 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 @@ -3092,6 +3092,14 @@ "env_variables": { "NCCL_GRAPH_STREAM_ORDERING": "0" } + }, + { + "name": "DDA_AllGatherThresholdAboveScratch_8Ranks", + "description": "RCCL_DDA_THRESHOLD above the IPC scratch: an AllGather whose total exceeds scratch falls back and succeeds instead of returning ncclInvalidArgument", + "test_filter": "DdaMPI_AllGatherThreshold.AboveScratchFallsBack", + "env_variables": { + "RCCL_DDA_THRESHOLD": "1073741824" + } } ] }, diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json index d2e1eaf6223d..8f548fcfb749 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json @@ -3277,6 +3277,14 @@ "env_variables": { "NCCL_GRAPH_STREAM_ORDERING": "0" } + }, + { + "name": "DDA_AllGatherThresholdAboveScratch_8Ranks", + "description": "RCCL_DDA_THRESHOLD above the IPC scratch: an AllGather whose total exceeds scratch falls back and succeeds instead of returning ncclInvalidArgument", + "test_filter": "DdaMPI_AllGatherThreshold.AboveScratchFallsBack", + "env_variables": { + "RCCL_DDA_THRESHOLD": "1073741824" + } } ] }, From 360d2639db93ec339bc277095bad40579c0d9a22 Mon Sep 17 00:00:00 2001 From: Praneeth Vallem Date: Thu, 8 Oct 2026 20:16:02 -0700 Subject: [PATCH 2/7] test(rccl): pin DDA IPC AllGather slice placement at 8 ranks The DDA IPC AllGather kernel (gfx942/gfx950, one process per GPU, 8 ranks) has no test that pins the backend and checks where each rank's slice lands. Coverage today is incidental: CeMPI_AllGather.EightRanks reaches it only when CE and symmetric kernels are off, and it does not assert which backend ran. Add DdaMPI_AllGather.SlicePlacement{OutOfPlace,InPlace}. Each case asserts rcclGetCollImplInfo reports DDA-IPC, fills rank r with r * 2^20 + i, pre-fills recv with a sentinel, and checks every slot at 1024 floats per rank (one block) and 8 MiB total (the gfx942 cap, full grid-stride loop). Failures are summed across ranks so rank 0 reports a peer's mismatch. Off-target arch, rank count, RCCL_DDA_ENABLE=0 and NCCL_LAUNCH_ORDER_IMPLICIT=1 skip on every rank; a missing DDA IPC init on gfx942/gfx950 at 8 ranks fails instead. Register the suite in test_categories_mpi.yaml and in the 8-rank runner configs for MI300X and MI355X. JIRA ID: AICOMRCCL-2772 --- projects/rccl/test/CMakeLists.txt | 1 + projects/rccl/test/DdaAllGatherMPITests.cpp | 215 ++++++++++++++++++ projects/rccl/test/test_categories_mpi.yaml | 1 + .../configs/mi300x_mellanox_ib.json | 5 + .../configs/mi355x_thor2_roce.json | 5 + 5 files changed, 227 insertions(+) create mode 100644 projects/rccl/test/DdaAllGatherMPITests.cpp diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index acac7f959602..916c7793af7a 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -1023,6 +1023,7 @@ if(BUILD_TESTS) DeviceApiMPITests.cpp CommNetMismatchMPITests.cpp test_sparse_send_recv.cpp + DdaAllGatherMPITests.cpp DdaAllGatherThresholdMPITests.cpp DdaAllReduceMPITests.cpp DdaFabricSimpleMPITests.cpp diff --git a/projects/rccl/test/DdaAllGatherMPITests.cpp b/projects/rccl/test/DdaAllGatherMPITests.cpp new file mode 100644 index 000000000000..9ef8213fd0d5 --- /dev/null +++ b/projects/rccl/test/DdaAllGatherMPITests.cpp @@ -0,0 +1,215 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +/** + * @file DdaAllGatherMPITests.cpp + * @brief Slice placement for the DDA IPC AllGather (gfx942/gfx950, 8 ranks). + * + * Run (example): + * mpirun -np 8 ./rccl-UnitTestsMPI --gtest_filter=DdaMPI_AllGather.* + */ + +#ifdef MPI_TESTS_ENABLED + +#include "DeviceBufferHelpers.hpp" +#include "MPIHelpers.hpp" +#include "MPITestBase.hpp" +#include "ResourceGuards.hpp" +#include "TestChecks.hpp" +#include "algorithms/dda/dda_init_detail.h" +#include "comm.h" +#include "rccl_common.h" + +#include +#include + +#include +#include +#include + +using namespace RCCLTestGuards; +using namespace RCCLTestHelpers; + +namespace +{ +constexpr int kDdaIpcRanks = nccl_dda_detail::kDdaNranks; + +// Keeps f(rank, i) unique per slot and exact in fp32 (7 * 2^20 + 2^18 < 2^24). +constexpr size_t kRankStride = size_t{1} << 20; + +// One block per rank, then 8 MiB total: the gfx942 cap and a full grid-stride. +constexpr size_t kSmallCount = 1024; +constexpr size_t kGridCount = 256 * 1024; + +// No rank sends a negative value, so an unwritten slot cannot verify. +constexpr float kUnwritten = -1.0f; + +float slicePattern(int rank, size_t i) +{ + return static_cast(static_cast(rank) * kRankStride + i); +} + +bool isDdaIpcArch() +{ + hipDeviceProp_t props{}; + if(hipGetDeviceProperties(&props, 0) != hipSuccess) + { + return false; + } + const std::string arch(props.gcnArchName); + return arch.find("gfx942") != std::string::npos + || arch.find("gfx950") != std::string::npos; +} + +// ASSERT_MPI_* skips healthy ranks, so rank 0 can hide a peer's failure. +int countFailingRanks(bool localOk) +{ + int failing = localOk ? 0 : 1; + MPI_Allreduce(MPI_IN_PLACE, &failing, 1, MPI_INT, MPI_SUM, MPI_COMM_WORLD); + return failing; +} +} // namespace + +class DdaMPI_AllGather : public MPITestBase +{ +protected: + int rank_{}; + int nRanks_{}; + + void SetUp() override + { + MPITestBase::SetUp(); + if(!validateTestPrerequisites(kDdaIpcRanks, kDdaIpcRanks, false, 1, 1)) + { + GTEST_SKIP() << "DDA IPC AllGather needs 8 ranks on one node"; + } + const bool optedOut = rcclParamDdaEnable() == 0; + std::string skip = mpiCoordinatedSkipReason( + !isDdaIpcArch() || optedOut, + optedOut ? "RCCL_DDA_ENABLE=0" : "DDA IPC needs gfx942 or gfx950"); + if(!skip.empty()) + { + GTEST_SKIP() << skip; + } + + const ncclResult_t res = createTestCommunicator(); + ASSERT_EQ(0, countFailingRanks(res == ncclSuccess)) + << "Rank " << getTestMpiRank() << ": " << ncclGetErrorString(res); + ncclComm_t comm = getActiveCommunicator(); + ncclCommUserRank(comm, &rank_); + ncclCommCount(comm, &nRanks_); + + // Implicit launch order disables DDA for the process lifetime. + skip = mpiCoordinatedSkipReason(comm->config.launchOrderImplicit == 1, + "NCCL_LAUNCH_ORDER_IMPLICIT=1"); + if(!skip.empty()) + { + GTEST_SKIP() << skip; + } + ASSERT_EQ(0, countFailingRanks(comm->ddaIpcBarrierState != nullptr)) + << "Rank " << rank_ << ": DDA IPC did not initialize"; + } + + void runPlacement(size_t count, bool inPlace) + { + const size_t totalCount = count * static_cast(nRanks_); + + void* recvBuf = nullptr; + ASSERT_EQ(0, countFailingRanks( + hipMalloc(&recvBuf, totalCount * sizeof(float)) + == hipSuccess)); + DeviceBufferAutoGuard recvGuard(recvBuf); + ASSERT_EQ(0, countFailingRanks( + initializeBufferWithPattern( + recvBuf, totalCount, + [](size_t) { return kUnwritten; }) + == hipSuccess)); + + void* sendBuf + = static_cast(recvBuf) + static_cast(rank_) * count; + void* separateSendBuf = nullptr; + if(!inPlace) + { + ASSERT_EQ(0, countFailingRanks( + hipMalloc(&separateSendBuf, count * sizeof(float)) + == hipSuccess)); + sendBuf = separateSendBuf; + } + DeviceBufferAutoGuard sendGuard(separateSendBuf); + ASSERT_EQ(0, countFailingRanks( + initializeBufferWithPattern( + sendBuf, count, + [this](size_t i) + { + return slicePattern(rank_, i); + }) + == hipSuccess)); + + // A kernel-path fallback also gathers correctly, so pin the backend. + int algo = -1, protocol = -1, maxChannels = -1; + const ncclResult_t queryRes = rcclGetCollImplInfo( + getActiveCommunicator(), ncclFuncAllGather, count, ncclFloat32, + ncclSum, sendBuf, recvBuf, /*graphCapturing=*/0, &algo, &protocol, + &maxChannels); + const char* algoName = nullptr; + rcclGetAlgoName(algo, &algoName); + ASSERT_EQ(0, countFailingRanks(queryRes == ncclSuccess + && algo == RCCL_DDA_IPC)) + << "Rank " << rank_ << ": " << count << " floats selected " + << (algoName ? algoName : "?") << ", not DDA-IPC"; + + const ncclResult_t agRes = ncclAllGather( + sendBuf, recvBuf, count, ncclFloat32, getActiveCommunicator(), + getActiveStream()); + const hipError_t syncErr = hipStreamSynchronize(getActiveStream()); + ASSERT_EQ(0, countFailingRanks(agRes == ncclSuccess + && syncErr == hipSuccess)) + << "Rank " << rank_ << ": " << ncclGetErrorString(agRes) << " / " + << hipGetErrorString(syncErr); + + auto [copyErr, host] = downloadBuffer(recvBuf, totalCount); + size_t mismatches = copyErr == hipSuccess ? 0 : totalCount; + size_t firstBad = totalCount; + for(size_t j = 0; copyErr == hipSuccess && j < totalCount; ++j) + { + if(host[j] != slicePattern(static_cast(j / count), j % count)) + { + firstBad = mismatches == 0 ? j : firstBad; + ++mismatches; + } + } + ASSERT_EQ(0, countFailingRanks(mismatches == 0)) + << "Rank " << rank_ << ": " << mismatches << " of " << totalCount + << " wrong, first at " << firstBad << " (count " << count << ")"; + } +}; + +// recvbuff[r * count + i] must hold f(r, i) on every rank. +TEST_F(DdaMPI_AllGather, SlicePlacementOutOfPlace) +{ + for(size_t count : {kSmallCount, kGridCount}) + { + runPlacement(count, /*inPlace=*/false); + if(HasFailure()) + { + return; + } + } +} + +TEST_F(DdaMPI_AllGather, SlicePlacementInPlace) +{ + for(size_t count : {kSmallCount, kGridCount}) + { + runPlacement(count, /*inPlace=*/true); + if(HasFailure()) + { + return; + } + } +} + +#endif // MPI_TESTS_ENABLED diff --git a/projects/rccl/test/test_categories_mpi.yaml b/projects/rccl/test/test_categories_mpi.yaml index ee1f8e68c607..4f4853162f9e 100644 --- a/projects/rccl/test/test_categories_mpi.yaml +++ b/projects/rccl/test/test_categories_mpi.yaml @@ -238,6 +238,7 @@ test_categories: - "DdaFabricSimple.*" - "DdaFabricBlockCap.*" - "DdaFabricBarrierStress.*" + - "DdaMPI_AllGather.*" exclude_windows: *mpi_exclude_windows labels: - "rccl" 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 0b97c71d6ae3..3aa6c9a1eb39 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 @@ -3093,6 +3093,11 @@ "NCCL_GRAPH_STREAM_ORDERING": "0" } }, + { + "name": "DDA_AllGatherSlicePlacement_8Ranks", + "description": "DDA IPC AllGather writes each rank's slice at the sender's offset on every rank, out of place and in place, at one-block and grid-stride sizes", + "test_filter": "DdaMPI_AllGather.SlicePlacement*" + }, { "name": "DDA_AllGatherThresholdAboveScratch_8Ranks", "description": "RCCL_DDA_THRESHOLD above the IPC scratch: an AllGather whose total exceeds scratch falls back and succeeds instead of returning ncclInvalidArgument", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json index 8f548fcfb749..ce92246ea731 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json @@ -3278,6 +3278,11 @@ "NCCL_GRAPH_STREAM_ORDERING": "0" } }, + { + "name": "DDA_AllGatherSlicePlacement_8Ranks", + "description": "DDA IPC AllGather writes each rank's slice at the sender's offset on every rank, out of place and in place, at one-block and grid-stride sizes", + "test_filter": "DdaMPI_AllGather.SlicePlacement*" + }, { "name": "DDA_AllGatherThresholdAboveScratch_8Ranks", "description": "RCCL_DDA_THRESHOLD above the IPC scratch: an AllGather whose total exceeds scratch falls back and succeeds instead of returning ncclInvalidArgument", From 1dbc938335e54dc9ddc7a8848dea24c2bf69a477 Mon Sep 17 00:00:00 2001 From: Praneeth Vallem Date: Thu, 8 Oct 2026 13:22:25 -0700 Subject: [PATCH 3/7] fix(rccl): keep DDA grid/block math in size_t past 64 GiB getGridAndBlockDims narrowed ceil(count / elementsPerThread) to uint32, so at 64 GiB per-rank payload it produced blockDim.x == 0 or a shrunken grid. Compute in size_t and narrow only after clamping to maxBlocks and 512 threads. JIRA ID: AICOMRCCL-2787 --- .../src/algorithms/dda/device/CollCommon.h | 10 +++++----- projects/rccl/test/DdaCollCommonTests.cpp | 20 +++++++++++++++++++ 2 files changed, 25 insertions(+), 5 deletions(-) diff --git a/projects/rccl/src/algorithms/dda/device/CollCommon.h b/projects/rccl/src/algorithms/dda/device/CollCommon.h index cbf0c9426126..d0795a8a9af7 100644 --- a/projects/rccl/src/algorithms/dda/device/CollCommon.h +++ b/projects/rccl/src/algorithms/dda/device/CollCommon.h @@ -162,8 +162,9 @@ static inline __device__ void allGather(T* const* __restrict__ ipcbuffs, T* __re } } -inline uint32_t divRoundUp(size_t a, size_t b) { - uint32_t y = static_cast((a + b - 1) / b); +// Returns size_t: at >= 64 GiB the thread and warp counts exceed uint32, so narrow only after clamping. +inline size_t divRoundUp(size_t a, size_t b) { + size_t y = (a + b - 1) / b; if (y == 0) { y = 1; } @@ -202,14 +203,13 @@ inline std::pair getGridAndBlockDims(size_t count, int typeSize, siz dim3 threads(0, 1, 1); dim3 blocks(0, 1, 1); if (ddaAlltoAllSingleBlockGrid(count, typeSize)) { - threads.x = divRoundUp(count, elementsPerWarp) * kThreadsPerWarp; + threads.x = static_cast(divRoundUp(count, elementsPerWarp) * kThreadsPerWarp); blocks.x = 1; } else { auto warpsRequired = divRoundUp(count, elementsPerWarp); blocks.x = calcBlockCount(divRoundUp(count, elementsPerThread), kThreadsPerBlock, maxBlocks); auto warpsPerBlock = divRoundUp(warpsRequired, blocks.x); - auto threadsPerBlock = std::min(kThreadsPerBlock, warpsPerBlock * kThreadsPerWarp); - threads.x = threadsPerBlock; + threads.x = static_cast(std::min(kThreadsPerBlock, warpsPerBlock * kThreadsPerWarp)); } return std::make_pair(blocks, threads); diff --git a/projects/rccl/test/DdaCollCommonTests.cpp b/projects/rccl/test/DdaCollCommonTests.cpp index da009c529e15..0154af178ca2 100644 --- a/projects/rccl/test/DdaCollCommonTests.cpp +++ b/projects/rccl/test/DdaCollCommonTests.cpp @@ -82,4 +82,24 @@ TEST(DdaCollCommon, getGridAndBlockDims_HalfPrecision) EXPECT_LE(gridBlock.second.x, 512u); } +// 64 GiB is the first size where count / elementsPerThread reaches 2^32; fabric scratch can be env-sized past it. +TEST(DdaCollCommon, getGridAndBlockDims_PastUint32ThreadCount) +{ + constexpr size_t kBytes = size_t{1} << 36; + for (int typeSize : {1, 2, 4}) + { + for (size_t maxBlocks : {size_t{1}, size_t{24}, size_t{256}}) + { + for (size_t count : {kBytes / typeSize - 1, kBytes / typeSize, kBytes / typeSize + 4096}) + { + SCOPED_TRACE(::testing::Message() << "typeSize=" << typeSize << " maxBlocks=" << maxBlocks + << " count=" << count); + auto gridBlock = dda::common::getGridAndBlockDims(count, typeSize, maxBlocks); + EXPECT_EQ(gridBlock.first.x, maxBlocks); + EXPECT_EQ(gridBlock.second.x, 512u); + } + } + } +} + } // namespace RcclUnitTesting From 5a23722dde6ab21dc012a1145156abc7e1ff778d Mon Sep 17 00:00:00 2001 From: Praneeth Vallem Date: Sat, 10 Oct 2026 14:26:54 -0700 Subject: [PATCH 4/7] test(rccl): share one fixture file for the DDA IPC AllGather MPI tests The slice placement and threshold fallback suites each carried their own copy of the includes, the gfx942/gfx950 check and the DDA skip logic. Move the threshold suite into DdaAllGatherMPITests.cpp and give both suites a common base fixture for the coordinated skips and the algo query. Suite names and runner rows are unchanged. JIRA ID: AICOMRCCL-2775, AICOMRCCL-2772 --- projects/rccl/test/CMakeLists.txt | 1 - projects/rccl/test/DdaAllGatherMPITests.cpp | 145 ++++++++++++++--- .../test/DdaAllGatherThresholdMPITests.cpp | 154 ------------------ 3 files changed, 126 insertions(+), 174 deletions(-) delete mode 100644 projects/rccl/test/DdaAllGatherThresholdMPITests.cpp diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index 916c7793af7a..6fff07365d9a 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -1024,7 +1024,6 @@ if(BUILD_TESTS) CommNetMismatchMPITests.cpp test_sparse_send_recv.cpp DdaAllGatherMPITests.cpp - DdaAllGatherThresholdMPITests.cpp DdaAllReduceMPITests.cpp DdaFabricSimpleMPITests.cpp SymmetricAbortCheckModeMPITests.cpp diff --git a/projects/rccl/test/DdaAllGatherMPITests.cpp b/projects/rccl/test/DdaAllGatherMPITests.cpp index 9ef8213fd0d5..fe7bcbb109ce 100644 --- a/projects/rccl/test/DdaAllGatherMPITests.cpp +++ b/projects/rccl/test/DdaAllGatherMPITests.cpp @@ -6,10 +6,13 @@ /** * @file DdaAllGatherMPITests.cpp - * @brief Slice placement for the DDA IPC AllGather (gfx942/gfx950, 8 ranks). + * @brief DDA IPC AllGather (gfx942/gfx950, 8 ranks): slice placement, and + * fallback when RCCL_DDA_THRESHOLD is above the IPC scratch. * * Run (example): * mpirun -np 8 ./rccl-UnitTestsMPI --gtest_filter=DdaMPI_AllGather.* + * mpirun -np 8 -x RCCL_DDA_THRESHOLD=1073741824 ./rccl-UnitTestsMPI + * --gtest_filter=DdaMPI_AllGatherThreshold.* */ #ifdef MPI_TESTS_ENABLED @@ -52,6 +55,12 @@ float slicePattern(int rank, size_t i) return static_cast(static_cast(rank) * kRankStride + i); } +// Unique per rank and exact in fp32 at any count (max 4095 * 8 + 7 < 2^24). +float rankPattern(int rank, size_t i) +{ + return static_cast((i % 4096) * 8 + static_cast(rank)); +} + bool isDdaIpcArch() { hipDeviceProp_t props{}; @@ -73,23 +82,55 @@ int countFailingRanks(bool localOk) } } // namespace -class DdaMPI_AllGather : public MPITestBase +class DdaIpcAllGatherMPITest : public MPITestBase { protected: - int rank_{}; - int nRanks_{}; - - void SetUp() override + // Empty when this job can reach DDA IPC; every rank gets the same answer. Call before comm init. + std::string preInitSkipReason() { - MPITestBase::SetUp(); if(!validateTestPrerequisites(kDdaIpcRanks, kDdaIpcRanks, false, 1, 1)) { - GTEST_SKIP() << "DDA IPC AllGather needs 8 ranks on one node"; + return "DDA IPC AllGather needs 8 ranks on one node"; } const bool optedOut = rcclParamDdaEnable() == 0; - std::string skip = mpiCoordinatedSkipReason( + return mpiCoordinatedSkipReason( !isDdaIpcArch() || optedOut, optedOut ? "RCCL_DDA_ENABLE=0" : "DDA IPC needs gfx942 or gfx950"); + } + + // Implicit launch order disables DDA for the process lifetime. + std::string launchOrderSkipReason() + { + return mpiCoordinatedSkipReason( + getActiveCommunicator()->config.launchOrderImplicit == 1, + "NCCL_LAUNCH_ORDER_IMPLICIT=1"); + } + + int selectedAlgo(const void* sendBuf, void* recvBuf, size_t count) + { + int algo = -1; + int protocol = -1; + int maxChannels = -1; + EXPECT_EQ(ncclSuccess, + rcclGetCollImplInfo(getActiveCommunicator(), + ncclFuncAllGather, count, ncclFloat32, + ncclSum, sendBuf, recvBuf, + /*graphCapturing=*/0, &algo, &protocol, + &maxChannels)); + return algo; + } +}; + +class DdaMPI_AllGather : public DdaIpcAllGatherMPITest +{ +protected: + int rank_{}; + int nRanks_{}; + + void SetUp() override + { + MPITestBase::SetUp(); + std::string skip = preInitSkipReason(); if(!skip.empty()) { GTEST_SKIP() << skip; @@ -102,9 +143,7 @@ class DdaMPI_AllGather : public MPITestBase ncclCommUserRank(comm, &rank_); ncclCommCount(comm, &nRanks_); - // Implicit launch order disables DDA for the process lifetime. - skip = mpiCoordinatedSkipReason(comm->config.launchOrderImplicit == 1, - "NCCL_LAUNCH_ORDER_IMPLICIT=1"); + skip = launchOrderSkipReason(); if(!skip.empty()) { GTEST_SKIP() << skip; @@ -149,15 +188,10 @@ class DdaMPI_AllGather : public MPITestBase == hipSuccess)); // A kernel-path fallback also gathers correctly, so pin the backend. - int algo = -1, protocol = -1, maxChannels = -1; - const ncclResult_t queryRes = rcclGetCollImplInfo( - getActiveCommunicator(), ncclFuncAllGather, count, ncclFloat32, - ncclSum, sendBuf, recvBuf, /*graphCapturing=*/0, &algo, &protocol, - &maxChannels); + const int algo = selectedAlgo(sendBuf, recvBuf, count); const char* algoName = nullptr; rcclGetAlgoName(algo, &algoName); - ASSERT_EQ(0, countFailingRanks(queryRes == ncclSuccess - && algo == RCCL_DDA_IPC)) + ASSERT_EQ(0, countFailingRanks(algo == RCCL_DDA_IPC)) << "Rank " << rank_ << ": " << count << " floats selected " << (algoName ? algoName : "?") << ", not DDA-IPC"; @@ -187,6 +221,10 @@ class DdaMPI_AllGather : public MPITestBase } }; +class DdaMPI_AllGatherThreshold : public DdaIpcAllGatherMPITest +{ +}; + // recvbuff[r * count + i] must hold f(r, i) on every rank. TEST_F(DdaMPI_AllGather, SlicePlacementOutOfPlace) { @@ -212,4 +250,73 @@ TEST_F(DdaMPI_AllGather, SlicePlacementInPlace) } } +// A total above scratch but under the threshold must fall back, not fail. +TEST_F(DdaMPI_AllGatherThreshold, AboveScratchFallsBack) +{ + const size_t threshold = 4 * static_cast(DDA_IPC_BUFFER_SIZE); + // RCCL caches the param on first read, so set it before comm init. + MPIHelpers::MpiEnvGuard thresholdGuard( + "RCCL_DDA_THRESHOLD", std::to_string(threshold).c_str()); + std::string skip = preInitSkipReason(); + if(!skip.empty()) + { + GTEST_SKIP() << skip; + } + skip = mpiCoordinatedSkipReason( + static_cast(rcclParamDdaThreshold()) != threshold, + "RCCL_DDA_THRESHOLD was read earlier; run alone or set it at launch"); + if(!skip.empty()) + { + GTEST_SKIP() << skip; + } + + ASSERT_MPI_EQ(ncclSuccess, createTestCommunicator()); + ncclComm* comm = getActiveCommunicator(); + int rank = -1; + int nRanks = 0; + ASSERT_MPI_EQ(ncclSuccess, ncclCommUserRank(comm, &rank)); + ASSERT_MPI_EQ(ncclSuccess, ncclCommCount(comm, &nRanks)); + skip = launchOrderSkipReason(); + if(!skip.empty()) + { + GTEST_SKIP() << skip; + } + + // 2x the per-rank share that fills scratch: 64 MiB each, 512 MiB total. + const size_t ranks = static_cast(nRanks); + const size_t fullCount = comm->ddaScratchBytes / (ranks * sizeof(float)); + const size_t count = 2 * fullCount; + const size_t totalCount = count * ranks; + ASSERT_MPI_TRUE(count * sizeof(float) <= comm->ddaScratchBytes); + ASSERT_MPI_TRUE(totalCount * sizeof(float) > comm->ddaScratchBytes); + ASSERT_MPI_TRUE(totalCount * sizeof(float) <= threshold); + + void* sendBuf = nullptr; + void* recvBuf = nullptr; + ASSERT_MPI_EQ(hipSuccess, hipMalloc(&sendBuf, count * sizeof(float))); + DeviceBufferAutoGuard sendGuard(sendBuf); + ASSERT_MPI_EQ(hipSuccess, hipMalloc(&recvBuf, totalCount * sizeof(float))); + DeviceBufferAutoGuard recvGuard(recvBuf); + + // Control: a total that exactly fills scratch still takes DDA IPC. + ASSERT_MPI_EQ(RCCL_DDA_IPC, selectedAlgo(sendBuf, recvBuf, fullCount)); + EXPECT_NE(RCCL_DDA_IPC, selectedAlgo(sendBuf, recvBuf, count)) + << "Rank " << rank << ": DDA IPC picked for a total above scratch"; + + ASSERT_MPI_EQ(hipSuccess, + initializeBufferWithPattern( + sendBuf, count, + [rank](size_t i) { return rankPattern(rank, i); })); + ASSERT_MPI_EQ(hipSuccess, + hipMemset(recvBuf, 0, totalCount * sizeof(float))); + ASSERT_MPI_EQ(ncclSuccess, + ncclAllGather(sendBuf, recvBuf, count, ncclFloat32, comm, + getActiveStream())); + ASSERT_MPI_EQ(hipSuccess, hipStreamSynchronize(getActiveStream())); + ASSERT_MPI_TRUE(verifyBufferData( + recvBuf, totalCount, [count](size_t j) { + return rankPattern(static_cast(j / count), j % count); + })); +} + #endif // MPI_TESTS_ENABLED diff --git a/projects/rccl/test/DdaAllGatherThresholdMPITests.cpp b/projects/rccl/test/DdaAllGatherThresholdMPITests.cpp deleted file mode 100644 index 3be2744b3e2f..000000000000 --- a/projects/rccl/test/DdaAllGatherThresholdMPITests.cpp +++ /dev/null @@ -1,154 +0,0 @@ -/************************************************************************* - * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ - -/** - * @file DdaAllGatherThresholdMPITests.cpp - * @brief DDA IPC AllGather with RCCL_DDA_THRESHOLD above the IPC scratch. - * - * Run (example): - * mpirun -np 8 -x RCCL_DDA_THRESHOLD=1073741824 ./rccl-UnitTestsMPI - * --gtest_filter=DdaMPI_AllGatherThreshold.* - */ - -#ifdef MPI_TESTS_ENABLED - -#include "DeviceBufferHelpers.hpp" -#include "MPIHelpers.hpp" -#include "MPITestBase.hpp" -#include "ResourceGuards.hpp" -#include "TestChecks.hpp" -#include "algorithms/dda/dda_init_detail.h" -#include "comm.h" -#include "rccl_common.h" - -#include -#include - -#include -#include - -using namespace RCCLTestGuards; -using namespace RCCLTestHelpers; - -namespace -{ -constexpr int kDdaIpcRanks = nccl_dda_detail::kDdaNranks; - -bool isDdaIpcArch() -{ - hipDeviceProp_t props{}; - if(hipGetDeviceProperties(&props, 0) != hipSuccess) - { - return false; - } - const std::string arch(props.gcnArchName); - return arch.find("gfx942") != std::string::npos - || arch.find("gfx950") != std::string::npos; -} - -// Unique per rank and exact in fp32 (max 4095 * 8 + 7 < 2^24). -float rankPattern(int rank, size_t i) -{ - return static_cast((i % 4096) * 8 + static_cast(rank)); -} -} // namespace - -class DdaMPI_AllGatherThreshold : public MPITestBase -{ -protected: - int selectedAlgo(const void* sendBuf, void* recvBuf, size_t count) - { - int algo = -1; - int protocol = -1; - int maxChannels = -1; - EXPECT_EQ(ncclSuccess, - rcclGetCollImplInfo(getActiveCommunicator(), - ncclFuncAllGather, count, ncclFloat32, - ncclSum, sendBuf, recvBuf, - /*graphCapturing=*/0, &algo, &protocol, - &maxChannels)); - return algo; - } -}; - -// A total above scratch but under the threshold must fall back, not fail. -TEST_F(DdaMPI_AllGatherThreshold, AboveScratchFallsBack) -{ - const size_t threshold = 4 * static_cast(DDA_IPC_BUFFER_SIZE); - // RCCL caches the param on first read, so set it before comm init. - MPIHelpers::MpiEnvGuard thresholdGuard( - "RCCL_DDA_THRESHOLD", std::to_string(threshold).c_str()); - if(!validateTestPrerequisites(kDdaIpcRanks, kDdaIpcRanks, false, 1, 1)) - { - GTEST_SKIP() << "DDA IPC AllGather needs 8 ranks on one node"; - } - const bool optedOut = rcclParamDdaEnable() == 0; - std::string skip = mpiCoordinatedSkipReason( - !isDdaIpcArch() || optedOut, - optedOut ? "RCCL_DDA_ENABLE=0" : "DDA IPC needs gfx942 or gfx950"); - if(!skip.empty()) - { - GTEST_SKIP() << skip; - } - skip = mpiCoordinatedSkipReason( - static_cast(rcclParamDdaThreshold()) != threshold, - "RCCL_DDA_THRESHOLD was read earlier; run alone or set it at launch"); - if(!skip.empty()) - { - GTEST_SKIP() << skip; - } - - ASSERT_MPI_EQ(ncclSuccess, createTestCommunicator()); - ncclComm* comm = getActiveCommunicator(); - int rank = -1; - int nRanks = 0; - ASSERT_MPI_EQ(ncclSuccess, ncclCommUserRank(comm, &rank)); - ASSERT_MPI_EQ(ncclSuccess, ncclCommCount(comm, &nRanks)); - skip = mpiCoordinatedSkipReason(comm->config.launchOrderImplicit == 1, - "NCCL_LAUNCH_ORDER_IMPLICIT=1"); - if(!skip.empty()) - { - GTEST_SKIP() << skip; - } - - // 2x the per-rank share that fills scratch: 64 MiB each, 512 MiB total. - const size_t ranks = static_cast(nRanks); - const size_t fullCount = comm->ddaScratchBytes / (ranks * sizeof(float)); - const size_t count = 2 * fullCount; - const size_t totalCount = count * ranks; - ASSERT_MPI_TRUE(count * sizeof(float) <= comm->ddaScratchBytes); - ASSERT_MPI_TRUE(totalCount * sizeof(float) > comm->ddaScratchBytes); - ASSERT_MPI_TRUE(totalCount * sizeof(float) <= threshold); - - void* sendBuf = nullptr; - void* recvBuf = nullptr; - ASSERT_MPI_EQ(hipSuccess, hipMalloc(&sendBuf, count * sizeof(float))); - DeviceBufferAutoGuard sendGuard(sendBuf); - ASSERT_MPI_EQ(hipSuccess, hipMalloc(&recvBuf, totalCount * sizeof(float))); - DeviceBufferAutoGuard recvGuard(recvBuf); - - // Control: a total that exactly fills scratch still takes DDA IPC. - ASSERT_MPI_EQ(RCCL_DDA_IPC, selectedAlgo(sendBuf, recvBuf, fullCount)); - EXPECT_NE(RCCL_DDA_IPC, selectedAlgo(sendBuf, recvBuf, count)) - << "Rank " << rank << ": DDA IPC picked for a total above scratch"; - - ASSERT_MPI_EQ(hipSuccess, - initializeBufferWithPattern( - sendBuf, count, - [rank](size_t i) { return rankPattern(rank, i); })); - ASSERT_MPI_EQ(hipSuccess, - hipMemset(recvBuf, 0, totalCount * sizeof(float))); - ASSERT_MPI_EQ(ncclSuccess, - ncclAllGather(sendBuf, recvBuf, count, ncclFloat32, comm, - getActiveStream())); - ASSERT_MPI_EQ(hipSuccess, hipStreamSynchronize(getActiveStream())); - ASSERT_MPI_TRUE(verifyBufferData( - recvBuf, totalCount, [count](size_t j) { - return rankPattern(static_cast(j / count), j % count); - })); -} - -#endif // MPI_TESTS_ENABLED From ee66cd4ab3d34c14c0c6f8277f1505748d51d7bb Mon Sep 17 00:00:00 2001 From: Praneeth Vallem Date: Mon, 12 Oct 2026 04:15:48 -0700 Subject: [PATCH 5/7] test(rccl): report DDA AllGather threshold failures from every rank AboveScratchFallsBack used ASSERT_MPI_*, which skips the passing ranks, so rank 0 could report a skip while a peer failed. Route its checks through countFailingRanks like the placement suite, guard each buffer before the shared allocation check so no rank leaks on a peer's failure, and fill the receive buffer with a negative sentinel so rank 0's zero values cannot pass unwritten. The threshold suite had no ctest category: DdaMPI_AllGather.* does not match it, and RCCL_DDA_THRESHOLD is cached on first read. Add a quick_env category that sets it at launch, as quick_env_pat_shared_qps_off does. JIRA ID: AICOMRCCL-2775, AICOMRCCL-2772 --- projects/rccl/test/DdaAllGatherMPITests.cpp | 87 ++++++++++++--------- projects/rccl/test/test_categories_mpi.yaml | 10 +++ 2 files changed, 62 insertions(+), 35 deletions(-) diff --git a/projects/rccl/test/DdaAllGatherMPITests.cpp b/projects/rccl/test/DdaAllGatherMPITests.cpp index fe7bcbb109ce..7b0f0853d806 100644 --- a/projects/rccl/test/DdaAllGatherMPITests.cpp +++ b/projects/rccl/test/DdaAllGatherMPITests.cpp @@ -169,15 +169,16 @@ class DdaMPI_AllGather : public DdaIpcAllGatherMPITest void* sendBuf = static_cast(recvBuf) + static_cast(rank_) * count; - void* separateSendBuf = nullptr; + void* separateSendBuf = nullptr; + const hipError_t sendAllocErr + = inPlace ? hipSuccess + : hipMalloc(&separateSendBuf, count * sizeof(float)); + DeviceBufferAutoGuard sendGuard(separateSendBuf); + ASSERT_EQ(0, countFailingRanks(sendAllocErr == hipSuccess)); if(!inPlace) { - ASSERT_EQ(0, countFailingRanks( - hipMalloc(&separateSendBuf, count * sizeof(float)) - == hipSuccess)); sendBuf = separateSendBuf; } - DeviceBufferAutoGuard sendGuard(separateSendBuf); ASSERT_EQ(0, countFailingRanks( initializeBufferWithPattern( sendBuf, count, @@ -270,12 +271,14 @@ TEST_F(DdaMPI_AllGatherThreshold, AboveScratchFallsBack) GTEST_SKIP() << skip; } - ASSERT_MPI_EQ(ncclSuccess, createTestCommunicator()); - ncclComm* comm = getActiveCommunicator(); - int rank = -1; - int nRanks = 0; - ASSERT_MPI_EQ(ncclSuccess, ncclCommUserRank(comm, &rank)); - ASSERT_MPI_EQ(ncclSuccess, ncclCommCount(comm, &nRanks)); + const ncclResult_t initRes = createTestCommunicator(); + ASSERT_EQ(0, countFailingRanks(initRes == ncclSuccess)) + << "Rank " << getTestMpiRank() << ": " << ncclGetErrorString(initRes); + ncclComm* comm = getActiveCommunicator(); + int rank = -1; + int nRanks = 0; + ASSERT_EQ(0, countFailingRanks(ncclCommUserRank(comm, &rank) == ncclSuccess + && ncclCommCount(comm, &nRanks) == ncclSuccess)); skip = launchOrderSkipReason(); if(!skip.empty()) { @@ -287,36 +290,50 @@ TEST_F(DdaMPI_AllGatherThreshold, AboveScratchFallsBack) const size_t fullCount = comm->ddaScratchBytes / (ranks * sizeof(float)); const size_t count = 2 * fullCount; const size_t totalCount = count * ranks; - ASSERT_MPI_TRUE(count * sizeof(float) <= comm->ddaScratchBytes); - ASSERT_MPI_TRUE(totalCount * sizeof(float) > comm->ddaScratchBytes); - ASSERT_MPI_TRUE(totalCount * sizeof(float) <= threshold); - - void* sendBuf = nullptr; - void* recvBuf = nullptr; - ASSERT_MPI_EQ(hipSuccess, hipMalloc(&sendBuf, count * sizeof(float))); + const size_t totalBytes = totalCount * sizeof(float); + ASSERT_EQ(0, countFailingRanks(count * sizeof(float) <= comm->ddaScratchBytes + && totalBytes > comm->ddaScratchBytes + && totalBytes <= threshold)) + << "Rank " << rank << ": scratch " << comm->ddaScratchBytes << " B, total " + << totalBytes << " B, threshold " << threshold << " B"; + + void* sendBuf = nullptr; + void* recvBuf = nullptr; + const hipError_t sendErr = hipMalloc(&sendBuf, count * sizeof(float)); DeviceBufferAutoGuard sendGuard(sendBuf); - ASSERT_MPI_EQ(hipSuccess, hipMalloc(&recvBuf, totalCount * sizeof(float))); + const hipError_t recvErr = hipMalloc(&recvBuf, totalBytes); DeviceBufferAutoGuard recvGuard(recvBuf); + ASSERT_EQ(0, countFailingRanks(sendErr == hipSuccess && recvErr == hipSuccess)); // Control: a total that exactly fills scratch still takes DDA IPC. - ASSERT_MPI_EQ(RCCL_DDA_IPC, selectedAlgo(sendBuf, recvBuf, fullCount)); - EXPECT_NE(RCCL_DDA_IPC, selectedAlgo(sendBuf, recvBuf, count)) + ASSERT_EQ(0, countFailingRanks(selectedAlgo(sendBuf, recvBuf, fullCount) + == RCCL_DDA_IPC)) + << "Rank " << rank << ": DDA IPC not picked for a total that fills scratch"; + EXPECT_EQ(0, countFailingRanks(selectedAlgo(sendBuf, recvBuf, count) + != RCCL_DDA_IPC)) << "Rank " << rank << ": DDA IPC picked for a total above scratch"; - ASSERT_MPI_EQ(hipSuccess, - initializeBufferWithPattern( - sendBuf, count, - [rank](size_t i) { return rankPattern(rank, i); })); - ASSERT_MPI_EQ(hipSuccess, - hipMemset(recvBuf, 0, totalCount * sizeof(float))); - ASSERT_MPI_EQ(ncclSuccess, - ncclAllGather(sendBuf, recvBuf, count, ncclFloat32, comm, - getActiveStream())); - ASSERT_MPI_EQ(hipSuccess, hipStreamSynchronize(getActiveStream())); - ASSERT_MPI_TRUE(verifyBufferData( - recvBuf, totalCount, [count](size_t j) { - return rankPattern(static_cast(j / count), j % count); - })); + const hipError_t sendInitErr = initializeBufferWithPattern( + sendBuf, count, [rank](size_t i) { return rankPattern(rank, i); }); + const hipError_t recvInitErr = initializeBufferWithPattern( + recvBuf, totalCount, [](size_t) { return kUnwritten; }); + ASSERT_EQ(0, countFailingRanks(sendInitErr == hipSuccess + && recvInitErr == hipSuccess)); + + const ncclResult_t agRes + = ncclAllGather(sendBuf, recvBuf, count, ncclFloat32, comm, getActiveStream()); + const hipError_t syncErr = hipStreamSynchronize(getActiveStream()); + ASSERT_EQ(0, countFailingRanks(agRes == ncclSuccess && syncErr == hipSuccess)) + << "Rank " << rank << ": " << ncclGetErrorString(agRes) << " / " + << hipGetErrorString(syncErr); + + size_t firstBad = totalCount; + const bool dataOk = verifyBufferData( + recvBuf, totalCount, + [count](size_t j) { return rankPattern(static_cast(j / count), j % count); }, + /*num_samples=*/0, /*tolerance=*/0.0, &firstBad); + ASSERT_EQ(0, countFailingRanks(dataOk)) + << "Rank " << rank << ": first mismatch at " << firstBad << " of " << totalCount; } #endif // MPI_TESTS_ENABLED diff --git a/projects/rccl/test/test_categories_mpi.yaml b/projects/rccl/test/test_categories_mpi.yaml index 4f4853162f9e..1129c6ee16be 100644 --- a/projects/rccl/test/test_categories_mpi.yaml +++ b/projects/rccl/test/test_categories_mpi.yaml @@ -215,6 +215,15 @@ test_categories: - "NCCL_DEBUG=INFO" - "NCCL_LAUNCH_MODE=GROUP" - "RCCL_PAT_SHARED_QPS=0" + quick_env_dda_allgather_threshold: + description: "DDA IPC AllGather fallback above scratch. RCCL_PARAM caches RCCL_DDA_THRESHOLD for the process lifetime, so AboveScratchFallsBack only runs when it is set before launch." + test_patterns: + - "DdaMPI_AllGatherThreshold.*" + exclude_windows: *mpi_exclude_windows + labels: *mpi_quick_labels + env_variables: + - "NCCL_DEBUG=INFO" + - "RCCL_DDA_THRESHOLD=1073741824" standard: description: "MPI regression — all transport / IB / implicit-launch / UBR / graph-capture suites (same patterns as comprehensive; lighter env)." test_patterns: &mpi_all_patterns @@ -426,6 +435,7 @@ execution_settings: quick_env_netib_xfer_trace: 1200 quick_env_netib_xfer_log_ranks: 1200 quick_env_pat_shared_qps_off: 1200 + quick_env_dda_allgather_threshold: 1200 standard: 1200 comprehensive: 1200 comprehensive_disable_env: 1200 From 804c02f77876b8790afcf92177d123192cfbd0b3 Mon Sep 17 00:00:00 2001 From: Praneeth Vallem Date: Mon, 12 Oct 2026 04:15:48 -0700 Subject: [PATCH 6/7] docs(rccl): note the DDA AllGather fallback and 64 GiB grid fixes JIRA ID: AICOMRCCL-2788, AICOMRCCL-2775, AICOMRCCL-2787 --- projects/rccl/CHANGELOG.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 7c6660ec1a5c..9b7ca5ceba1f 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -38,6 +38,8 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it. ### Resolved issues +* Fixed `ncclAllGather` returning `ncclInvalidArgument` instead of falling back to another algorithm when `RCCL_DDA_THRESHOLD` is raised above the DDA scratch size and the gathered total does not fit the scratch. The DDA IPC and fabric AllGather eligibility checks now test the gathered total, as their launchers do. +* Fixed DDA kernels launching with `blockDim.x=0` for payloads of 64 GiB or more, which `RCCL_DDA_THRESHOLD` can make reachable on the DDA fabric path. * Fixed profiler overhead when no profiler plugin is loaded or only kernel-channel events are enabled. * Fixed profiler API events reporting rank 0 instead of the originating communicator rank. * Fixed P2P IPC registration reuse producing out-of-bounds remote addresses when a registered allocation spans multiple cuMem segments. From 1ebfe047ba85efdd03c7ceb161917a9c9c14ddc8 Mon Sep 17 00:00:00 2001 From: Praneeth Vallem Date: Fri, 16 Oct 2026 13:53:35 -0700 Subject: [PATCH 7/7] test(rccl): tighten DDA AllGather MPI test buffers and patterns Guard recvBuf before the cross-rank alloc check so a peer's failed hipMalloc no longer leaks it on the other ranks. Fold i / 4096 into rankPattern so a slice shifted by a multiple of 4096 floats no longer verifies clean. Share one file-local mismatch counter between runPlacement and AboveScratchFallsBack. --- projects/rccl/test/DdaAllGatherMPITests.cpp | 62 ++++++++++++++------- 1 file changed, 41 insertions(+), 21 deletions(-) diff --git a/projects/rccl/test/DdaAllGatherMPITests.cpp b/projects/rccl/test/DdaAllGatherMPITests.cpp index 7b0f0853d806..e340226bc83b 100644 --- a/projects/rccl/test/DdaAllGatherMPITests.cpp +++ b/projects/rccl/test/DdaAllGatherMPITests.cpp @@ -55,10 +55,10 @@ float slicePattern(int rank, size_t i) return static_cast(static_cast(rank) * kRankStride + i); } -// Unique per rank and exact in fp32 at any count (max 4095 * 8 + 7 < 2^24). +// Unique per rank and exact in fp32 (max 4095 * 8 + 7 < 2^24); i / 4096 catches a slice shifted by 4096 floats. float rankPattern(int rank, size_t i) { - return static_cast((i % 4096) * 8 + static_cast(rank)); + return static_cast(((i % 4096 + i / 4096) % 4096) * 8 + static_cast(rank)); } bool isDdaIpcArch() @@ -80,6 +80,28 @@ int countFailingRanks(bool localOk) MPI_Allreduce(MPI_IN_PLACE, &failing, 1, MPI_INT, MPI_SUM, MPI_COMM_WORLD); return failing; } + +// Mismatches in recvBuf against expected(j); all count as wrong if the copy fails. +template +size_t countMismatches(const void* recvBuf, size_t totalCount, Expected expected, size_t* firstBad) +{ + auto [copyErr, host] = downloadBuffer(recvBuf, totalCount); + *firstBad = totalCount; + if(copyErr != hipSuccess) + { + return totalCount; + } + size_t mismatches = 0; + for(size_t j = 0; j < totalCount; ++j) + { + if(host[j] != expected(j)) + { + *firstBad = mismatches == 0 ? j : *firstBad; + ++mismatches; + } + } + return mismatches; +} } // namespace class DdaIpcAllGatherMPITest : public MPITestBase @@ -156,11 +178,11 @@ class DdaMPI_AllGather : public DdaIpcAllGatherMPITest { const size_t totalCount = count * static_cast(nRanks_); - void* recvBuf = nullptr; - ASSERT_EQ(0, countFailingRanks( - hipMalloc(&recvBuf, totalCount * sizeof(float)) - == hipSuccess)); + void* recvBuf = nullptr; + const hipError_t recvAllocErr + = hipMalloc(&recvBuf, totalCount * sizeof(float)); DeviceBufferAutoGuard recvGuard(recvBuf); + ASSERT_EQ(0, countFailingRanks(recvAllocErr == hipSuccess)); ASSERT_EQ(0, countFailingRanks( initializeBufferWithPattern( recvBuf, totalCount, @@ -205,17 +227,14 @@ class DdaMPI_AllGather : public DdaIpcAllGatherMPITest << "Rank " << rank_ << ": " << ncclGetErrorString(agRes) << " / " << hipGetErrorString(syncErr); - auto [copyErr, host] = downloadBuffer(recvBuf, totalCount); - size_t mismatches = copyErr == hipSuccess ? 0 : totalCount; - size_t firstBad = totalCount; - for(size_t j = 0; copyErr == hipSuccess && j < totalCount; ++j) - { - if(host[j] != slicePattern(static_cast(j / count), j % count)) + size_t firstBad = totalCount; + const size_t mismatches = countMismatches( + recvBuf, totalCount, + [count](size_t j) { - firstBad = mismatches == 0 ? j : firstBad; - ++mismatches; - } - } + return slicePattern(static_cast(j / count), j % count); + }, + &firstBad); ASSERT_EQ(0, countFailingRanks(mismatches == 0)) << "Rank " << rank_ << ": " << mismatches << " of " << totalCount << " wrong, first at " << firstBad << " (count " << count << ")"; @@ -327,13 +346,14 @@ TEST_F(DdaMPI_AllGatherThreshold, AboveScratchFallsBack) << "Rank " << rank << ": " << ncclGetErrorString(agRes) << " / " << hipGetErrorString(syncErr); - size_t firstBad = totalCount; - const bool dataOk = verifyBufferData( + size_t firstBad = totalCount; + const size_t mismatches = countMismatches( recvBuf, totalCount, [count](size_t j) { return rankPattern(static_cast(j / count), j % count); }, - /*num_samples=*/0, /*tolerance=*/0.0, &firstBad); - ASSERT_EQ(0, countFailingRanks(dataOk)) - << "Rank " << rank << ": first mismatch at " << firstBad << " of " << totalCount; + &firstBad); + ASSERT_EQ(0, countFailingRanks(mismatches == 0)) + << "Rank " << rank << ": " << mismatches << " of " << totalCount + << " wrong, first at " << firstBad; } #endif // MPI_TESTS_ENABLED