From 3fdc8b03cb6f695714343c122f357268f29910a5 Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Fri, 11 Sep 2026 11:40:12 -0600 Subject: [PATCH 01/10] Correct the P2P LL threshold to maintain existing SIMPLE/LL crossover tuning on AMD systems. --- projects/rccl/src/enqueue/enqueue.cc | 17 ++++++++++++++++- 1 file changed, 16 insertions(+), 1 deletion(-) diff --git a/projects/rccl/src/enqueue/enqueue.cc b/projects/rccl/src/enqueue/enqueue.cc index 679a42e266e..7026e1c321f 100644 --- a/projects/rccl/src/enqueue/enqueue.cc +++ b/projects/rccl/src/enqueue/enqueue.cc @@ -1522,7 +1522,22 @@ static ncclResult_t addP2pToPlan(struct ncclComm* comm, struct ncclKernelPlan* p } else { bool lat = useLL128OptIn ? hasLL128[dir] : hasLL[dir]; ssize_t latencyThreshold = useLL128OptIn ? ncclParamP2pLL128Threshold() : ncclParamP2pLLThreshold(); - if (!(useLL128OptIn && latencyThreshold == 0)) lat &= bytes[dir] <= nChannels[dir] * latencyThreshold; + ssize_t latencyBytesMax = nChannels[dir] * latencyThreshold; + // nChannels here is the SIMPLE estimate, which collapses to 1 for latency-bound sizes and so + // hides the latency protocol on topologies that expose many P2P channels per peer. Single-node + // all-XGMI gfx942/gfx950 exposes 8, where the SIMPLE estimate costs ~11% bus bandwidth on + // 8-rank alltoall between 128 KiB and 512 KiB total (16-64 KiB per peer). Widen eligibility to + // the smallest SIMPLE part there, capped at the measured 64 KiB/peer crossover so 128 KiB/peer + // and above keep SIMPLE (going past it regresses 1-4 MiB total). This covers every P2P op -- + // send/recv, alltoall, gather, scatter and the p2p-based all-gather/reduce-scatter -- since + // they all size their transfers here. A user-set NCCL_P2P_LL_THRESHOLD keeps the stock rule, as + // does the LL128 P2P path, whose crossover has not been measured on single node. + static bool const userSetP2pLLThreshold = ncclGetEnv("NCCL_P2P_LL_THRESHOLD") != nullptr; + if (comm->nNodes <= 1 && comm->isAllNvlink && !useLL128OptIn && + (comm->cudaArch == 940 || comm->cudaArch == 950) && !userSetP2pLLThreshold) { + latencyBytesMax = std::max(latencyBytesMax, std::min(comm->p2pChunkSize / 8, 65536)); + } + if (!(useLL128OptIn && latencyThreshold == 0)) lat &= bytes[dir] <= latencyBytesMax; protocol[dir] = lat ? (useLL128OptIn ? NCCL_PROTO_LL128 : NCCL_PROTO_LL) : NCCL_PROTO_SIMPLE; } } From b92da4042d72367c47b528a33da15243ed787a5e Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Mon, 26 Oct 2026 15:54:46 -0600 Subject: [PATCH 02/10] docs(rccl): shorten the single-node P2P LL threshold comment The selection rule is unchanged. State the gfx942/gfx950 cap and the exceptions for a set NCCL_P2P_LL_THRESHOLD and the LL128 path. --- projects/rccl/src/enqueue/enqueue.cc | 12 +++--------- 1 file changed, 3 insertions(+), 9 deletions(-) diff --git a/projects/rccl/src/enqueue/enqueue.cc b/projects/rccl/src/enqueue/enqueue.cc index 7026e1c321f..6406f47e1f6 100644 --- a/projects/rccl/src/enqueue/enqueue.cc +++ b/projects/rccl/src/enqueue/enqueue.cc @@ -1523,15 +1523,9 @@ static ncclResult_t addP2pToPlan(struct ncclComm* comm, struct ncclKernelPlan* p bool lat = useLL128OptIn ? hasLL128[dir] : hasLL[dir]; ssize_t latencyThreshold = useLL128OptIn ? ncclParamP2pLL128Threshold() : ncclParamP2pLLThreshold(); ssize_t latencyBytesMax = nChannels[dir] * latencyThreshold; - // nChannels here is the SIMPLE estimate, which collapses to 1 for latency-bound sizes and so - // hides the latency protocol on topologies that expose many P2P channels per peer. Single-node - // all-XGMI gfx942/gfx950 exposes 8, where the SIMPLE estimate costs ~11% bus bandwidth on - // 8-rank alltoall between 128 KiB and 512 KiB total (16-64 KiB per peer). Widen eligibility to - // the smallest SIMPLE part there, capped at the measured 64 KiB/peer crossover so 128 KiB/peer - // and above keep SIMPLE (going past it regresses 1-4 MiB total). This covers every P2P op -- - // send/recv, alltoall, gather, scatter and the p2p-based all-gather/reduce-scatter -- since - // they all size their transfers here. A user-set NCCL_P2P_LL_THRESHOLD keeps the stock rule, as - // does the LL128 P2P path, whose crossover has not been measured on single node. + // SIMPLE's channel count is 1 here, so the per-channel LL threshold never selects LL on + // single-node gfx942/gfx950. Allow LL up to p2pChunkSize/8, capped at 64 KiB/peer. A set + // NCCL_P2P_LL_THRESHOLD, and the LL128 path, keep the stock rule. static bool const userSetP2pLLThreshold = ncclGetEnv("NCCL_P2P_LL_THRESHOLD") != nullptr; if (comm->nNodes <= 1 && comm->isAllNvlink && !useLL128OptIn && (comm->cudaArch == 940 || comm->cudaArch == 950) && !userSetP2pLLThreshold) { From 3947d1785db1b2e2ecb0c9d8ecce31c796841f81 Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Tue, 27 Oct 2026 09:14:32 -0600 Subject: [PATCH 03/10] Added tests. --- projects/rccl/test/SendRecvTests.cpp | 164 +++++++++++++++++++++++++++ 1 file changed, 164 insertions(+) diff --git a/projects/rccl/test/SendRecvTests.cpp b/projects/rccl/test/SendRecvTests.cpp index 1988240215c..4f42e83d68d 100644 --- a/projects/rccl/test/SendRecvTests.cpp +++ b/projects/rccl/test/SendRecvTests.cpp @@ -4,6 +4,7 @@ * See LICENSE.txt for license information ************************************************************************/ #include "TestBed.hpp" +#include "common/ProcessIsolatedTestRunner.hpp" #include #include @@ -87,6 +88,36 @@ namespace RcclUnitTesting return found; } + // True if any P2P protocol line for this byte count reports protocol. + // "LL " does not match "LL128". A self-copy in the same alltoall stays SIMPLE, so a + // size can legitimately log both protocols; callers decide which combination they need. + static bool DebugLogsHaveProtocolForBytes(const std::string& globPattern, const char* protocol, long bytes) + { + const std::string needle = std::string("RCCL P2P SendRecv protocol=") + protocol + + " dir="; + const std::string bytesNeedle = "bytes=" + std::to_string(bytes) + " "; + glob_t g{}; + bool found = false; + if (glob(globPattern.c_str(), 0, nullptr, &g) == 0) + { + for (size_t i = 0; i < g.gl_pathc && !found; ++i) + { + std::ifstream f(g.gl_pathv[i]); + std::string line; + while (std::getline(f, line)) + { + if (line.find(needle) != std::string::npos && line.find(bytesNeedle) != std::string::npos) + { + found = true; + break; + } + } + } + } + globfree(&g); + return found; + } + static bool DebugLogsContainNeedle(const std::string& globPattern, const char* needle) { glob_t g{}; @@ -1108,4 +1139,137 @@ namespace RcclUnitTesting EXPECT_TRUE(RunGfx1250FourRankRing(testBed, sendBytes, recvBytes)); } } + + // Per-peer alltoall sizes whose P2P protocol the single-node gfx94/gfx95 LL cap must produce. + // A self-copy logs SIMPLE at every size, so "LL" means the cross-peer op selected LL, and + // "Simple" means no cross-peer op selected LL. + struct P2pLlCrossoverExpect { + long bytes; + bool expectLL; + }; + + // Two-rank alltoall on the legacy-LL path. The SIMPLE channel estimate is 1 at and below + // 64 KiB/peer, which is what made NCCL 2.29's per-channel threshold skip LL. Each case runs + // in its own process (see the tests below) because the eligibility check latches + // NCCL_P2P_LL_THRESHOLD on first use. + static void RunSingleNodeGfx94LlCrossover(P2pLlCrossoverExpect const* expect, int nExpect) + { + std::string arch = DeviceGcnArchName(); + if (arch.find("gfx94") == std::string::npos && arch.find("gfx95") == std::string::npos) { + GTEST_SKIP() << "Skipping... single-node gfx942/gfx950 LL crossover is gfx94/gfx95 only (detected " + << arch << ")."; + } + // Inherited pins change the SIMPLE estimate or the 64 KiB cap. + unsetenv("NCCL_P2P_NVL_CHUNKSIZE"); + unsetenv("NCCL_MAX_P2P_NCHANNELS"); + unsetenv("NCCL_MIN_P2P_NCHANNELS"); + unsetenv("NCCL_PROTO"); + + std::string const debugGlob = "/tmp/rccl_p2p_ll_crossover_" + std::to_string(getpid()) + ".*"; + RemoveGlobbedFiles(debugGlob); + setenv("NCCL_DEBUG", "INFO", 1); + setenv("NCCL_DEBUG_SUBSYS", "COLL", 1); + setenv("NCCL_DEBUG_FILE", ("/tmp/rccl_p2p_ll_crossover_" + std::to_string(getpid()) + ".%p").c_str(), 1); + { + TestBed testBed; + if (testBed.ev.maxGpus < 2) { + testBed.Finalize(); + RemoveGlobbedFiles(debugGlob); + unsetenv("NCCL_DEBUG_FILE"); + unsetenv("NCCL_DEBUG_SUBSYS"); + unsetenv("NCCL_DEBUG"); + GTEST_SKIP() << "Skipping... P2P LL crossover needs at least 2 GPUs (detected " + << testBed.ev.maxGpus << ")."; + } + std::vector dataTypes; + testBed.GetSupportedDataTypes(dataTypes, {ncclInt8}); + if (dataTypes.empty()) { + testBed.Finalize(); + RemoveGlobbedFiles(debugGlob); + unsetenv("NCCL_DEBUG_FILE"); + unsetenv("NCCL_DEBUG_SUBSYS"); + unsetenv("NCCL_DEBUG"); + GTEST_SKIP() << "Skipping... test datatypes excluded by UT_DATATYPES."; + } + + // Two ranks: the SIMPLE estimate is already 1 at these sizes, and gfx94 DDA alltoall + // requires 8 ranks, so this stays on the P2P planner the fix changes. + int const nRanks = 2; + testBed.InitComms(nRanks); + if (::testing::Test::HasFatalFailure()) { + testBed.Finalize(); + return; + } + bool isCorrect = true; + for (int i = 0; i < nExpect && isCorrect && !::testing::Test::HasFailure(); ++i) { + size_t const total = static_cast(expect[i].bytes) * static_cast(nRanks); + testBed.SetCollectiveArgs(ncclCollAlltoAll, ncclInt8, total, total); + testBed.AllocateMem(false, false); + testBed.PrepareData(); + testBed.ExecuteCollectives(); + testBed.ValidateResults(isCorrect); + testBed.DeallocateMem(); + } + testBed.DestroyComms(); + testBed.Finalize(); + EXPECT_TRUE(isCorrect); + } + if (::testing::Test::HasFailure()) { + RemoveGlobbedFiles(debugGlob); + unsetenv("NCCL_DEBUG_FILE"); + unsetenv("NCCL_DEBUG_SUBSYS"); + unsetenv("NCCL_DEBUG"); + return; + } + for (int i = 0; i < nExpect; ++i) { + bool const sawLL = DebugLogsHaveProtocolForBytes(debugGlob, "LL", expect[i].bytes); + bool const sawSimple = DebugLogsHaveProtocolForBytes(debugGlob, "Simple", expect[i].bytes); + if (expect[i].expectLL) { + EXPECT_TRUE(sawLL) << "per-peer bytes " << expect[i].bytes + << " expected LL; the SIMPLE channel estimate must not hide it"; + } else { + EXPECT_TRUE(sawSimple) << "per-peer bytes " << expect[i].bytes << " expected SIMPLE"; + EXPECT_FALSE(sawLL) << "per-peer bytes " << expect[i].bytes + << " must stay SIMPLE (the 64 KiB/peer cap)"; + } + } + RemoveGlobbedFiles(debugGlob); + unsetenv("NCCL_DEBUG_FILE"); + unsetenv("NCCL_DEBUG_SUBSYS"); + unsetenv("NCCL_DEBUG"); + } + + // Default threshold. 16/32/64 KiB per peer must be LL; 64 KiB+1 and 128 KiB-512 KiB must stay + // SIMPLE. Isolated so a prior test's NCCL_P2P_LL_THRESHOLD cannot latch the check. + TEST(SendRecv, SingleNodeGfx94LlCrossover) + { + RUN_ISOLATED_TEST_WITH_ENV( + "SingleNodeGfx94LlCrossover", + []() { + unsetenv("NCCL_P2P_LL_THRESHOLD"); + P2pLlCrossoverExpect const expect[] = { + {16384, true}, {32768, true}, {65536, true}, + {65537, false}, {131072, false}, {262144, false}, {524288, false}, + }; + RunSingleNodeGfx94LlCrossover(expect, sizeof(expect) / sizeof(expect[0])); + }, + {{"NCCL_ALLOC_P2P_NET_LL_BUFFERS", "0"}, + {"NCCL_P2P_LL128_ENABLE", "0"}, + {"RCCL_DDA_ENABLE", "0"}}); + } + + // An explicit NCCL_P2P_LL_THRESHOLD keeps the stock per-channel rule, so 16 KiB/peer stays SIMPLE. + TEST(SendRecv, SingleNodeGfx94LlCrossoverHonorsUserThreshold) + { + RUN_ISOLATED_TEST_WITH_ENV( + "SingleNodeGfx94LlCrossoverHonorsUserThreshold", + []() { + P2pLlCrossoverExpect const expect[] = {{16384, false}}; + RunSingleNodeGfx94LlCrossover(expect, 1); + }, + {{"NCCL_P2P_LL_THRESHOLD", "4096"}, + {"NCCL_ALLOC_P2P_NET_LL_BUFFERS", "0"}, + {"NCCL_P2P_LL128_ENABLE", "0"}, + {"RCCL_DDA_ENABLE", "0"}}); + } } From 79ea4ac3528cad346df3c2768427e1c0689c725c Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Tue, 27 Oct 2026 15:19:36 -0600 Subject: [PATCH 04/10] test(rccl): move gfx94 P2P LL crossover coverage to host enqueue microtests Replace the process-isolated SendRecvTests that latched NCCL_P2P_LL_THRESHOLD with focused EnqueueMicrotest cases around addP2pToPlan. This lets the test verify the widened LL cap (up to 64 KiB/peer) and the user-threshold fallback without relying on debug-log parsing or process isolation. Also make the user-threshold check read NCCL_P2P_LL_THRESHOLD on every call, so a later set or unset is visible. --- projects/rccl/src/enqueue/enqueue.cc | 5 +- projects/rccl/test/SendRecvTests.cpp | 164 ------------------------ projects/rccl/test/host/enqueue-test.cc | 102 ++++++++++++++- 3 files changed, 104 insertions(+), 167 deletions(-) diff --git a/projects/rccl/src/enqueue/enqueue.cc b/projects/rccl/src/enqueue/enqueue.cc index 6406f47e1f6..2d8090eb3bc 100644 --- a/projects/rccl/src/enqueue/enqueue.cc +++ b/projects/rccl/src/enqueue/enqueue.cc @@ -1525,8 +1525,9 @@ static ncclResult_t addP2pToPlan(struct ncclComm* comm, struct ncclKernelPlan* p ssize_t latencyBytesMax = nChannels[dir] * latencyThreshold; // SIMPLE's channel count is 1 here, so the per-channel LL threshold never selects LL on // single-node gfx942/gfx950. Allow LL up to p2pChunkSize/8, capped at 64 KiB/peer. A set - // NCCL_P2P_LL_THRESHOLD, and the LL128 path, keep the stock rule. - static bool const userSetP2pLLThreshold = ncclGetEnv("NCCL_P2P_LL_THRESHOLD") != nullptr; + // NCCL_P2P_LL_THRESHOLD, and the LL128 path, keep the stock rule. The env is read on + // each call so a later set or unset is visible. + bool const userSetP2pLLThreshold = ncclGetEnv("NCCL_P2P_LL_THRESHOLD") != nullptr; if (comm->nNodes <= 1 && comm->isAllNvlink && !useLL128OptIn && (comm->cudaArch == 940 || comm->cudaArch == 950) && !userSetP2pLLThreshold) { latencyBytesMax = std::max(latencyBytesMax, std::min(comm->p2pChunkSize / 8, 65536)); diff --git a/projects/rccl/test/SendRecvTests.cpp b/projects/rccl/test/SendRecvTests.cpp index 4f42e83d68d..1988240215c 100644 --- a/projects/rccl/test/SendRecvTests.cpp +++ b/projects/rccl/test/SendRecvTests.cpp @@ -4,7 +4,6 @@ * See LICENSE.txt for license information ************************************************************************/ #include "TestBed.hpp" -#include "common/ProcessIsolatedTestRunner.hpp" #include #include @@ -88,36 +87,6 @@ namespace RcclUnitTesting return found; } - // True if any P2P protocol line for this byte count reports protocol. - // "LL " does not match "LL128". A self-copy in the same alltoall stays SIMPLE, so a - // size can legitimately log both protocols; callers decide which combination they need. - static bool DebugLogsHaveProtocolForBytes(const std::string& globPattern, const char* protocol, long bytes) - { - const std::string needle = std::string("RCCL P2P SendRecv protocol=") + protocol + - " dir="; - const std::string bytesNeedle = "bytes=" + std::to_string(bytes) + " "; - glob_t g{}; - bool found = false; - if (glob(globPattern.c_str(), 0, nullptr, &g) == 0) - { - for (size_t i = 0; i < g.gl_pathc && !found; ++i) - { - std::ifstream f(g.gl_pathv[i]); - std::string line; - while (std::getline(f, line)) - { - if (line.find(needle) != std::string::npos && line.find(bytesNeedle) != std::string::npos) - { - found = true; - break; - } - } - } - } - globfree(&g); - return found; - } - static bool DebugLogsContainNeedle(const std::string& globPattern, const char* needle) { glob_t g{}; @@ -1139,137 +1108,4 @@ namespace RcclUnitTesting EXPECT_TRUE(RunGfx1250FourRankRing(testBed, sendBytes, recvBytes)); } } - - // Per-peer alltoall sizes whose P2P protocol the single-node gfx94/gfx95 LL cap must produce. - // A self-copy logs SIMPLE at every size, so "LL" means the cross-peer op selected LL, and - // "Simple" means no cross-peer op selected LL. - struct P2pLlCrossoverExpect { - long bytes; - bool expectLL; - }; - - // Two-rank alltoall on the legacy-LL path. The SIMPLE channel estimate is 1 at and below - // 64 KiB/peer, which is what made NCCL 2.29's per-channel threshold skip LL. Each case runs - // in its own process (see the tests below) because the eligibility check latches - // NCCL_P2P_LL_THRESHOLD on first use. - static void RunSingleNodeGfx94LlCrossover(P2pLlCrossoverExpect const* expect, int nExpect) - { - std::string arch = DeviceGcnArchName(); - if (arch.find("gfx94") == std::string::npos && arch.find("gfx95") == std::string::npos) { - GTEST_SKIP() << "Skipping... single-node gfx942/gfx950 LL crossover is gfx94/gfx95 only (detected " - << arch << ")."; - } - // Inherited pins change the SIMPLE estimate or the 64 KiB cap. - unsetenv("NCCL_P2P_NVL_CHUNKSIZE"); - unsetenv("NCCL_MAX_P2P_NCHANNELS"); - unsetenv("NCCL_MIN_P2P_NCHANNELS"); - unsetenv("NCCL_PROTO"); - - std::string const debugGlob = "/tmp/rccl_p2p_ll_crossover_" + std::to_string(getpid()) + ".*"; - RemoveGlobbedFiles(debugGlob); - setenv("NCCL_DEBUG", "INFO", 1); - setenv("NCCL_DEBUG_SUBSYS", "COLL", 1); - setenv("NCCL_DEBUG_FILE", ("/tmp/rccl_p2p_ll_crossover_" + std::to_string(getpid()) + ".%p").c_str(), 1); - { - TestBed testBed; - if (testBed.ev.maxGpus < 2) { - testBed.Finalize(); - RemoveGlobbedFiles(debugGlob); - unsetenv("NCCL_DEBUG_FILE"); - unsetenv("NCCL_DEBUG_SUBSYS"); - unsetenv("NCCL_DEBUG"); - GTEST_SKIP() << "Skipping... P2P LL crossover needs at least 2 GPUs (detected " - << testBed.ev.maxGpus << ")."; - } - std::vector dataTypes; - testBed.GetSupportedDataTypes(dataTypes, {ncclInt8}); - if (dataTypes.empty()) { - testBed.Finalize(); - RemoveGlobbedFiles(debugGlob); - unsetenv("NCCL_DEBUG_FILE"); - unsetenv("NCCL_DEBUG_SUBSYS"); - unsetenv("NCCL_DEBUG"); - GTEST_SKIP() << "Skipping... test datatypes excluded by UT_DATATYPES."; - } - - // Two ranks: the SIMPLE estimate is already 1 at these sizes, and gfx94 DDA alltoall - // requires 8 ranks, so this stays on the P2P planner the fix changes. - int const nRanks = 2; - testBed.InitComms(nRanks); - if (::testing::Test::HasFatalFailure()) { - testBed.Finalize(); - return; - } - bool isCorrect = true; - for (int i = 0; i < nExpect && isCorrect && !::testing::Test::HasFailure(); ++i) { - size_t const total = static_cast(expect[i].bytes) * static_cast(nRanks); - testBed.SetCollectiveArgs(ncclCollAlltoAll, ncclInt8, total, total); - testBed.AllocateMem(false, false); - testBed.PrepareData(); - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - testBed.DeallocateMem(); - } - testBed.DestroyComms(); - testBed.Finalize(); - EXPECT_TRUE(isCorrect); - } - if (::testing::Test::HasFailure()) { - RemoveGlobbedFiles(debugGlob); - unsetenv("NCCL_DEBUG_FILE"); - unsetenv("NCCL_DEBUG_SUBSYS"); - unsetenv("NCCL_DEBUG"); - return; - } - for (int i = 0; i < nExpect; ++i) { - bool const sawLL = DebugLogsHaveProtocolForBytes(debugGlob, "LL", expect[i].bytes); - bool const sawSimple = DebugLogsHaveProtocolForBytes(debugGlob, "Simple", expect[i].bytes); - if (expect[i].expectLL) { - EXPECT_TRUE(sawLL) << "per-peer bytes " << expect[i].bytes - << " expected LL; the SIMPLE channel estimate must not hide it"; - } else { - EXPECT_TRUE(sawSimple) << "per-peer bytes " << expect[i].bytes << " expected SIMPLE"; - EXPECT_FALSE(sawLL) << "per-peer bytes " << expect[i].bytes - << " must stay SIMPLE (the 64 KiB/peer cap)"; - } - } - RemoveGlobbedFiles(debugGlob); - unsetenv("NCCL_DEBUG_FILE"); - unsetenv("NCCL_DEBUG_SUBSYS"); - unsetenv("NCCL_DEBUG"); - } - - // Default threshold. 16/32/64 KiB per peer must be LL; 64 KiB+1 and 128 KiB-512 KiB must stay - // SIMPLE. Isolated so a prior test's NCCL_P2P_LL_THRESHOLD cannot latch the check. - TEST(SendRecv, SingleNodeGfx94LlCrossover) - { - RUN_ISOLATED_TEST_WITH_ENV( - "SingleNodeGfx94LlCrossover", - []() { - unsetenv("NCCL_P2P_LL_THRESHOLD"); - P2pLlCrossoverExpect const expect[] = { - {16384, true}, {32768, true}, {65536, true}, - {65537, false}, {131072, false}, {262144, false}, {524288, false}, - }; - RunSingleNodeGfx94LlCrossover(expect, sizeof(expect) / sizeof(expect[0])); - }, - {{"NCCL_ALLOC_P2P_NET_LL_BUFFERS", "0"}, - {"NCCL_P2P_LL128_ENABLE", "0"}, - {"RCCL_DDA_ENABLE", "0"}}); - } - - // An explicit NCCL_P2P_LL_THRESHOLD keeps the stock per-channel rule, so 16 KiB/peer stays SIMPLE. - TEST(SendRecv, SingleNodeGfx94LlCrossoverHonorsUserThreshold) - { - RUN_ISOLATED_TEST_WITH_ENV( - "SingleNodeGfx94LlCrossoverHonorsUserThreshold", - []() { - P2pLlCrossoverExpect const expect[] = {{16384, false}}; - RunSingleNodeGfx94LlCrossover(expect, 1); - }, - {{"NCCL_P2P_LL_THRESHOLD", "4096"}, - {"NCCL_ALLOC_P2P_NET_LL_BUFFERS", "0"}, - {"NCCL_P2P_LL128_ENABLE", "0"}, - {"RCCL_DDA_ENABLE", "0"}}); - } } diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index 5109cc0179f..17f03eda48d 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -4361,6 +4361,66 @@ struct Gfx1250AddP2pScene { return nullptr; } }; + +// Single-node all-XGMI gfx942/gfx950 legacy-LL gate. cudaArch is 100*major+10*minor, +// so gfx942 is 940; BatchPlanComm's default 942 does not enter the branch. +// p2pChunkSize is the single-node default (512 KiB), so the SIMPLE estimate is 1 +// channel at or below 64 KiB/peer. nChannelsMax is 1 on purpose: the bug is the +// protocol choice, not how many channels the LL recompute would then stripe. +struct Gfx94P2pLlScene { + BatchPlanComm bp; + std::unique_ptr topo{new ncclTopoSystem{}}; + std::vector peers; + std::vector peerSlots; + char llBuf[8]{}; + char sendMem[8]{}; + char recvMem[8]{}; + int planTotalTasks[2]{1, 1}; + + explicit Gfx94P2pLlScene(int cudaArch) : bp(/*nNodes=*/1, cudaArch) { + InstallSendRecvDevFuncIds(); + ncclComm* c = bp.c(); + c->rank = 0; + c->nRanks = 2; + c->nNodes = 1; + c->cudaArch = cudaArch; + c->isAllNvlink = 1; + c->allocP2pNetLLBuffers = 0; + c->p2pnChannels = 1; + c->p2pnChannelsPerPeer = 1; + c->p2pChannelShiftSize = 0; + c->p2pChunkSize = 1 << 19; + c->buffSizes[NCCL_PROTO_LL] = 1 << 20; + c->buffSizes[NCCL_PROTO_LL128] = 1 << 20; + c->buffSizes[NCCL_PROTO_SIMPLE] = 1 << 20; + SetSingleGpuArch(topo.get(), cudaArch == 950 ? "gfx950" : "gfx942"); + topo->ll128Enabled = false; + c->topo = topo.get(); + + peers.resize(2); + peerSlots.resize(2); + for (int i = 0; i < 2; ++i) { + peerSlots[i] = &peers[i]; + for (int ci = 0; ci < NCCL_MAX_CONNS; ++ci) + peers[i].send[ci].conn.buffs[NCCL_PROTO_LL] = peers[i].recv[ci].conn.buffs[NCCL_PROTO_LL] = llBuf; + } + c->channels[0].peers = peerSlots.data(); + } + + ncclResult_t add(ssize_t bytes) { + struct ncclTaskP2p* tasks[2] = {nullptr, nullptr}; + return addP2pToPlan(bp.c(), bp.p(), /*nChannelsMin=*/1, /*nChannelsMax=*/1, /*p2pRound=*/0, + /*sendRank=*/1, sendMem, bytes, /*recvRank=*/1, recvMem, bytes, + /*sendOpCount=*/0, /*recvOpCount=*/0, planTotalTasks, tasks); + } + + ncclDevWorkP2p* p2pWork() { + for (auto* node = bp.p()->workQueue.head; node != nullptr; node = node->next) { + if (node->workType == ncclDevWorkTypeP2p) return reinterpret_cast(node + 1); + } + return nullptr; + } +}; } // namespace TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx1250Enable1Mixed2KiB8KiB_OneLl128WorkItem) { @@ -4553,6 +4613,47 @@ TEST_F(EnqueueMicrotest, AddPersistentPlanRef_DropsNonLiveRefsAndKeepsLiveOnes) ncclPersistentPlanRefDrop(registered); } +// Pin the numeric cutoff through the redirected ncclParamP2pLLThreshold (g_loadParam, +// no cache). Presence of NCCL_P2P_LL_THRESHOLD is a separate ncclGetEnv read: the +// widen applies only when that name is absent, so an explicit 4096 stays on the +// stock per-channel rule even though the number matches the default. +void PinP2pLlThreshold(bool userSet) { + if (userSet) SetMicroEnv("NCCL_P2P_LL_THRESHOLD", "4096"); + else SetMicroEnvAbsent("NCCL_P2P_LL_THRESHOLD"); + SetParam("P2P_LL_THRESHOLD", 4096); +} + +void ExpectGfx94RecvProtoLL(int cudaArch, ssize_t bytes, int protoLL) { + Gfx94P2pLlScene sc(cudaArch); + ASSERT_EQ(ncclSuccess, sc.add(bytes)) << "arch=" << cudaArch << " bytes=" << bytes; + auto* work = sc.p2pWork(); + ASSERT_NE(nullptr, work) << "arch=" << cudaArch << " bytes=" << bytes; + EXPECT_EQ(protoLL, int(work->recvProtoLL)) << "arch=" << cudaArch << " bytes=" << bytes; + EXPECT_EQ(protoLL, int(work->sendProtoLL)) << "arch=" << cudaArch << " bytes=" << bytes; +} + +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94UnsetLlThreshold_WidensLlThrough64KiB) { + PinP2pLlThreshold(/*userSet=*/false); + const struct { + ssize_t bytes; + int protoLL; + } cases[] = { + {16 << 10, 1}, {32 << 10, 1}, {64 << 10, 1}, {(64 << 10) + 1, 0}, + {128 << 10, 0}, {256 << 10, 0}, {512 << 10, 0}, + }; + for (const auto& c : cases) ExpectGfx94RecvProtoLL(940, c.bytes, c.protoLL); + // gfx950 shares the gate. 942 is not 100*9+10*4, so the default BatchPlanComm + // arch must stay SIMPLE at 16 KiB. + ExpectGfx94RecvProtoLL(950, 16 << 10, 1); + ExpectGfx94RecvProtoLL(942, 16 << 10, 0); +} + +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94UserLlThreshold_KeepsPerChannelRule) { + PinP2pLlThreshold(/*userSet=*/true); + ExpectGfx94RecvProtoLL(940, 16 << 10, 0); + ExpectGfx94RecvProtoLL(950, 16 << 10, 0); +} + // =========================================================================== // rawTaskAppend / ncclCeInitJob (the NCCL_ENQUEUE_REARCH_ENABLE=1 path), then argument validation arms. // =========================================================================== @@ -5442,4 +5543,3 @@ TEST_F(EnqueueMicrotest, TopoGetAlgoInfo_Nvls_UsesNvlsChannelsCappedByCommOnlyFo EXPECT_EQ(c.algo, task.algorithm); EXPECT_EQ(c.nMaxChannels, task.nMaxChannels) << c.algo << " nNodes=" << c.nNodes; } -} From 222660089d696f0821008b0658e2e8ef8327716c Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Thu, 5 Nov 2026 08:16:49 -0600 Subject: [PATCH 05/10] test(rccl): lock the gfx94 P2P LL widen gates nNodes, isAllNvlink, and the LL128 opt-in each keep 32 KiB on SIMPLE. Dropping any of those checks would select LL, and the new cases fail. Co-authored-by: Cursor --- projects/rccl/test/host/enqueue-test.cc | 62 +++++++++++++++++++------ 1 file changed, 48 insertions(+), 14 deletions(-) diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index 17f03eda48d..ba8d89e3bac 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -4367,25 +4367,36 @@ struct Gfx1250AddP2pScene { // p2pChunkSize is the single-node default (512 KiB), so the SIMPLE estimate is 1 // channel at or below 64 KiB/peer. nChannelsMax is 1 on purpose: the bug is the // protocol choice, not how many channels the LL recompute would then stripe. +struct Gfx94P2pLlOpts { + int cudaArch = 940; + int nNodes = 1; + int isAllNvlink = 1; + bool ll128OptIn = false; +}; + struct Gfx94P2pLlScene { BatchPlanComm bp; std::unique_ptr topo{new ncclTopoSystem{}}; std::vector peers; std::vector peerSlots; char llBuf[8]{}; + char ll128Buf[8]{}; char sendMem[8]{}; char recvMem[8]{}; int planTotalTasks[2]{1, 1}; - explicit Gfx94P2pLlScene(int cudaArch) : bp(/*nNodes=*/1, cudaArch) { + explicit Gfx94P2pLlScene(Gfx94P2pLlOpts opts) : bp(opts.nNodes, opts.cudaArch) { InstallSendRecvDevFuncIds(); ncclComm* c = bp.c(); c->rank = 0; c->nRanks = 2; - c->nNodes = 1; - c->cudaArch = cudaArch; - c->isAllNvlink = 1; - c->allocP2pNetLLBuffers = 0; + c->nNodes = opts.nNodes; + c->maxLocalRanks = 2; // nNodes > 1 divides the P2P channel base by this + c->cudaArch = opts.cudaArch; + c->isAllNvlink = opts.isAllNvlink; + c->allocP2pNetLLBuffers = opts.ll128OptIn; + c->ll128LineElems = 8; + c->ll128DataElems = 7; c->p2pnChannels = 1; c->p2pnChannelsPerPeer = 1; c->p2pChannelShiftSize = 0; @@ -4393,16 +4404,18 @@ struct Gfx94P2pLlScene { c->buffSizes[NCCL_PROTO_LL] = 1 << 20; c->buffSizes[NCCL_PROTO_LL128] = 1 << 20; c->buffSizes[NCCL_PROTO_SIMPLE] = 1 << 20; - SetSingleGpuArch(topo.get(), cudaArch == 950 ? "gfx950" : "gfx942"); - topo->ll128Enabled = false; + SetSingleGpuArch(topo.get(), opts.cudaArch == 950 ? "gfx950" : "gfx942"); + topo->ll128Enabled = opts.ll128OptIn; c->topo = topo.get(); peers.resize(2); peerSlots.resize(2); for (int i = 0; i < 2; ++i) { peerSlots[i] = &peers[i]; - for (int ci = 0; ci < NCCL_MAX_CONNS; ++ci) + for (int ci = 0; ci < NCCL_MAX_CONNS; ++ci) { peers[i].send[ci].conn.buffs[NCCL_PROTO_LL] = peers[i].recv[ci].conn.buffs[NCCL_PROTO_LL] = llBuf; + peers[i].send[ci].conn.buffs[NCCL_PROTO_LL128] = peers[i].recv[ci].conn.buffs[NCCL_PROTO_LL128] = ll128Buf; + } } c->channels[0].peers = peerSlots.data(); } @@ -4623,13 +4636,19 @@ void PinP2pLlThreshold(bool userSet) { SetParam("P2P_LL_THRESHOLD", 4096); } -void ExpectGfx94RecvProtoLL(int cudaArch, ssize_t bytes, int protoLL) { - Gfx94P2pLlScene sc(cudaArch); - ASSERT_EQ(ncclSuccess, sc.add(bytes)) << "arch=" << cudaArch << " bytes=" << bytes; +void ExpectGfx94RecvProtoLL(Gfx94P2pLlOpts opts, ssize_t bytes, int protoLL) { + Gfx94P2pLlScene sc(opts); + ASSERT_EQ(ncclSuccess, sc.add(bytes)) << "arch=" << opts.cudaArch << " bytes=" << bytes; auto* work = sc.p2pWork(); - ASSERT_NE(nullptr, work) << "arch=" << cudaArch << " bytes=" << bytes; - EXPECT_EQ(protoLL, int(work->recvProtoLL)) << "arch=" << cudaArch << " bytes=" << bytes; - EXPECT_EQ(protoLL, int(work->sendProtoLL)) << "arch=" << cudaArch << " bytes=" << bytes; + ASSERT_NE(nullptr, work) << "arch=" << opts.cudaArch << " bytes=" << bytes; + EXPECT_EQ(protoLL, int(work->recvProtoLL)) << "arch=" << opts.cudaArch << " bytes=" << bytes; + EXPECT_EQ(protoLL, int(work->sendProtoLL)) << "arch=" << opts.cudaArch << " bytes=" << bytes; +} + +void ExpectGfx94RecvProtoLL(int cudaArch, ssize_t bytes, int protoLL) { + Gfx94P2pLlOpts opts; + opts.cudaArch = cudaArch; + ExpectGfx94RecvProtoLL(opts, bytes, protoLL); } TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94UnsetLlThreshold_WidensLlThrough64KiB) { @@ -4654,6 +4673,21 @@ TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94UserLlThreshold_KeepsPerChannelRule) ExpectGfx94RecvProtoLL(950, 16 << 10, 0); } +// 32 KiB is inside the widen (stock cutoff is 4 KiB). Each gate that is off must +// stay SIMPLE; dropping nNodes, isAllNvlink, or !useLL128OptIn would pick LL. +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94WidenGates_StaySimple) { + PinP2pLlThreshold(/*userSet=*/false); + Gfx94P2pLlOpts multiNode; + multiNode.nNodes = 2; + ExpectGfx94RecvProtoLL(multiNode, 32 << 10, 0); + Gfx94P2pLlOpts pcie; + pcie.isAllNvlink = 0; + ExpectGfx94RecvProtoLL(pcie, 32 << 10, 0); + Gfx94P2pLlOpts ll128; + ll128.ll128OptIn = true; + ExpectGfx94RecvProtoLL(ll128, 32 << 10, 0); +} + // =========================================================================== // rawTaskAppend / ncclCeInitJob (the NCCL_ENQUEUE_REARCH_ENABLE=1 path), then argument validation arms. // =========================================================================== From 4445e454728bb5686b0dfb821582d2b66f3f0c26 Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Thu, 5 Nov 2026 08:17:13 -0600 Subject: [PATCH 06/10] test(rccl): separate the P2P LL cap from the chunk-size term A 256 KiB chunk stops LL at 32 KiB. A 1 MiB chunk stays capped at 64 KiB. The default 512 KiB chunk makes those two bounds equal. Co-authored-by: Cursor --- projects/rccl/test/host/enqueue-test.cc | 19 ++++++++++++++++++- 1 file changed, 18 insertions(+), 1 deletion(-) diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index ba8d89e3bac..8bc5f66cbe6 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -4372,6 +4372,7 @@ struct Gfx94P2pLlOpts { int nNodes = 1; int isAllNvlink = 1; bool ll128OptIn = false; + int p2pChunkSize = 1 << 19; }; struct Gfx94P2pLlScene { @@ -4400,7 +4401,7 @@ struct Gfx94P2pLlScene { c->p2pnChannels = 1; c->p2pnChannelsPerPeer = 1; c->p2pChannelShiftSize = 0; - c->p2pChunkSize = 1 << 19; + c->p2pChunkSize = opts.p2pChunkSize; c->buffSizes[NCCL_PROTO_LL] = 1 << 20; c->buffSizes[NCCL_PROTO_LL128] = 1 << 20; c->buffSizes[NCCL_PROTO_SIMPLE] = 1 << 20; @@ -4688,6 +4689,22 @@ TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94WidenGates_StaySimple) { ExpectGfx94RecvProtoLL(ll128, 32 << 10, 0); } +// The default 512 KiB chunk makes p2pChunkSize/8 equal the 64 KiB cap, so either +// operand of the min would pass. A smaller chunk stops at chunk/8. A larger one +// stays at the cap. +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94ChunkSize_BoundsTheWiden) { + PinP2pLlThreshold(/*userSet=*/false); + Gfx94P2pLlOpts small; + small.p2pChunkSize = 256 << 10; + ExpectGfx94RecvProtoLL(small, 32 << 10, 1); + ExpectGfx94RecvProtoLL(small, (32 << 10) + 1, 0); + Gfx94P2pLlOpts large; + large.p2pChunkSize = 1 << 20; + ExpectGfx94RecvProtoLL(large, 64 << 10, 1); + ExpectGfx94RecvProtoLL(large, (64 << 10) + 1, 0); + ExpectGfx94RecvProtoLL(large, 128 << 10, 0); +} + // =========================================================================== // rawTaskAppend / ncclCeInitJob (the NCCL_ENQUEUE_REARCH_ENABLE=1 path), then argument validation arms. // =========================================================================== From 04100f63b2114da261784933504261f23d088a7e Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Thu, 5 Nov 2026 08:19:56 -0600 Subject: [PATCH 07/10] fix(rccl): treat an empty P2P LL threshold as unset ncclLoadParam already ignores NCCL_P2P_LL_THRESHOLD when the value is empty and uses the default. The widen now uses that same rule, and keeps reading it on each call so a later set or unset stays visible. Co-authored-by: Cursor --- projects/rccl/src/enqueue/enqueue.cc | 13 +++++++++---- projects/rccl/test/host/enqueue-test.cc | 14 +++++++++++--- 2 files changed, 20 insertions(+), 7 deletions(-) diff --git a/projects/rccl/src/enqueue/enqueue.cc b/projects/rccl/src/enqueue/enqueue.cc index 2d8090eb3bc..57741267ad4 100644 --- a/projects/rccl/src/enqueue/enqueue.cc +++ b/projects/rccl/src/enqueue/enqueue.cc @@ -1291,6 +1291,12 @@ static ncclResult_t scheduleCollTasksToPlan(struct ncclComm* comm, struct ncclKe // relatively early -- internode alltoall sweeps show it regressing past ~16-32 KiB/channel. Default // 4 KiB: the largest cutoff that avoids the legacy-LL mid-size regressions while keeping its win. NCCL_PARAM(P2pLLThreshold, "P2P_LL_THRESHOLD", 4096); +// Non-empty only. ncclLoadParam treats "" as the default, so an empty value is not a user set. +// Read on each call: this is not the parameter's cached load, and a later set or unset must show. +static bool p2pLLThresholdUserSet() { + const char* str = ncclGetEnv("NCCL_P2P_LL_THRESHOLD"); + return str && str[0] != '\0'; +} // Separate, independent threshold for the LL128 latency path (used on gfx942/gfx950 with // NCCL_ALLOC_P2P_NET_LL_BUFFERS=1 and comm LL128 enabled, and on gfx1250 when // NCCL_P2P_LL128_ENABLE=1). Default 16 KiB: internode alltoall/scatter/gather sweeps (1-16 nodes) @@ -1524,10 +1530,9 @@ static ncclResult_t addP2pToPlan(struct ncclComm* comm, struct ncclKernelPlan* p ssize_t latencyThreshold = useLL128OptIn ? ncclParamP2pLL128Threshold() : ncclParamP2pLLThreshold(); ssize_t latencyBytesMax = nChannels[dir] * latencyThreshold; // SIMPLE's channel count is 1 here, so the per-channel LL threshold never selects LL on - // single-node gfx942/gfx950. Allow LL up to p2pChunkSize/8, capped at 64 KiB/peer. A set - // NCCL_P2P_LL_THRESHOLD, and the LL128 path, keep the stock rule. The env is read on - // each call so a later set or unset is visible. - bool const userSetP2pLLThreshold = ncclGetEnv("NCCL_P2P_LL_THRESHOLD") != nullptr; + // single-node gfx942/gfx950. Allow LL up to p2pChunkSize/8, capped at 64 KiB/peer. A + // non-empty NCCL_P2P_LL_THRESHOLD, and the LL128 path, keep the stock rule. + bool const userSetP2pLLThreshold = p2pLLThresholdUserSet(); if (comm->nNodes <= 1 && comm->isAllNvlink && !useLL128OptIn && (comm->cudaArch == 940 || comm->cudaArch == 950) && !userSetP2pLLThreshold) { latencyBytesMax = std::max(latencyBytesMax, std::min(comm->p2pChunkSize / 8, 65536)); diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index 8bc5f66cbe6..12aabce9319 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -4628,9 +4628,9 @@ TEST_F(EnqueueMicrotest, AddPersistentPlanRef_DropsNonLiveRefsAndKeepsLiveOnes) } // Pin the numeric cutoff through the redirected ncclParamP2pLLThreshold (g_loadParam, -// no cache). Presence of NCCL_P2P_LL_THRESHOLD is a separate ncclGetEnv read: the -// widen applies only when that name is absent, so an explicit 4096 stays on the -// stock per-channel rule even though the number matches the default. +// no cache). Whether the user set NCCL_P2P_LL_THRESHOLD is a separate ncclGetEnv +// read: the widen applies when that name is absent or empty, so an explicit 4096 +// stays on the stock per-channel rule even though the number matches the default. void PinP2pLlThreshold(bool userSet) { if (userSet) SetMicroEnv("NCCL_P2P_LL_THRESHOLD", "4096"); else SetMicroEnvAbsent("NCCL_P2P_LL_THRESHOLD"); @@ -4674,6 +4674,14 @@ TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94UserLlThreshold_KeepsPerChannelRule) ExpectGfx94RecvProtoLL(950, 16 << 10, 0); } +// NCCL_P2P_LL_THRESHOLD= is present but empty. ncclLoadParam uses the default, so +// the widen still applies. A nullptr-only check would treat this as user-set. +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94EmptyLlThreshold_Widens) { + SetMicroEnv("NCCL_P2P_LL_THRESHOLD", ""); + SetParam("P2P_LL_THRESHOLD", 4096); + ExpectGfx94RecvProtoLL(940, 16 << 10, 1); +} + // 32 KiB is inside the widen (stock cutoff is 4 KiB). Each gate that is off must // stay SIMPLE; dropping nNodes, isAllNvlink, or !useLL128OptIn would pick LL. TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94WidenGates_StaySimple) { From e0c4de8c6f9cdb90fedbcae6f3632580884e9999 Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Thu, 5 Nov 2026 08:23:03 -0600 Subject: [PATCH 08/10] docs(rccl): record the gfx94 P2P LL threshold widen NCCL_P2P_LL_THRESHOLD is documented in the user guide, not in env-variables.rst. Describe the unset single-node cap and add the changelog entry for the regression fix. Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 1 + projects/rccl/docs/userguide/source/env.rst | 4 +++- 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index bf2ae721ea6..97b4acb203b 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -84,6 +84,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Fixed a crash (`HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION`) on gfx110x and gfx120x when the Tree algorithm runs over shared memory with more channels than the Tree graph search found, for example with `NCCL_ALGO=Tree` and `NCCL_P2P_DISABLE=1`. Channels added after `ncclTopoPreset` had no intra-node tree links; `ncclTopoPostset` now sets the tree and CollNet chain links on every channel, as it already did for rings. * Fixed the host one-sided RMA API being unavailable on multi-node communicators whose ranks had fused NICs. NIC fusion merges devices for the regular transport only; GIN and RMA select from their own topology node types and never receive a fused device. * Fixed Anvil SDMA GIN small-message AllToAll latency after `ncclGinFenceLevel` support. A sub-threshold `gin.put`, a `gin.putValue`, or a standalone `gin.signal` emitted two system-scope fences per peer: one in `fenceBeforeSignal` and one in `signalPeer`. `signalPeer` now emits a single system-scope release immediately before the remote atomic, rather than an agent-scope release before it and a system-scope fence after it, which also stops payload stores from passing the signal cell. A clean-queue IPC `SignalInc` skips `fenceBeforeSignal` entirely, since it would only repeat that release. The SDMA `quiet()` before a signal now runs only when the call itself issued SDMA; previously any call carrying a signal or counter quieted the peer queue, including an IPC put after an earlier SDMA put to the same peer, which now relies on `Flush` to drain that queue. The `NCCL_GIN_ANVIL_SDMA_IPC_AGENT_FENCE` debug knob is removed. Each `Flush` thread drains the peers it owns and then clears exactly the bits it drained, over a mask disjoint from every other thread's. Previously `Flush` loaded a snapshot, quieted it, and cleared the whole word on the coop's leading thread, on every rank, which could drop a concurrent mark whose doorbell had not yet fired and then skip the quiet that mark needed. Clearing after the drain also keeps the invariant that a clear bit means a drained queue, so a second coop flushing the same context cannot return from `Flush` while the first is still inside `quiet()`. +* Fixed P2P protocol selection on single-node all-XGMI gfx942 and gfx950. The NCCL 2.29 channel estimate counted SIMPLE channels, which is 1 below 64 KiB per peer, so the 4 KiB `NCCL_P2P_LL_THRESHOLD` never selected legacy LL and transfers in that range lost the low-latency path. With the variable unset or empty, legacy LL is now used up to the smaller of the SIMPLE part size (`NCCL_P2P_NVL_CHUNKSIZE / 8`, 64 KiB by default) and 64 KiB per peer. A non-empty `NCCL_P2P_LL_THRESHOLD` still uses the per-channel rule, and the LL128 path is unchanged. Other architectures, PCIe-only pairs, and multi-node communicators are unchanged. ### Known issues diff --git a/projects/rccl/docs/userguide/source/env.rst b/projects/rccl/docs/userguide/source/env.rst index 9f579057e0b..9211cd0cd45 100644 --- a/projects/rccl/docs/userguide/source/env.rst +++ b/projects/rccl/docs/userguide/source/env.rst @@ -1945,9 +1945,11 @@ NCCL_P2P_LL_THRESHOLD The ``NCCL_P2P_LL_THRESHOLD`` is the maximum per-channel message size (in bytes) at or below which RCCL uses the legacy LL low-latency protocol for P2P (send/recv) operations. Above this size, RCCL uses the SIMPLE protocol. This threshold applies when the P2P latency path uses the legacy LL protocol (i.e. when the LL128 path is not active for the communicator). Because legacy LL carries roughly 2x wire overhead, it stops beating SIMPLE at relatively small sizes, so this defaults lower than ``NCCL_P2P_LL128_THRESHOLD``. +On single-node all-XGMI gfx942 and gfx950, an unset or empty value does not use that per-channel product. The SIMPLE channel estimate is 1 for messages up to the SIMPLE part size, so ``nChannels * 4096`` would not select LL. RCCL instead allows legacy LL up to the smaller of that part size (``NCCL_P2P_NVL_CHUNKSIZE / 8``, 64 KiB with the default 512 KiB chunk) and 64 KiB per peer. A non-empty value keeps the per-channel rule on every architecture. The LL128 path uses ``NCCL_P2P_LL128_THRESHOLD`` instead and is not widened. + Values accepted ^^^^^^^^^^^^^^^ -Decimal number. Default is 4096. +Decimal number. Default is 4096. An empty value is treated as unset. NCCL_P2P_LL128_THRESHOLD ------------------------ From a4479a205368413ed7d7587a0da826996eaf9139 Mon Sep 17 00:00:00 2001 From: Corey Derochie <161367113+corey-derochie-amd@users.noreply.github.com> Date: Thu, 5 Nov 2026 08:45:38 -0600 Subject: [PATCH 09/10] Fix formatting in enqueue-test.cc --- projects/rccl/test/host/enqueue-test.cc | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index 12aabce9319..d5e9a4e09ca 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -5602,3 +5602,4 @@ TEST_F(EnqueueMicrotest, TopoGetAlgoInfo_Nvls_UsesNvlsChannelsCappedByCommOnlyFo EXPECT_EQ(c.algo, task.algorithm); EXPECT_EQ(c.nMaxChannels, task.nMaxChannels) << c.algo << " nNodes=" << c.nNodes; } +} From 0d2e0a8dc656d2840ebd13d9a3572dabf3b74c4e Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Fri, 6 Nov 2026 11:36:27 -0600 Subject: [PATCH 10/10] test(rccl): keep the multi-node P2P LL gate off the unfaked PXN path nNodes > 1 asks rcclParamPxnOptQpUsage, which the host fake aborts on. Disable PXN for that case, matching the production default, so the channel base is computed and the test can check the protocol. Co-authored-by: Cursor --- projects/rccl/test/host/enqueue-test.cc | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index d5e9a4e09ca..4520ed36497 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -4686,9 +4686,15 @@ TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94EmptyLlThreshold_Widens) { // stay SIMPLE; dropping nNodes, isAllNvlink, or !useLL128OptIn would pick LL. TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94WidenGates_StaySimple) { PinP2pLlThreshold(/*userSet=*/false); - Gfx94P2pLlOpts multiNode; - multiNode.nNodes = 2; - ExpectGfx94RecvProtoLL(multiNode, 32 << 10, 0); + // nNodes > 1 takes the channel base through rcclParamPxnOptQpUsage, which the + // host fake aborts on. PXN disabled is the production default and skips that + // call. The widen still requires a single node. + { + ScopedHook pxnDisable(g_pxnDisable, [](struct ncclComm*) { return 1; }); + Gfx94P2pLlOpts multiNode; + multiNode.nNodes = 2; + ExpectGfx94RecvProtoLL(multiNode, 32 << 10, 0); + } Gfx94P2pLlOpts pcie; pcie.isAllNvlink = 0; ExpectGfx94RecvProtoLL(pcie, 32 << 10, 0);