diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 669b91f8a19b..fc1810158aaf 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -54,6 +54,8 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: ### Resolved issues * Fixed topology GPU indices that followed the HIP device order instead of the PCI order. This reordered rings and NIC selection, causing an `ncclGather` performance regression on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU. +* 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. 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/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/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index 5118040dd3b5..dd6bf1617505 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -1041,6 +1041,7 @@ if(BUILD_TESTS) DeviceApiMPITests.cpp CommNetMismatchMPITests.cpp test_sparse_send_recv.cpp + DdaAllGatherMPITests.cpp DdaAllReduceMPITests.cpp DdaFabricSimpleMPITests.cpp SymmetricAbortCheckModeMPITests.cpp diff --git a/projects/rccl/test/DdaAllGatherMPITests.cpp b/projects/rccl/test/DdaAllGatherMPITests.cpp new file mode 100644 index 000000000000..e340226bc83b --- /dev/null +++ b/projects/rccl/test/DdaAllGatherMPITests.cpp @@ -0,0 +1,359 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +/** + * @file DdaAllGatherMPITests.cpp + * @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 + +#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); +} + +// 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 + i / 4096) % 4096) * 8 + static_cast(rank)); +} + +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; +} + +// 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 +{ +protected: + // Empty when this job can reach DDA IPC; every rank gets the same answer. Call before comm init. + std::string preInitSkipReason() + { + if(!validateTestPrerequisites(kDdaIpcRanks, kDdaIpcRanks, false, 1, 1)) + { + return "DDA IPC AllGather needs 8 ranks on one node"; + } + const bool optedOut = rcclParamDdaEnable() == 0; + 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; + } + + 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_); + + skip = launchOrderSkipReason(); + 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; + 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, + [](size_t) { return kUnwritten; }) + == hipSuccess)); + + void* sendBuf + = static_cast(recvBuf) + static_cast(rank_) * count; + 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) + { + sendBuf = 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. + const int algo = selectedAlgo(sendBuf, recvBuf, count); + const char* algoName = nullptr; + rcclGetAlgoName(algo, &algoName); + ASSERT_EQ(0, countFailingRanks(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); + + size_t firstBad = totalCount; + const size_t mismatches = countMismatches( + recvBuf, totalCount, + [count](size_t j) + { + 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 << ")"; + } +}; + +class DdaMPI_AllGatherThreshold : public DdaIpcAllGatherMPITest +{ +}; + +// 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; + } + } +} + +// 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; + } + + 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()) + { + 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; + 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); + 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_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"; + + 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 size_t mismatches = countMismatches( + recvBuf, totalCount, + [count](size_t j) { return rankPattern(static_cast(j / count), j % count); }, + &firstBad); + ASSERT_EQ(0, countFailingRanks(mismatches == 0)) + << "Rank " << rank << ": " << mismatches << " of " << totalCount + << " wrong, first at " << firstBad; +} + +#endif // MPI_TESTS_ENABLED 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 diff --git a/projects/rccl/test/DdaFabricEligibilityTests.cpp b/projects/rccl/test/DdaFabricEligibilityTests.cpp index 2a8215c65b2a..77e97fba11e1 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/test/test_categories_mpi.yaml b/projects/rccl/test/test_categories_mpi.yaml index c696cda978aa..3fb7d8d90311 100644 --- a/projects/rccl/test/test_categories_mpi.yaml +++ b/projects/rccl/test/test_categories_mpi.yaml @@ -218,6 +218,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 @@ -242,6 +251,7 @@ test_categories: - "DdaFabricSimple.*" - "DdaFabricBlockCap.*" - "DdaFabricBarrierStress.*" + - "DdaMPI_AllGather.*" exclude_windows: *mpi_exclude_windows labels: - "rccl" @@ -429,6 +439,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 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 fab6109311d6..09e6d70aa1ec 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 @@ -3162,6 +3162,19 @@ "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", + "test_filter": "DdaMPI_AllGatherThreshold.AboveScratchFallsBack", + "env_variables": { + "RCCL_DDA_THRESHOLD": "1073741824" + } + }, { "name": "HostApi_DenseAllToAllPutSignal_8Ranks", "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", 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 d6c347dfd306..8949533f8d04 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 @@ -3333,6 +3333,19 @@ "env_variables": { "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", + "test_filter": "DdaMPI_AllGatherThreshold.AboveScratchFallsBack", + "env_variables": { + "RCCL_DDA_THRESHOLD": "1073741824" + } } ] },