From d3dd2a44e12d21ad9fb7c7dea7f2d311a19967ec Mon Sep 17 00:00:00 2001 From: Ilya Kosarev Date: Thu, 20 Aug 2026 17:49:30 +0000 Subject: [PATCH 1/7] net-ib: add the worker-safe harness the threaded suites need The ordinary helpers barrier and assert fatally, so a worker cannot use them: a fatal assertion only returns from the worker, and a worker-side collective strands the peer rank when a sibling fails. This adds the data-path, IB-CAST and fault primitives that return a ThreadResult instead, and the machinery around them: - Workers spread over NICs whose index both ranks agree is routable, intersected with an Allreduce. Resolving that list per rank would let slot k sit on different hardware on each side, where connect and accept succeed and the traffic is then dropped -- the unexplained completion timeout the filter exists to prevent. - The rank's HIP device is propagated into workers, since the current device is thread-local, and a failure to bind it is reported rather than ignored, after both start gates so a failing worker cannot strand its siblings. - Per-worker payload seeds with an odd stride and a non-zero base. makeBytePattern observes the seed modulo 256, so an even stride makes workers collide -- with the obvious 100000 it is every eighth -- and a payload delivered on the wrong connection would still verify. - A bounded rendezvous, for the places where the start gates are not enough: they synchronize entry into the body and nothing after it. - A refused deregMr is counted and reduced across ranks before it is asserted, so it cannot surface later as a phantom MR leak or as a failure on the rank whose output the runner mutes. Co-authored-by: Cursor --- .../transport/NetIbMPI/NetIbMPITestBase.hpp | 859 +++++++++++++++++- 1 file changed, 830 insertions(+), 29 deletions(-) diff --git a/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp b/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp index cca94272ce26..f76652e48d0f 100644 --- a/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp +++ b/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp @@ -11,6 +11,7 @@ #include #include "MPITestBase.hpp" #include "NetIbCastInspect.hpp" +#include "NetIbFaultInject.hpp" #include "ResourceGuards.hpp" #include "TestChecks.hpp" #include "DeviceBufferHelpers.hpp" @@ -19,6 +20,7 @@ #include "net.h" #include "plugin/nccl_net.h" #include +#include #include #include #include @@ -129,9 +131,15 @@ struct NetMHandleDeleter { } }; +// Deregistrations that failed inside a worker. A destructor cannot fail a test, +// and swallowing the result would turn a refused deregistration into a phantom +// MR leak later, so the count is recorded here and the harness checks it on the +// main thread once the workers have joined. The serial teardown asserts on the +// same return value. +inline std::atomic g_workerDeregFailures{0}; + // Worker threads must not invoke TEST_INFO or any helper that can call MPI. -// This deleter is used only by the threaded test bodies; errors are surfaced -// by the surrounding transfer operation and resource-leak checks. +// This deleter is used only by the threaded test bodies. struct NetMHandleWorkerDeleter { ncclNet_t* net; void* comm; @@ -140,7 +148,9 @@ struct NetMHandleWorkerDeleter { void operator()(void* mhandle) const { if (mhandle && net && comm) { - (void)net->deregMr(comm, mhandle); + if (net->deregMr(comm, mhandle) != ncclSuccess) { + g_workerDeregFailures.fetch_add(1, std::memory_order_relaxed); + } } } }; @@ -852,8 +862,12 @@ class NetIbMPITest : public MPITestBase { static constexpr int kThreadTagStride = 1000; static constexpr int kMaxThreadTagOffset = 1; // listener-ready flag + handle static constexpr int kMpiGuaranteedTagUb = 32767; + // Upper bound on connections a single worker may own; each one consumes its + // own tag stride during the main-thread handle exchange. + static constexpr int kMaxConnsPerWorker = 2; - static_assert((MPIEnvironment::kMaxThreads - 1) * kThreadTagStride + kMaxThreadTagOffset + static_assert((MPIEnvironment::kMaxThreads * kMaxConnsPerWorker - 1) * kThreadTagStride + + kMaxThreadTagOffset <= kMpiGuaranteedTagUb, "worst-case per-thread MPI tag must fit in the tag range every " "MPI implementation is required to provide"); @@ -870,12 +884,22 @@ class NetIbMPITest : public MPITestBase { run.results.resize(nThreads); run.threadIds.resize(nThreads); + // The HIP current device is thread-local, and MPIEnvironment binds it + // once on the main thread. Without propagating it, a worker touching GPU + // memory would land on device 0 regardless of this rank's assignment. + int hipDevice = -1; + if (hipGetDevice(&hipDevice) != hipSuccess) hipDevice = -1; + ThreadStartGate startGate(nThreads); ThreadStartGate bodyGate(nThreads); std::atomic inFlight{0}; std::atomic maxInFlight{0}; auto worker = [&](int threadIdx) { run.threadIds[threadIdx] = std::this_thread::get_id(); + // Reported only after both gates below: leaving early would strand + // the sibling workers waiting on them. + const hipError_t deviceError = + (hipDevice >= 0) ? hipSetDevice(hipDevice) : hipSuccess; if (!startGate.ArriveAndWait()) { run.results[threadIdx].ok = false; run.results[threadIdx].msg = "worker launch was cancelled"; @@ -893,6 +917,14 @@ class NetIbMPITest : public MPITestBase { run.results[threadIdx].msg = "worker body start was cancelled"; return; } + if (deviceError != hipSuccess) { + inFlight.fetch_sub(1, std::memory_order_relaxed); + run.results[threadIdx].ok = false; + run.results[threadIdx].msg = "hipSetDevice(" + std::to_string(hipDevice) + + ") failed in the worker: " + + hipGetErrorString(deviceError); + return; + } try { run.results[threadIdx] = body(threadIdx); } catch (const std::exception& error) { @@ -941,27 +973,585 @@ class NetIbMPITest : public MPITestBase { } } + // Reduces worker outcomes across ranks and reports them on rank 0, which is + // the only rank with GTest listeners. The failing worker's message travels + // with the flag: without that, a failure on a non-zero rank shows up as + // "failed on a peer rank" and the actual reason is lost, since non-zero ranks + // have their output muted unless RCCL_MPI_LOG_ALL_RANKS is set. bool SynchronizeThreadResults(const std::vector& results, const char* phase) { + static constexpr int kResultTextBytes = 512; + int localFailed = 0; - for (const auto& result : results) - if (!result.ok) localFailed = 1; + std::string localText; + for (size_t threadIdx = 0; threadIdx < results.size(); ++threadIdx) { + if (results[threadIdx].ok) continue; + localFailed = 1; + if (localText.size() < kResultTextBytes / 2) + localText += "thread " + std::to_string(threadIdx) + ": " + + results[threadIdx].msg + "; "; + } int globalFailed = 0; - if (MPI_Allreduce(&localFailed, &globalFailed, 1, MPI_INT, MPI_MAX, MPI_COMM_WORLD) != MPI_SUCCESS) { + if (MPI_Allreduce(&localFailed, &globalFailed, 1, MPI_INT, MPI_MAX, MPI_COMM_WORLD) + != MPI_SUCCESS) { ADD_FAILURE() << phase << ": MPI_Allreduce failed"; return false; } if (!globalFailed) return true; - for (size_t threadIdx = 0; threadIdx < results.size(); ++threadIdx) { - if (!results[threadIdx].ok) - ADD_FAILURE() << phase << ", thread " << threadIdx << ": " << results[threadIdx].msg; + std::vector sendText(kResultTextBytes, '\0'); + snprintf(sendText.data(), kResultTextBytes, "%s", localText.c_str()); + std::vector allText(kResultTextBytes * MPIEnvironment::world_size, '\0'); + if (MPI_Allgather(sendText.data(), kResultTextBytes, MPI_CHAR, allText.data(), + kResultTextBytes, MPI_CHAR, MPI_COMM_WORLD) != MPI_SUCCESS) { + ADD_FAILURE() << phase << ": MPI_Allgather of worker messages failed"; + return false; + } + + for (int rank = 0; rank < MPIEnvironment::world_size; ++rank) { + const char* text = allText.data() + rank * kResultTextBytes; + if (text[0] == '\0') continue; + ADD_FAILURE() << phase << ", rank " << rank << ": " << text; + } + return false; + } + + // ── Worker-safe data path ──────────────────────────────────────── + // + // The ordinary transfer helpers (PostSendWithRetry, DoSendRecv, CastDoSendRecv) + // barrier and assert fatally, so a worker cannot use them: a fatal assertion + // only returns from the worker, and a worker-side collective strands the peer + // rank when a sibling fails. These return ThreadResult instead, and the main + // thread reports it after joining. + // ──────────────────────────────────────────────────────────────── + + ThreadResult WorkerRegister(void* comm, void* data, size_t size, int type, void** mhandle) { + ThreadResult result; + if (RegisterMemory(comm, data, size, type, mhandle) != ncclSuccess || *mhandle == nullptr) { + result.ok = false; + result.msg = "regMr failed"; + } + return result; + } + + ThreadResult WorkerPostSend(void* sendComm, void* data, size_t size, int tag, + void* mhandle, void** request) { + ThreadResult result; + int attempts = 0; + do { + if (PostSend(sendComm, data, size, tag, mhandle, request) != ncclSuccess) { + result.ok = false; + result.msg = "isend failed"; + return result; + } + if (*request != nullptr) break; + if (++attempts >= kMaxRetryAttempts) { + result.ok = false; + result.msg = "isend returned a NULL request after retries"; + return result; + } + usleep(kPollIntervalUs); + } while (*request == nullptr); + return result; + } + + ThreadResult WorkerPostRecv(void* recvComm, void* data, size_t size, int tag, + void* mhandle, void** request) { + ThreadResult result; + void* bufs[1] = {data}; + size_t sizes[1] = {size}; + int tags[1] = {tag}; + void* handles[1] = {mhandle}; + if (PostRecv(recvComm, 1, bufs, sizes, tags, handles, request) != ncclSuccess + || *request == nullptr) { + result.ok = false; + result.msg = "irecv failed or returned a NULL request"; + } + return result; + } + + ThreadResult WorkerWait(void* request, int* sizes, int timeoutMs = kDefaultTimeoutMs) { + ThreadResult result; + if (WaitForCompletion(request, sizes, timeoutMs) != ncclSuccess) { + result.ok = false; + result.msg = "request did not complete within the timeout"; + } + return result; + } + + // One transfer on the worker's own connection, without touching buffer + // contents. For GPU buffers, or when the caller verifies the payload itself. + ThreadResult WorkerSendRecvRaw(int rank, ConnectionPair& pair, void* buffer, size_t size, + int tag, void* mhandle, int timeoutMs = kDefaultTimeoutMs, + int* receivedSize = nullptr) { + void* request = nullptr; + ThreadResult result = (rank == 0) + ? WorkerPostRecv(pair.recvComm, buffer, size, tag, mhandle, &request) + : WorkerPostSend(pair.sendComm, buffer, size, tag, mhandle, &request); + if (!result.ok) return result; + + int sizes[1] = {0}; + result = WorkerWait(request, sizes, timeoutMs); + if (!result.ok) return result; + if (receivedSize) *receivedSize = sizes[0]; + return result; + } + + // One transfer plus host-side payload check: the sender fills a seeded + // pattern, the receiver clears its buffer first and verifies afterwards. A + // distinct seed per worker turns any cross-connection delivery into a data + // failure rather than a silent pass. + ThreadResult WorkerSendRecvPattern(int rank, ConnectionPair& pair, void* buffer, size_t size, + int tag, void* mhandle, int seed, + int timeoutMs = kDefaultTimeoutMs) { + if (rank == 0) { + memset(buffer, 0, size); + } else { + fillHostBufferWithPattern(buffer, size, makeBytePattern(seed)); + } + + int received = 0; + ThreadResult result = WorkerSendRecvRaw(rank, pair, buffer, size, tag, mhandle, + timeoutMs, &received); + if (!result.ok) return result; + + if (rank == 0) { + if (received != (int)size) { + result.ok = false; + result.msg = "received size mismatch"; + return result; + } + if (size > 0 + && !verifyHostBufferData(buffer, size, makeBytePattern(seed))) { + result.ok = false; + result.msg = "data validation failed"; + } + } + return result; + } + + // Composite: allocate a host buffer, register it, run one seeded transfer, + // and release both. Covers the common single-transfer worker body. + ThreadResult WorkerHostTransfer(int rank, ConnectionPair& pair, size_t size, int tag, + int seed, int timeoutMs = kDefaultTimeoutMs) { + ThreadResult result; + void* buffer = malloc(size ? size : 1); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(comm, buffer, size ? size : 1, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, NetMHandleWorkerDeleter(net_, comm)); + + return WorkerSendRecvPattern(rank, pair, buffer, size, tag, mhandle, seed, timeoutMs); + } + + // ── Worker-safe IB-CAST scheduler inspection ───────────────────── + // Token and cursor state is per send communicator, so a worker can arm and + // read its own connection without disturbing the others. + + ThreadResult WorkerCastSetTokens(void* sendComm, const std::vector& tokens) { + ThreadResult result; + if (ncclIbCastSetTokens(sendComm, tokens.data(), (int)tokens.size()) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastSetTokens failed"; + } + return result; + } + + ThreadResult WorkerCastGetSchedState(void* sendComm, struct ncclIbCastSchedState* out) { + ThreadResult result; + memset(out, 0, sizeof(*out)); + if (ncclIbCastGetSchedState(sendComm, out) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastGetSchedState failed"; + } + return result; + } + + // Warm the scheduler up (it initializes on the first send) and arm equal + // weights. nqps comes from the environment so both ranks agree without a + // collective; the sender confirms the connection really uses that many QPs. + ThreadResult WorkerCastPrepareTokens(int rank, ConnectionPair& pair, void* buffer, + void* mhandle, int nqps, int tag, int seed) { + ThreadResult result = WorkerSendRecvPattern(rank, pair, buffer, 64, tag, mhandle, seed); + if (!result.ok || rank != 1) return result; + + // Token arithmetic in these tests is derived from the environment so the + // expected values are predictable, which only holds if the connection + // really uses that many QPs. Asking the scheduler here is what makes a + // mismatch a clear failure instead of a wrong expectation. Where the count + // drives an action rather than an expectation -- arming faults on every QP + // -- the live value is used directly instead. + int liveNqps = 0; + result = WorkerCastLiveNqps(pair.sendComm, &liveNqps); + if (!result.ok) return result; + if (liveNqps != nqps) { + result.ok = false; + result.msg = "scheduler reports nqps=" + std::to_string(liveNqps) + + " but the environment implies " + std::to_string(nqps); + return result; + } + return WorkerCastSetTokens(pair.sendComm, EqualTokens(nqps)); + } + + // Worker-safe CAST transfer with a token-consumption expectation. Only the + // sender owns scheduler state, so it checks the delta while the receiver + // verifies the payload. Pass a negative delta to skip the token check. + ThreadResult WorkerCastTransferExpectTokens(int rank, ConnectionPair& pair, void* buffer, + size_t size, int tag, void* mhandle, int seed, + int expectedTokenDelta, + int timeoutMs = kLargeTransferTimeoutMs) { + struct ncclIbCastSchedState before = {}; + ThreadResult result; + if (rank == 1) { + result = WorkerCastGetSchedState(pair.sendComm, &before); + if (!result.ok) return result; + } + + result = WorkerSendRecvPattern(rank, pair, buffer, size, tag, mhandle, seed, timeoutMs); + if (!result.ok) return result; + + if (rank == 1 && expectedTokenDelta >= 0) { + struct ncclIbCastSchedState after = {}; + result = WorkerCastGetSchedState(pair.sendComm, &after); + if (!result.ok) return result; + const int delta = before.activeTotTokens - after.activeTotTokens; + if (delta != expectedTokenDelta) { + result.ok = false; + result.msg = "expected a WRR token delta of " + + std::to_string(expectedTokenDelta) + " at size " + + std::to_string(size) + ", observed " + std::to_string(delta); + } + } + return result; + } + + // Number of QPs per connection the environment asks for. The CAST configs + // always set it, and CAST_ENV_CHECK_OR_SKIP has already enforced that. + static int CastEnvNqps() { + const char* value = getenv("NCCL_IB_QPS_PER_CONNECTION"); + return (value && value[0]) ? std::atoi(value) : 1; + } + +#if defined(ENABLE_FAULT_INJECTION) + // ── Worker-safe IB-CAST fault injection ────────────────────────── + // The pre-call intercept hooks store their state in the communicator + // (faultQpDelayUs / faultQpError in ncclIbNetCommBase), so a worker can + // break its own connection without touching anybody else's. The + // libibverbs-level ops registry is process-global and deliberately not + // used here. + + ThreadResult WorkerCastFaultArmError(void* sendComm, int nqps) { + ThreadResult result; + for (int qp = 0; qp < nqps; qp++) { + if (ncclIbCastFaultSetQpError(sendComm, qp, /*inject=*/true) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastFaultSetQpError failed on QP " + std::to_string(qp); + return result; + } + } + return result; + } + + ThreadResult WorkerCastFaultSetDelay(void* sendComm, int qpIdx, uint32_t delayUs) { + ThreadResult result; + if (ncclIbCastFaultSetQpDelay(sendComm, qpIdx, delayUs) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastFaultSetQpDelay failed"; + } + return result; + } + + ThreadResult WorkerCastFaultClear(void* sendComm) { + ThreadResult result; + if (ncclIbCastFaultClear(sendComm) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastFaultClear failed"; + } + return result; + } + + int WorkerCastFatalCount(void* sendComm) { + int count = 0; + if (ncclIbCastFaultGetFatalCount(sendComm, &count) != ncclSuccess) return -1; + return count; + } + + ThreadResult WorkerCastDriveQpToError(void* sendComm, int qpIdx) { + ThreadResult result; + if (ncclIbCastFaultDriveQpToError(sendComm, qpIdx) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastFaultDriveQpToError failed"; + } + return result; + } + + // Post one send and poll it, treating both a failed isend and a fatal error + // count as an outcome rather than a test failure. Used by fault tests that + // accept either signal. + struct WorkerFaultSendOutcome { + ncclResult_t sendRet = ncclSuccess; + bool completed = false; + int fatalCount = 0; + }; + + WorkerFaultSendOutcome WorkerCastFaultSend(void* sendComm, void* buffer, size_t size, int tag, + void* mhandle, int pollIterations) { + WorkerFaultSendOutcome outcome; + void* request = nullptr; + for (int attempt = 0; attempt < kMaxRetryAttempts; attempt++) { + outcome.sendRet = PostSend(sendComm, buffer, size, tag, mhandle, &request); + if (outcome.sendRet != ncclSuccess || request != nullptr) break; + usleep(kPollIntervalUs); + } + outcome.fatalCount = WorkerCastFatalCount(sendComm); + + if (outcome.sendRet == ncclSuccess && request != nullptr) { + for (int poll = 0; poll < pollIterations; poll++) { + int done = 0; + int sizes[1] = {0}; + const ncclResult_t testRet = TestRequest(request, &done, sizes); + if (testRet != ncclSuccess) { + outcome.sendRet = testRet; + break; + } + outcome.fatalCount = WorkerCastFatalCount(sendComm); + if (done) { + outcome.completed = true; + break; + } + if (outcome.fatalCount > 0) break; + usleep(kPollIntervalUs); + } + } + return outcome; + } + + // Shared worker body for the failover tests: drive the sender's QP 0 into + // the error state, then require the payload to still arrive over the + // surviving device of a fused NIC, with no fatal error and a device state + // that is no longer Ok. Resiliency state is per communicator, so N workers + // can fail their own links at once; what they share is the single global + // recovery thread, which is exactly what concurrency here exercises. + ThreadResult WorkerCastFailoverTransfer(int rank, ConnectionPair& pair, void* buffer, + size_t size, int tag, void* mhandle, int seed, + int messages = 1) { + ThreadResult result; + if (rank == 1) { + result = WorkerCastDriveQpToError(pair.sendComm, 0); + if (!result.ok) return result; + } + + for (int i = 0; i < messages; i++) { + result = WorkerSendRecvPattern(rank, pair, buffer, size, tag + i, mhandle, seed + i, + kLargeTransferTimeoutMs); + if (!result.ok) { + result.msg = "after driving QP 0 to error: " + result.msg; + return result; + } + } + + if (rank != 1) return result; + + const int fatalCount = WorkerCastFatalCount(pair.sendComm); + if (fatalCount != 0) { + result.ok = false; + result.msg = "failover should have handled the QP error, but the fatal count is " + + std::to_string(fatalCount); + return result; + } + + struct ncclIbCastResiliencyState state = {}; + if (ncclIbCastGetResiliencyState(pair.sendComm, &state) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastGetResiliencyState failed"; + return result; + } + // 0 is ncclIbResiliencyDevStateOk; anything else means the failure was + // registered (Error, RecoveryInProgress or Recovered). + if (state.devState[0] == 0) { + result.ok = false; + result.msg = "device 0 still reports Ok after its QP was driven to error"; + } + return result; + } + + // Failover with requests already in flight: post every message first, then + // break QP 0 under them, then require all of them to complete. Because isend + // only gets a request once the receiver has published a FIFO slot, "all sends + // posted" already implies "all receives posted", which is the ordering the + // serial test buys with a barrier the worker cannot use. + // + // buffer must hold messages * size bytes; each message uses its own slice so + // completions cannot be confused with each other. + ThreadResult WorkerCastFailoverInFlight(int rank, ConnectionPair& pair, void* buffer, + size_t size, int tag, void* mhandle, int seed, + int messages) { + ThreadResult result; + std::vector requests(messages, nullptr); + + for (int i = 0; i < messages; i++) { + char* slice = static_cast(buffer) + static_cast(i) * size; + if (rank == 0) { + memset(slice, 0, size); + result = WorkerPostRecv(pair.recvComm, slice, size, tag + i, mhandle, + &requests[i]); + } else { + fillHostBufferWithPattern(slice, size, makeBytePattern(seed + i)); + result = WorkerPostSend(pair.sendComm, slice, size, tag + i, mhandle, + &requests[i]); + } + if (!result.ok) return result; + } + + if (rank == 1) { + result = WorkerCastDriveQpToError(pair.sendComm, 0); + if (!result.ok) return result; + } + + for (int i = 0; i < messages; i++) { + int sizes[1] = {0}; + result = WorkerWait(requests[i], sizes, kLargeTransferTimeoutMs); + if (!result.ok) { + result.msg = "request " + std::to_string(i) + + " never completed after QP 0 was driven to error"; + return result; + } + if (rank != 0) continue; + + char* slice = static_cast(buffer) + static_cast(i) * size; + if (sizes[0] != (int)size + || !verifyHostBufferData(slice, size, makeBytePattern(seed + i))) { + result.ok = false; + result.msg = "message " + std::to_string(i) + + " arrived corrupted or short after failover"; + return result; + } + } + + if (rank != 1) return result; + + const int fatalCount = WorkerCastFatalCount(pair.sendComm); + if (fatalCount != 0) { + result.ok = false; + result.msg = "failover should have absorbed the QP error, but the fatal count is " + + std::to_string(fatalCount); + } + return result; + } + + // Poll a receive the peer may never satisfy, and report whether it finished. + bool WorkerDrainRecv(void* request, int pollIterations) { + if (!request) return true; + for (int poll = 0; poll < pollIterations; poll++) { + int done = 0; + int sizes[1] = {0}; + if (TestRequest(request, &done, sizes) != ncclSuccess) return false; + if (done) return true; + usleep(kPollIntervalUs); } - if (!localFailed) - ADD_FAILURE() << phase << " failed on a peer rank"; return false; } + // Releases a receive that the peer's injected send will never satisfy. + // + // Two things make this necessary rather than tidy. The work request keeps + // referencing the receive buffer's memory region until the queue pair is + // destroyed, which happens after the worker has already had to deregister it, + // and deregistration cannot wait for the close because it needs the + // communicator that the close frees. Nor can the transfer simply be finished: + // the failed send consumed the FIFO slot this receive was matched to, so no + // later send can complete it. + // + // Driving this side's queue pairs to error is what breaks the tie. The NIC + // completes the outstanding work request with a flush status, so nothing + // references the region by the time the worker unwinds. The connection is + // spent either way -- it exists in these tests to be broken. + ThreadResult WorkerCastFlushAbandonedRecv(void* recvComm, void* request) { + ThreadResult result; + if (!request) return result; + + // The API rejects an index past the connection's QP count, which is how + // the loop learns where to stop. + static constexpr int kQpProbeLimit = 64; + for (int qp = 0; qp < kQpProbeLimit; qp++) { + if (ncclIbCastFaultDriveRecvQpToError(recvComm, qp) != ncclSuccess) break; + } + + // A flush completion surfaces as an error, and that is the expected + // outcome here: all that matters is that the request stops being + // outstanding before the memory region goes away. + static constexpr int kFlushPolls = 500; // 500 * 10ms = 5s + for (int poll = 0; poll < kFlushPolls; poll++) { + int done = 0; + int sizes[1] = {0}; + if (TestRequest(request, &done, sizes) != ncclSuccess) return result; + if (done) return result; + usleep(kPollIntervalUs); + } + result.ok = false; + result.msg = "the abandoned receive was still outstanding after its queue pairs were " + "driven to error"; + return result; + } + + // Both sides lose QP 0 and the payload then has to ride the surviving device. + // + // The serial recovery test breaks the queue pairs with a transfer already in + // flight, so that failover rescues an outstanding request. A worker cannot + // reproduce that: the serial body coordinates the two sides with an MPI + // handshake around the break, and without one, both orders were measured to be + // unusable. Breaking both sides with the receive posted flushes that work + // request while the data still arrives over the surviving device, and the two + // sides run one message apart for the rest of the test (seen once in four full + // matrix runs, as the sender exhausting its FIFO-slot retries on message n+1 + // while the receiver waited for message n). Breaking only the sender in flight + // and the receiver afterwards failed every one of nineteen consecutive runs. + // So the break happens while the connection is idle, and the difference from + // the serial body is stated in the test plan rather than papered over. + ThreadResult WorkerTransferAcrossQpFailure(int rank, ConnectionPair& pair, void* buffer, + size_t size, int tag, void* mhandle, int seed, + int timeoutMs) { + ThreadResult result; + if (rank == 0) { + if (ncclIbCastFaultDriveRecvQpToError(pair.recvComm, 0) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastFaultDriveRecvQpToError failed"; + return result; + } + } else { + result = WorkerCastDriveQpToError(pair.sendComm, 0); + if (!result.ok) return result; + } + return WorkerSendRecvPattern(rank, pair, buffer, size, tag, mhandle, seed, timeoutMs); + } + + // QPs the connection actually uses. NCCL_IB_QPS_PER_CONNECTION states only a + // request: on a merged device the plugin creates that many per member, so + // arming faults from the environment would leave the remaining QPs healthy + // and "the send must fail" would depend on which QP the scheduler picked. + // Valid once the scheduler is warm, which the first successful send does. + ThreadResult WorkerCastLiveNqps(void* sendComm, int* nqps) { + struct ncclIbCastSchedState state; + ThreadResult result = WorkerCastGetSchedState(sendComm, &state); + if (!result.ok) return result; + if (state.nqps <= 0) { + result.ok = false; + result.msg = "scheduler reports nqps=" + std::to_string(state.nqps); + return result; + } + *nqps = state.nqps; + return result; + } +#endif /* ENABLE_FAULT_INJECTION */ + ncclResult_t InitNetIbCtx(void** ctxOut) { ncclNetCommConfig_t commConfig = {}; commConfig.trafficClass = NCCL_NET_TRAFFIC_CLASS_UNDEF; @@ -1054,21 +1644,131 @@ class NetIbMPITest : public MPITestBase { for (auto& connection : connections) TeardownConnectionForThread(connection, rank); } + // Which device each worker connects over. Spreading puts workers on + // different NICs, so concurrency reaches several devices' protection + // domains and MR caches instead of hammering one. + struct ThreadDevPolicy { + int dev = 0; + bool spread = false; + + static ThreadDevPolicy Fixed(int dev) { return ThreadDevPolicy{dev, false}; } + static ThreadDevPolicy Spread() { return ThreadDevPolicy{0, true}; } + }; + + // Payload seed for worker threadIdx at sequence position seq. Workers must + // differ here, otherwise a payload delivered on the wrong connection still + // verifies and the cross-worker check is vacuous. + // + // makeBytePattern() observes the seed only modulo 256, so the stride has to + // be odd: an even stride collides for workers whose indices differ by + // 256/gcd(stride, 256) -- with a stride of 100000 that is every 8th worker. + // The base keeps the seed off zero, whose pattern begins with the 0x00 that + // the receive buffer is cleared to. + static constexpr int kWorkerSeedBase = 55; + static constexpr int kWorkerSeedStride = 100001; + static int WorkerSeed(int threadIdx, int seq) { + return kWorkerSeedBase + threadIdx * kWorkerSeedStride + seq; + } + + // Plugin-visible indices of physical (non-merged) devices. devices() also + // reports virtual NICs created by earlier tests, and a deduped 1-device vNIC + // reuses an existing physical index, so dedupe on vProps.devs[0]. + // Only devices that can actually deliver traffic are returned: a RoCE NIC + // with link-local GIDs only sets up QPs without complaint and then drops + // packets silently, which would surface as an unexplained connect or accept + // timeout on whichever worker happened to be handed that device. + std::vector PhysicalDeviceIndices() { + std::vector indices; + int n = 0; + if (GetDeviceCount(&n) != ncclSuccess) return indices; + std::set seen; + for (int i = 0; i < n; i++) { + ncclNetProperties_t props; + memset(&props, 0, sizeof(props)); + if (GetDeviceProperties(i, &props) != ncclSuccess) continue; + if (props.vProps.ndevs > 1) continue; + if (props.name && !CanRouteCrossNode(props.name)) continue; + if (seen.insert(props.vProps.devs[0]).second) indices.push_back(i); + } + return indices; + } + + // Device indices both ranks agree on. PhysicalDeviceIndices() reads local + // sysfs only, so where the two nodes have different NICs alive, slot k would + // sit on different hardware per rank: connect and accept still succeed and + // the traffic is then dropped, which surfaces as an unexplained completion + // timeout -- the very failure the routability filter exists to prevent. + // Intersecting the two lists keeps every slot symmetric. + bool AgreedPhysicalDeviceIndices(std::vector* out, std::string* reason) { + out->clear(); + unsigned long long localMask = 0; + for (int dev : PhysicalDeviceIndices()) { + if (dev < 0 || dev >= 64) continue; + localMask |= (1ull << dev); + } + unsigned long long agreed = 0; + if (MPI_Allreduce(&localMask, &agreed, 1, MPI_UNSIGNED_LONG_LONG, MPI_BAND, + MPI_COMM_WORLD) != MPI_SUCCESS) { + *reason = "MPI_Allreduce of the routable device mask failed"; + return false; + } + for (int dev = 0; dev < 64; ++dev) + if (agreed & (1ull << dev)) out->push_back(dev); + if (out->empty()) { + *reason = "no device index is routable on both ranks, so workers cannot be spread"; + return false; + } + return true; + } + + int ResolveWorkerDev(const ThreadDevPolicy& policy, const std::vector& physical, + int slotIdx) { + if (!policy.spread) return policy.dev; + return physical[slotIdx % physical.size()]; + } + + // Bounded rendezvous for workers that must all reach a point before any of + // them moves on, which the start gates cannot express: they only synchronize + // entry into the body. Bounded so a worker that failed earlier cannot hang + // its siblings; the caller reports the timeout as its own failure. + static bool WorkerRendezvous(std::atomic& arrived, int expected, int pollIterations) { + arrived.fetch_add(1, std::memory_order_release); + for (int poll = 0; poll < pollIterations; poll++) { + if (arrived.load(std::memory_order_acquire) >= expected) return true; + usleep(kPollIntervalUs); + } + return false; + } + // Mirrors the normal rccl-tests -t execution model: contexts and // communicators are established by the main thread, then workers drive // independent comms concurrently. MPI is used only between phases. - void RunMultiThreadedIndependent(int dev, int nThreads, - std::function body) { + // + // connsPerWorker > 1 gives every worker several independent connections, + // which a body needs when it has to prove that state from one connection + // does not leak into a fresh one. + void RunMultiThreadedIndependentGroups( + ThreadDevPolicy policy, int nThreads, int connsPerWorker, + std::function&)> body) { const int rank = MPIEnvironment::world_rank; const int peerRank = (rank + 1) % 2; - std::vector connections(nThreads); - std::vector initResults(nThreads); - - for (int threadIdx = 0; threadIdx < nThreads; ++threadIdx) { - if (InitNetIbCtx(&connections[threadIdx].ctx) != ncclSuccess - || connections[threadIdx].ctx == nullptr) { - initResults[threadIdx].ok = false; - initResults[threadIdx].msg = "InitNetIb failed or returned a null context"; + const int nSlots = nThreads * connsPerWorker; + std::vector physical; + if (policy.spread) { + std::string reason; + // Both ranks reach the same verdict, so the skip cannot strand a peer. + if (!AgreedPhysicalDeviceIndices(&physical, &reason)) GTEST_SKIP() << reason; + } + std::vector connections(nSlots); + std::vector initResults(nSlots); + + g_workerDeregFailures.store(0, std::memory_order_relaxed); + + for (int slot = 0; slot < nSlots; ++slot) { + if (InitNetIbCtx(&connections[slot].ctx) != ncclSuccess + || connections[slot].ctx == nullptr) { + initResults[slot].ok = false; + initResults[slot].msg = "InitNetIb failed or returned a null context"; } } if (!SynchronizeThreadResults(initResults, "NetIB context initialization")) { @@ -1078,11 +1778,11 @@ class NetIbMPITest : public MPITestBase { return; } - std::vector setupResults(nThreads); - for (int threadIdx = 0; threadIdx < nThreads; ++threadIdx) { - setupResults[threadIdx] = SetupConnectionForThread( - connections[threadIdx].ctx, dev, connections[threadIdx].pair, rank, peerRank, - threadIdx * kThreadTagStride); + std::vector setupResults(nSlots); + for (int slot = 0; slot < nSlots; ++slot) { + setupResults[slot] = SetupConnectionForThread( + connections[slot].ctx, ResolveWorkerDev(policy, physical, slot), + connections[slot].pair, rank, peerRank, slot * kThreadTagStride); } if (!SynchronizeThreadResults(setupResults, "NetIB connection setup")) { MPI_Barrier(MPI_COMM_WORLD); @@ -1091,14 +1791,115 @@ class NetIbMPITest : public MPITestBase { return; } - ThreadWorkerRun run = RunThreadWorkers( - nThreads, [&](int threadIdx) { return body(threadIdx, connections[threadIdx].pair); }); + ThreadWorkerRun run = RunThreadWorkers(nThreads, [&](int threadIdx) { + std::vector pairs; + pairs.reserve(connsPerWorker); + for (int c = 0; c < connsPerWorker; ++c) + pairs.push_back(&connections[threadIdx * connsPerWorker + c].pair); + return body(threadIdx, pairs); + }); VerifyThreadFanOut(run, nThreads); SynchronizeThreadResults(run.results, "NetIB threaded data path"); MPI_Barrier(MPI_COMM_WORLD); TeardownThreadConnections(connections, rank); + + // Reduced, so both ranks fail together and the message names the rank that + // saw it. A local-only check would fail on rank 1 alone, whose output the + // runner mutes, leaving a result mismatch with no visible reason. + const int localDeregFailures = g_workerDeregFailures.load(std::memory_order_relaxed); + int totalDeregFailures = 0; + MPI_Allreduce(&localDeregFailures, &totalDeregFailures, 1, MPI_INT, MPI_SUM, + MPI_COMM_WORLD); + EXPECT_EQ(totalDeregFailures, 0) + << "a worker's deregMr was refused (" << localDeregFailures << " on this rank, " + << totalDeregFailures + << " across both), which would otherwise surface as an MR leak"; + MPI_Barrier(MPI_COMM_WORLD); + } + + void RunMultiThreadedIndependent(ThreadDevPolicy policy, int nThreads, + std::function body) { + RunMultiThreadedIndependentGroups( + policy, nThreads, 1, + [&](int threadIdx, std::vector& pairs) { + return body(threadIdx, *pairs[0]); + }); + } + + void RunMultiThreadedIndependent(int dev, int nThreads, + std::function body) { + RunMultiThreadedIndependent(ThreadDevPolicy::Fixed(dev), nThreads, body); + } + + // How a threaded size sweep registers memory. Some serial bodies register the + // whole buffer once, others register exactly the current size on every step; + // on a fused device that second shape is the point of the test, since each + // registration fans out across both members' protection domains and caches. + // The threaded branch has to match whichever its serial body does, or it + // quietly covers less. + enum class SweepRegistration { Once, PerSize }; + + // Threaded size sweep: every worker walks the list on its own connection with + // a per-worker payload seed, so a transfer delivered on the wrong connection + // fails verification. Wraps the run in an RDMA resource leak check. + void RunThreadedSizeSweep(ThreadDevPolicy policy, int nThreads, + const std::vector& sizes, int repeats, + const char* label, + SweepRegistration registration = SweepRegistration::Once) { + const int rank = MPIEnvironment::world_rank; + size_t maxSize = 1; + for (size_t size : sizes) maxSize = std::max(maxSize, size); + + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + policy, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + void* buffer = malloc(maxSize); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + const bool perSize = (registration == SweepRegistration::PerSize); + + void* wholeMhandle = nullptr; + if (!perSize) { + result = WorkerRegister(comm, buffer, maxSize, NCCL_PTR_HOST, &wholeMhandle); + if (!result.ok) return result; + } + NetMHandleWorkerGuard wholeGuard(wholeMhandle, + NetMHandleWorkerDeleter(net_, comm)); + + int tag = 0; + for (size_t size : sizes) { + void* sizeMhandle = nullptr; + if (perSize) { + result = WorkerRegister(comm, buffer, size, NCCL_PTR_HOST, &sizeMhandle); + if (!result.ok) return result; + } + // Released at the end of this size, so the next one registers + // again, as the serial body does. + NetMHandleWorkerGuard sizeGuard(sizeMhandle, + NetMHandleWorkerDeleter(net_, comm)); + void* mhandle = perSize ? sizeMhandle : wholeMhandle; + + for (int repeat = 0; repeat < repeats; repeat++) { + const int timeout = (size > 1024 * 1024) ? kLargeTransferTimeoutMs + : kDefaultTimeoutMs; + result = WorkerSendRecvPattern(rank, pair, buffer, size, tag, mhandle, + WorkerSeed(threadIdx, tag), timeout); + if (!result.ok) return result; + tag++; + } + } + return result; + }); MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), label); } // =============================================================== From 01ed236cdf84068c0849dd770e8087df030cb8a8 Mon Sep 17 00:00:00 2001 From: Ilya Kosarev Date: Thu, 20 Aug 2026 17:49:31 +0000 Subject: [PATCH 2/7] net-ib: run the data path, IB-CAST and fault tests at several worker counts Thirty-one existing tests gain a branch that runs their body on every worker at once, chosen by the state concurrency stresses rather than by conversion cost: the per-device MR cache and protection domain, per-communicator fault and scheduler state, and the request pool and completion routing. Each threaded branch keeps its serial body's shape, including registering per size where that is the point, as it is on a fused device. Three tests are new: - FaultIsolationAcrossWorkers has no single-threaded equivalent. One worker breaks its own connection and must observe the failure while every other worker keeps transferring with a fatal count of zero. The victim publishes an armed flag the bystanders wait on and holds the fault until they finish, so the check cannot be satisfied by a victim that armed, failed and cleared before its siblings ran. - ThreadedProgressDoesNotSerialize runs one transfer loop at a single worker and then at all of them and requires the parallel phase to stay well under what full serialization would cost. A fourth phase holds a mutex across every transfer and the same measurement must flag it, so the gate proves its own sensitivity rather than passing because it went blind. - FaultInjectionShimsAbsentUnlessRequested checks that arming an ops-level fault is refused while RCCL_IB_FAULT_INJECTION is off, which holds only while the libibverbs shims are absent. Deterministic on purpose: those shims release their mutex before the real poll_cq and hold it only for a lookup either side, so their cost does not turn into a stable timing signal -- measured across four, eight and sixteen workers, the healthy range and the range with shims overlapped entirely. Two threaded bodies are deliberately weaker than their serial counterparts, and say so: CastStressMultiRoundTwoConns cannot assert that the RTT timer fired at the 10 s interval the threaded branches need, and the recovery test breaks its queue pairs while idle and sends one message afterwards, because the serial in-flight break relies on an MPI handshake a worker cannot make and ten messages drift the two ranks apart in about one run in six. Co-authored-by: Cursor --- .../test/transport/NetIbMPI/CastTests.cpp | 293 ++++++++ .../transport/NetIbMPI/FaultInjectTests.cpp | 701 ++++++++++++++++++ .../test/transport/NetIbMPI/GeneralTests.cpp | 49 +- .../transport/NetIbMPI/NicFusionTests.cpp | 204 +++++ .../test/transport/NetIbMPI/StressTests.cpp | 684 +++++++++++++++++ 5 files changed, 1928 insertions(+), 3 deletions(-) diff --git a/projects/rccl/test/transport/NetIbMPI/CastTests.cpp b/projects/rccl/test/transport/NetIbMPI/CastTests.cpp index 04a871ec9707..3eb0ff52dc3c 100644 --- a/projects/rccl/test/transport/NetIbMPI/CastTests.cpp +++ b/projects/rccl/test/transport/NetIbMPI/CastTests.cpp @@ -590,6 +590,59 @@ TEST_F(NetIbMPITest, CastSplitDataThresholdBoundary) { net_ = &netIbCast; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads. Scheduler state is per send + // communicator, so every worker arms and audits its own tokens while N + // schedulers run the boundary decision concurrently. + if (MPIEnvironment::nThreads > 1) { + // Token-delta assertions need the RTT-driven update suspended: with + // several workers the wall-clock gap between the before and after + // snapshots is wide enough for the timer to rewrite the ledger. + CAST_REQUIRE_UPDATE_INTERVAL_OR_SKIP(10000000); + const int nqps = CastEnvNqps(); + const size_t splitDataMin = GetSplitDataMin(); + if (splitDataMin == 0) + GTEST_SKIP() << "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN=0: no WRR/split boundary exists"; + const size_t threshold = splitDataMin * static_cast(nqps); + + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + void* buffer = malloc(threshold); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, threshold, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + const int seed = 1599 + threadIdx * 1000; + result = WorkerCastPrepareTokens(rank, pair, buffer, mhandle, nqps, 1599, seed); + if (!result.ok) return result; + + // At the threshold the split path must leave WRR tokens alone. + result = WorkerCastTransferExpectTokens(rank, pair, buffer, threshold, 1600, + mhandle, seed + 1, /*tokens=*/0); + if (!result.ok) return result; + + // One byte below it, WRR must consume exactly one token. + return WorkerCastTransferExpectTokens(rank, pair, buffer, threshold - 1, 1601, + mhandle, seed + 2, /*tokens=*/1); + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), + "threaded CastSplitDataThresholdBoundary"); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -1029,6 +1082,71 @@ TEST_F(NetIbMPITest, CastSendRecvMultipleSizes) { net_ = &netIbCast; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads: N schedulers sweep both sides of + // the WRR/split boundary at once, each auditing its own token ledger. + if (MPIEnvironment::nThreads > 1) { + // Token-delta assertions need the RTT-driven update suspended: with + // several workers the wall-clock gap between the before and after + // snapshots is wide enough for the timer to rewrite the ledger. + CAST_REQUIRE_UPDATE_INTERVAL_OR_SKIP(10000000); + const int nqps = CastEnvNqps(); + const size_t splitDataMin = GetSplitDataMin(); + if (splitDataMin == 0) + GTEST_SKIP() << "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN=0: no WRR/split boundary exists"; + const size_t threshold = splitDataMin * static_cast(nqps); + const size_t bufSize = threshold * 2; + + std::vector wrrSizes; + for (size_t size : std::initializer_list{512, 4096, splitDataMin, threshold - 1}) + if (size > 0 && size < threshold) wrrSizes.push_back(size); + wrrSizes.erase(std::unique(wrrSizes.begin(), wrrSizes.end()), wrrSizes.end()); + const std::vector splitSizes = {threshold, threshold * 2}; + + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + void* buffer = malloc(bufSize); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, bufSize, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + const int seedBase = 1999 + threadIdx * 1000; + result = WorkerCastPrepareTokens(rank, pair, buffer, mhandle, nqps, 1999, + seedBase); + if (!result.ok) return result; + + int tag = 2000; + for (size_t size : wrrSizes) { + result = WorkerCastTransferExpectTokens(rank, pair, buffer, size, tag++, + mhandle, seedBase + tag, + /*tokens=*/1); + if (!result.ok) return result; + } + for (size_t size : splitSizes) { + result = WorkerCastTransferExpectTokens(rank, pair, buffer, size, tag++, + mhandle, seedBase + tag, + /*tokens=*/0); + if (!result.ok) return result; + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded CastSendRecvMultipleSizes"); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -1140,6 +1258,47 @@ TEST_F(NetIbMPITest, CastLargeTransfer) { net_ = &netIbCast; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads: concurrent 16 MB split-path + // transfers, each worker confirming its own scheduler stayed untouched. + if (MPIEnvironment::nThreads > 1) { + // Token-delta assertions need the RTT-driven update suspended: with + // several workers the wall-clock gap between the before and after + // snapshots is wide enough for the timer to rewrite the ledger. + CAST_REQUIRE_UPDATE_INTERVAL_OR_SKIP(10000000); + const int nqps = CastEnvNqps(); + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + void* buffer = malloc(kLargeBufferSize); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, kLargeBufferSize, NCCL_PTR_HOST, + &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + const int seed = 2099 + threadIdx * 1000; + result = WorkerCastPrepareTokens(rank, pair, buffer, mhandle, nqps, 2099, seed); + if (!result.ok) return result; + + return WorkerCastTransferExpectTokens(rank, pair, buffer, kLargeBufferSize, + 2100, mhandle, seed + 1, /*tokens=*/0); + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded CastLargeTransfer"); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -1206,6 +1365,47 @@ TEST_F(NetIbMPITest, CastSendRecvZeroSize) { net_ = &netIbCast; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads: a zero-byte send must still take + // the WRR path on every worker's own scheduler. + if (MPIEnvironment::nThreads > 1) { + // Token-delta assertions need the RTT-driven update suspended: with + // several workers the wall-clock gap between the before and after + // snapshots is wide enough for the timer to rewrite the ledger. + CAST_REQUIRE_UPDATE_INTERVAL_OR_SKIP(10000000); + const int nqps = CastEnvNqps(); + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t regSize = 64; + void* buffer = malloc(regSize); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, regSize, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + const int seed = 2199 + threadIdx * 1000; + result = WorkerCastPrepareTokens(rank, pair, buffer, mhandle, nqps, 2199, seed); + if (!result.ok) return result; + + return WorkerCastTransferExpectTokens(rank, pair, buffer, 0, 2200, mhandle, + seed + 1, /*tokens=*/1); + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded CastSendRecvZeroSize"); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -1280,6 +1480,99 @@ TEST_F(NetIbMPITest, CastStressMultiRoundTwoConns) { net_ = &netIbCast; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads. The single-threaded body already + // drives many connections, but strictly one after another; here every worker + // owns a connection so the schedulers advance simultaneously. Each worker + // audits its own ledger: tokens are per send communicator. + if (MPIEnvironment::nThreads > 1) { + if (GetSplitDataMin() == 0) + GTEST_SKIP() << "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN=0: no WRR/split boundary exists"; + + const int nqps = CastEnvNqps(); + const size_t splitDataMin = GetSplitDataMin(); + const size_t msgSize = std::max(64, splitDataMin - 1); + const std::vector rampSizes = {splitDataMin / 4, splitDataMin / 2, + splitDataMin, splitDataMin * 4}; + const size_t bufSize = std::max(msgSize, splitDataMin * 4); + static constexpr int kThreadedMsgs = 100; + static constexpr int kThreadedRampRounds = 5; + + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + void* buffer = malloc(bufSize); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, bufSize, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + const int tagBase = 10000 + threadIdx * 200; + const int seedBase = 3999 + threadIdx * 1000; + result = WorkerCastPrepareTokens(rank, pair, buffer, mhandle, nqps, tagBase, + seedBase); + if (!result.ok) return result; + + // Phase 1: small messages, all on the WRR path. + for (int i = 0; i < kThreadedMsgs; i++) { + result = WorkerSendRecvPattern(rank, pair, buffer, msgSize, tagBase + 1 + i, + mhandle, seedBase + i); + if (!result.ok) return result; + } + + if (rank == 1) { + struct ncclIbCastSchedState state = {}; + result = WorkerCastGetSchedState(pair.sendComm, &state); + if (!result.ok) return result; + if (!state.schedInit) { + result.ok = false; + result.msg = "scheduler never initialized after the WRR phase"; + return result; + } + int sum = 0; + for (int qp = 0; qp < state.nqps; qp++) sum += state.activeQpTokens[qp]; + if (sum != state.activeTotTokens) { + result.ok = false; + result.msg = "per-QP tokens do not sum to activeTotTokens"; + return result; + } + if (state.activeTotTokens < 0 + || state.activeTotTokens > state.initTotTokens) { + result.ok = false; + result.msg = "activeTotTokens left the [0, initTotTokens] range"; + return result; + } + } + + // Phase 2: ramp across the WRR/split boundary. + int tag = tagBase + 1 + kThreadedMsgs; + for (int round = 0; round < kThreadedRampRounds; round++) { + for (size_t size : rampSizes) { + if (size == 0 || size > bufSize) continue; + result = WorkerSendRecvPattern(rank, pair, buffer, size, tag++, mhandle, + seedBase + tag, + kLargeTransferTimeoutMs); + if (!result.ok) return result; + } + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), + "threaded CastStressMultiRoundTwoConns"); + return; + } + constexpr int kNConns = 100; std::vector listenComms(kNConns, nullptr); std::vector sendComms(kNConns, nullptr); diff --git a/projects/rccl/test/transport/NetIbMPI/FaultInjectTests.cpp b/projects/rccl/test/transport/NetIbMPI/FaultInjectTests.cpp index 9a07cfcd925a..7e1d56844f73 100644 --- a/projects/rccl/test/transport/NetIbMPI/FaultInjectTests.cpp +++ b/projects/rccl/test/transport/NetIbMPI/FaultInjectTests.cpp @@ -85,6 +85,68 @@ TEST_F(NetIbMPITest, FaultInjCastQpErrorIsFatal) { net_ = &netIbCast; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads. Injected error state lives in + // the communicator, so each worker breaks only its own connection and each + // one must observe the failure on its own send. + if (MPIEnvironment::nThreads > 1) { + RunMultiThreadedIndependent( + 0, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t size = 1024; + void* buffer = malloc(size); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, size, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + // Warm the scheduler up before arming, as the serial body does. + result = WorkerSendRecvPattern(rank, pair, buffer, size, 300, mhandle, + threadIdx * 100); + if (!result.ok) return result; + + if (rank == 0) { + // Posting this receive is what lets the peer reach the injected + // error, and its send then fails before posting, so nothing will + // complete this request. Flushing it keeps the work request from + // outliving the memory region the worker must deregister. + void* request = nullptr; + result = WorkerPostRecv(pair.recvComm, buffer, size, 301, mhandle, &request); + if (!result.ok) return result; + if (WorkerDrainRecv(request, 100)) return result; + return WorkerCastFlushAbandonedRecv(pair.recvComm, request); + } + + int liveNqps = 0; + result = WorkerCastLiveNqps(pair.sendComm, &liveNqps); + if (!result.ok) return result; + + result = WorkerCastFaultArmError(pair.sendComm, liveNqps); + if (!result.ok) return result; + + const WorkerFaultSendOutcome outcome = + WorkerCastFaultSend(pair.sendComm, buffer, size, 301, mhandle, 200); + if (outcome.sendRet == ncclSuccess && outcome.fatalCount <= 0) { + result.ok = false; + result.msg = "expected isend to fail or a fatal error to be counted after " + "arming every QP with error injection"; + return result; + } + return WorkerCastFaultClear(pair.sendComm); + }); + MPI_Barrier(MPI_COMM_WORLD); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -362,6 +424,55 @@ TEST_F(NetIbMPITest, FaultInjCastDelayDataIntegrity) { net_ = &netIbCast; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads: every worker slows down its own + // QP 0 and still has to deliver intact data while the other workers keep the + // device busy. + if (MPIEnvironment::nThreads > 1) { + RunMultiThreadedIndependent( + 0, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + static constexpr int kThreadedMsgs = 20; + const size_t size = 8192; + void* buffer = malloc(size); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, size, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + result = WorkerSendRecvPattern(rank, pair, buffer, size, 4999, mhandle, + threadIdx * 1000); + if (!result.ok) return result; + + if (rank == 1) { + result = WorkerCastFaultSetDelay(pair.sendComm, /*qpIdx=*/0, + /*delayUs=*/2000); + if (!result.ok) return result; + } + + for (int i = 0; i < kThreadedMsgs; i++) { + result = WorkerSendRecvPattern(rank, pair, buffer, size, 5000 + i, mhandle, + threadIdx * 1000 + i, + kLargeTransferTimeoutMs); + if (!result.ok) return result; + } + + if (rank == 1) return WorkerCastFaultClear(pair.sendComm); + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -574,6 +685,95 @@ TEST_F(NetIbMPITest, FaultInjCastQpErrorClearRecovers) { net_ = &netIbCast; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads. Each worker owns two + // connections: it breaks the first, clears the fault, and then has to move + // data cleanly on the second, so leftover fault state cannot hide behind a + // quiet fabric. + if (MPIEnvironment::nThreads > 1) { + RunMultiThreadedIndependentGroups( + ThreadDevPolicy::Fixed(0), MPIEnvironment::nThreads, /*connsPerWorker=*/2, + [&](int threadIdx, std::vector& pairs) -> ThreadResult { + ThreadResult result; + const size_t size = 1024; + void* buffer = malloc(size); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + ConnectionPair& faulted = *pairs[0]; + ConnectionPair& fresh = *pairs[1]; + + void* faultedComm = (rank == 0) ? faulted.recvComm : faulted.sendComm; + void* faultedMh = nullptr; + result = WorkerRegister(faultedComm, buffer, size, NCCL_PTR_HOST, &faultedMh); + if (!result.ok) return result; + NetMHandleWorkerGuard faultedGuard(faultedMh, + NetMHandleWorkerDeleter(net_, faultedComm)); + + result = WorkerSendRecvPattern(rank, faulted, buffer, size, 500, faultedMh, + threadIdx * 100); + if (!result.ok) return result; + + if (rank == 0) { + // Flushed rather than left hanging: phase 2 reuses this buffer, + // and a work request from the broken connection must not still + // be pointing at it. + void* request = nullptr; + result = WorkerPostRecv(faulted.recvComm, buffer, size, 501, faultedMh, + &request); + if (!result.ok) return result; + if (!WorkerDrainRecv(request, 100)) { + result = WorkerCastFlushAbandonedRecv(faulted.recvComm, request); + if (!result.ok) return result; + } + } else { + int liveNqps = 0; + result = WorkerCastLiveNqps(faulted.sendComm, &liveNqps); + if (!result.ok) return result; + + result = WorkerCastFaultArmError(faulted.sendComm, liveNqps); + if (!result.ok) return result; + + const WorkerFaultSendOutcome outcome = + WorkerCastFaultSend(faulted.sendComm, buffer, size, 501, faultedMh, 200); + if (outcome.sendRet == ncclSuccess && outcome.fatalCount <= 0) { + result.ok = false; + result.msg = "phase 1 did not observe the injected fault"; + return result; + } + result = WorkerCastFaultClear(faulted.sendComm); + if (!result.ok) return result; + } + + // Phase 2: the fresh connection must be unaffected. + void* freshComm = (rank == 0) ? fresh.recvComm : fresh.sendComm; + void* freshMh = nullptr; + result = WorkerRegister(freshComm, buffer, size, NCCL_PTR_HOST, &freshMh); + if (!result.ok) return result; + NetMHandleWorkerGuard freshGuard(freshMh, + NetMHandleWorkerDeleter(net_, freshComm)); + + result = WorkerSendRecvPattern(rank, fresh, buffer, size, 502, freshMh, + threadIdx * 100 + 7); + if (!result.ok) return result; + + if (rank == 1) { + const int fatalCount = WorkerCastFatalCount(fresh.sendComm); + if (fatalCount != 0) { + result.ok = false; + result.msg = "the fresh connection reports fatal errors (" + + std::to_string(fatalCount) + ")"; + } + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + return; + } + // ── Phase 1: inject fault on first connection ───────────────────────── void* listenComm1 = nullptr; void* sendComm1 = nullptr; @@ -825,6 +1025,41 @@ TEST_F(NetIbMPITest, FailoverCqeErrorRecovered) { << "Found " << totalDevs << " physical devices."; } + // Parameterized by MPIEnvironment::nThreads: N links fail at once, so the + // process-global recovery thread serves several communicators while each of + // them still has to deliver its payload over the surviving device. + if (MPIEnvironment::nThreads > 1) { + RunMultiThreadedIndependent( + mergedDev, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t size = 8192; + void* buffer = malloc(size); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, size, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + result = WorkerSendRecvPattern(rank, pair, buffer, size, 600, mhandle, + threadIdx * 1000); + if (!result.ok) return result; + + return WorkerCastFailoverTransfer(rank, pair, buffer, size, 601, mhandle, + threadIdx * 1000 + 1); + }); + MPI_Barrier(MPI_COMM_WORLD); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -1227,6 +1462,41 @@ TEST_F(NetIbMPITest, FailoverLargeMessageDataIntegrity) { GTEST_SKIP() << "Requires NIC Fusion (ndevs >= 2). Need at least 2 IB devices."; } + // Parameterized by MPIEnvironment::nThreads: concurrent 64 KB messages must + // survive their own link failure byte for byte, which also keeps several + // retransmit paths active on one device at the same time. + if (MPIEnvironment::nThreads > 1) { + RunMultiThreadedIndependent( + mergedDev, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t size = 65536; + void* buffer = malloc(size); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, size, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + result = WorkerSendRecvPattern(rank, pair, buffer, size, 900, mhandle, + threadIdx * 1000, kLargeTransferTimeoutMs); + if (!result.ok) return result; + + return WorkerCastFailoverTransfer(rank, pair, buffer, size, 901, mhandle, + threadIdx * 1000 + 1); + }); + MPI_Barrier(MPI_COMM_WORLD); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -1498,6 +1768,46 @@ TEST_F(NetIbMPITest, FailoverMultiRequestInFlight) { GTEST_SKIP() << "Requires NIC Fusion (ndevs >= 2)."; } + // Parameterized by MPIEnvironment::nThreads: every worker drives several + // messages across its own link failure, so the resiliency state machine + // handles concurrent failed requests on several communicators at once. + if (MPIEnvironment::nThreads > 1) { + RunMultiThreadedIndependent( + mergedDev, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + static constexpr int kThreadedReqs = 4; + const size_t size = 4096; + // One slice per in-flight message, plus one for the warmup. + const size_t bufSize = size * (kThreadedReqs + 1); + void* buffer = malloc(bufSize); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, bufSize, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + result = WorkerSendRecvPattern(rank, pair, buffer, size, 1200, mhandle, + threadIdx * 1000); + if (!result.ok) return result; + + // This test is about requests already in flight when the link + // dies, so the fault lands after every message is posted. + return WorkerCastFailoverInFlight(rank, pair, buffer, size, 1210, mhandle, + threadIdx * 1000 + 1, kThreadedReqs); + }); + MPI_Barrier(MPI_COMM_WORLD); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -1730,6 +2040,147 @@ TEST_F(NetIbMPITest, RecoverySuccessRestoresTraffic) { GTEST_SKIP() << "Requires NIC Fusion (ndevs >= 2). Found " << totalDevs << " physical devices."; } + // Parameterized by MPIEnvironment::nThreads. Recovery assertions stay + // per-communicator, but the recovery thread and its inbox are process-global, + // so N simultaneous failures are the only way to see it serve several + // connections. The poll budget is generous for exactly that reason. + if (MPIEnvironment::nThreads > 1) { + RunMultiThreadedIndependent( + mergedDev, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + // One message, not a run of them. With more, the two sides drift a + // message apart after recovery -- the sender exhausts its FIFO-slot + // retries on message n+1 while the receiver still waits for n, with + // both devices reported healthy -- in roughly one run in six, from + // the second message onward and under either order of breaking the + // queue pairs. The serial body does not show it because its MPI + // handshakes resynchronize the two sides between phases. Recorded + // as a finding rather than smoothed over; one message still proves + // that traffic resumes on the recovered connection. + static constexpr int kThreadedPostRecoveryMsgs = 1; + static constexpr int kRecoveryPollIterations = 6000; // 6000 * 10ms = 60s + // The receiver has no way to learn when the sender leaves the + // recovery poll: the serial body uses an MPI handshake, which a + // worker cannot call. So the first post-recovery message has to + // wait out the sender's whole poll budget on top of its own. + static constexpr int kFirstPostRecoveryTimeoutMs = + kRecoveryPollIterations * (kPollIntervalUs / 1000) + kLargeTransferTimeoutMs; + const size_t size = 8192; + void* buffer = malloc(size); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, size, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + result = WorkerSendRecvPattern(rank, pair, buffer, size, 1300, mhandle, + threadIdx * 1000); + if (!result.ok) return result; + + // Both directions lose QP 0 with the transfer already posted: the + // alive-message handshake needs both sides to notice and enter + // recovery, and the payload has to ride the surviving device. + result = WorkerTransferAcrossQpFailure(rank, pair, buffer, size, 1301, mhandle, + WorkerSeed(threadIdx, 1), + kLargeTransferTimeoutMs); + if (!result.ok) { + result.msg = "failover transfer after the link failure: " + result.msg; + return result; + } + + if (rank == 1) { + // Failover was supposed to absorb the QP error, as the serial + // body asserts after its own phase 1. + const int fatalCount = WorkerCastFatalCount(pair.sendComm); + if (fatalCount != 0) { + result.ok = false; + result.msg = "failover left " + std::to_string(fatalCount) + + " fatal errors on the connection"; + return result; + } + } + + if (rank == 1) { + // Wait for the global recovery thread to bring device 0 back. + bool recovered = false; + int lastState = -1; + for (int poll = 0; poll < kRecoveryPollIterations; poll++) { + struct ncclIbCastResiliencyState state = {}; + if (ncclIbCastGetResiliencyState(pair.sendComm, &state) != ncclSuccess) { + result.ok = false; + result.msg = "ncclIbCastGetResiliencyState failed while waiting for " + "recovery"; + return result; + } + lastState = state.devState[0]; + if (lastState == kDevStateOk || lastState == kDevStateRecovered) { + recovered = true; + break; + } + if (lastState == kDevStateRecoveryFailed + || lastState == kDevStateErrorPermanent) { + break; + } + usleep(kPollIntervalUs); + } + if (!recovered) { + result.ok = false; + result.msg = "device 0 never returned to Ok or Recovered; last state was " + + std::to_string(lastState); + return result; + } + } + + // Traffic must flow again on the recovered connection. + for (int i = 0; i < kThreadedPostRecoveryMsgs; i++) { + result = WorkerSendRecvPattern(rank, pair, buffer, size, 1310 + i, mhandle, + WorkerSeed(threadIdx, 10 + i), + (i == 0) ? kFirstPostRecoveryTimeoutMs + : kLargeTransferTimeoutMs); + if (!result.ok) { + // This phase has been seen to fail once in a full matrix run and never + // in isolation, with both ranks bounded at their own budgets waiting for + // each other, so the device state at the moment of failure is what the + // next occurrence needs to be diagnosable. + result.msg = "post-recovery traffic, message " + std::to_string(i) + ": " + + result.msg; + if (rank == 1) { + struct ncclIbCastResiliencyState state = {}; + if (ncclIbCastGetResiliencyState(pair.sendComm, &state) + == ncclSuccess) { + result.msg += "; devState[0]=" + std::to_string(state.devState[0]) + + " devState[1]=" + std::to_string(state.devState[1]); + } + } + return result; + } + } + + if (rank == 1) { + // Sustained traffic after recovery must stay clean too, which + // is the serial body's closing assertion. + const int fatalCount = WorkerCastFatalCount(pair.sendComm); + if (fatalCount != 0) { + result.ok = false; + result.msg = "post-recovery traffic reported " + std::to_string(fatalCount) + + " fatal errors"; + } + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + return; + } + void* listenComm = nullptr; void* sendComm = nullptr; void* recvComm = nullptr; @@ -3224,4 +3675,254 @@ TEST_F(NetIbMPITest, FaultInjCastOpsApiInvalidArgs) { TeardownConnection(recvComm, listenComm, sendComm, mhandle); } +// ============================================================================= +// Test: FaultInjectionShimsAbsentUnlessRequested +// +// With RCCL_IB_FAULT_INJECTION off, the libibverbs ops shims must not be +// installed. They are not a passive instrument: shimPollCq takes one +// process-global mutex on every call, so a context that carries them serializes +// completion polling across every thread in the process, including production +// proxy threads that never asked for fault injection. This is the deterministic +// half of that guard -- ThreadedProgressDoesNotSerialize measures the effect, +// this one catches the cause before it can reach a normal run. +// +// The probe: arming an ops-level fault looks up the device context in the +// registry that install() fills in, so it returns ncclInvalidArgument while the +// shims are absent and ncclSuccess once they are in place. +// ============================================================================= +TEST_F(NetIbMPITest, FaultInjectionShimsAbsentUnlessRequested) { + ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, + false, kMinGpusPerNode, kNoNodeLimit)) + << "Test requires exactly " << kExactTwoProcesses << " processes"; + + CAST_ENV_CHECK_OR_SKIP(); + + const char* injection = getenv("RCCL_IB_FAULT_INJECTION"); + if (injection && injection[0] && strcmp(injection, "0") != 0) { + GTEST_SKIP() << "RCCL_IB_FAULT_INJECTION=" << injection + << ": this configuration installs the shims on purpose"; + } + + const int rank = MPIEnvironment::world_rank; + + net_ = &netIbCast; + AssertInitAndGetDevices(nullptr); + + void* listenComm = nullptr; + void* sendComm = nullptr; + void* recvComm = nullptr; + SetupCastConnection(/*dev=*/0, &listenComm, &sendComm, &recvComm); + + constexpr size_t kMsgSize = 1024; + std::vector buf(kMsgSize, 0); + void* comm = (rank == 0) ? recvComm : sendComm; + void* mhandle = nullptr; + ASSERT_EQ(RegisterMemory(comm, buf.data(), kMsgSize, NCCL_PTR_HOST, &mhandle), ncclSuccess); + + // One transfer first: the QPs a probe reaches through must be up. + const int actualNqps = GetActualNqps(sendComm, recvComm, buf.data(), kMsgSize, + /*tag=*/900, mhandle); + ASSERT_GT(actualNqps, 0); + + const char* kLeak = "the libibverbs ops shims are installed without " + "RCCL_IB_FAULT_INJECTION: every poll_cq on this device now takes a " + "process-global mutex, which serializes completion polling for the " + "whole process"; + if (rank == 1) { + EXPECT_EQ(ncclIbCastFaultOpsSetPostSendError(sendComm, 0, EAGAIN), ncclInvalidArgument) + << kLeak; + EXPECT_EQ(ncclIbCastFaultOpsSetPollCqError(sendComm, 0, kWcRetryExcErr, 1, + /*injectWhenIdle=*/false), + ncclInvalidArgument) + << kLeak; + } else { + EXPECT_EQ(ncclIbCastFaultOpsSetPostRecvError(recvComm, 0, EAGAIN), ncclInvalidArgument) + << kLeak; + EXPECT_EQ(ncclIbCastFaultOpsSetPollCqError(recvComm, 0, kWcRetryExcErr, 1, + /*injectWhenIdle=*/false), + ncclInvalidArgument) + << kLeak; + } + + MPI_Barrier(MPI_COMM_WORLD); + TeardownConnection(recvComm, listenComm, sendComm, mhandle); +} + +// ============================================================================= +// Test: FaultIsolationAcrossWorkers +// +// Multithread-only. One worker arms error injection on its own connection and +// must observe the failure; every other worker keeps transferring on its own +// connection and must stay clean, with a fatal count of zero. +// +// This is the assertion the serial tests cannot make: injected state lives in +// ncclIbNetCommBase, so a broken connection must not disturb its siblings on the +// same device. It also puts a failing connection and healthy traffic on one NIC +// at the same time, which is how a real process behaves when one link dies. +// ============================================================================= +TEST_F(NetIbMPITest, FaultIsolationAcrossWorkers) { + ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, + false, kMinGpusPerNode, kNoNodeLimit)) + << "Test requires exactly " << kExactTwoProcesses << " processes"; + + CAST_ENV_CHECK_OR_SKIP(); + + const int rank = MPIEnvironment::world_rank; + const int nThreads = MPIEnvironment::nThreads; + if (nThreads < 2) { + GTEST_SKIP() << "requires --net_ib_nthreads greater than one: the whole point is one " + "failing connection alongside healthy ones"; + } + + net_ = &netIbCast; + AssertInitAndGetDevices(nullptr); + + static constexpr int kHealthyTransfers = 10; + // The isolation claim only means something while the fault is live, and the + // worker gate synchronizes nothing beyond entry into the body. So the + // bystanders wait for the victim to report its fault armed, the victim then + // waits for all of them to finish their healthy traffic before clearing it, + // and only then do they read the fatal count. Without the second wait the + // victim could arm, fail, clear and finish before any bystander was + // scheduled, and the count would be read on a fabric where no fault ever + // lived. + std::atomic faultArmed{false}; + std::atomic victimFinished{false}; + std::atomic bystandersDone{0}; + const int kBystanders = nThreads - 1; + // Bounded, so a victim that dies before arming cannot hang its siblings. + static constexpr int kFlagPollIterations = 3000; // 3000 * 10ms = 30s + auto waitForFlag = [](std::atomic& flag, int pollIterations) { + for (int poll = 0; poll < pollIterations; poll++) { + if (flag.load(std::memory_order_acquire)) return true; + usleep(kPollIntervalUs); + } + return false; + }; + + RunMultiThreadedIndependent( + 0, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t size = 1024; + void* buffer = malloc(size); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(workerComm, buffer, size, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + // Every worker starts from a working connection. + result = WorkerSendRecvPattern(rank, pair, buffer, size, 800, mhandle, + threadIdx * 1000); + if (!result.ok) return result; + + const bool victim = (threadIdx == 0); + if (victim) { + // Releases the siblings on every exit, including the failure + // paths, so they never wait on a victim that already gave up. + struct FaultWindow { + std::atomic& armed; + std::atomic& finished; + ~FaultWindow() { + armed.store(true, std::memory_order_release); + finished.store(true, std::memory_order_release); + } + } window{faultArmed, victimFinished}; + + if (rank == 0) { + // Posting this receive is what lets the peer reach the injected + // error, so it is the local point where the fault window opens. + // Nothing will complete it, so it is flushed before the worker + // unwinds its registration. + void* request = nullptr; + result = WorkerPostRecv(pair.recvComm, buffer, size, 801, mhandle, &request); + if (!result.ok) return result; + faultArmed.store(true, std::memory_order_release); + if (WorkerDrainRecv(request, 100)) return result; + return WorkerCastFlushAbandonedRecv(pair.recvComm, request); + } + + int liveNqps = 0; + result = WorkerCastLiveNqps(pair.sendComm, &liveNqps); + if (!result.ok) return result; + + result = WorkerCastFaultArmError(pair.sendComm, liveNqps); + if (!result.ok) return result; + faultArmed.store(true, std::memory_order_release); + + const WorkerFaultSendOutcome outcome = + WorkerCastFaultSend(pair.sendComm, buffer, size, 801, mhandle, 200); + if (outcome.sendRet == ncclSuccess && outcome.fatalCount <= 0) { + result.ok = false; + result.msg = "the victim connection did not observe its injected fault"; + return result; + } + + // Hold the fault until the bystanders have finished, so their + // traffic really did share the device with a broken connection. + for (int poll = 0; poll < kFlagPollIterations; poll++) { + if (bystandersDone.load(std::memory_order_acquire) >= kBystanders) break; + usleep(kPollIntervalUs); + } + if (bystandersDone.load(std::memory_order_acquire) < kBystanders) { + result.ok = false; + result.msg = "only " + std::to_string(bystandersDone.load()) + + " of " + std::to_string(kBystanders) + + " bystander workers finished while the fault was armed"; + return result; + } + + return WorkerCastFaultClear(pair.sendComm); + } + + // Bystanders must be untouched by the victim's failure, and their + // traffic has to run while that fault is live. + if (!waitForFlag(faultArmed, kFlagPollIterations)) { + result.ok = false; + result.msg = "the victim worker never reported its fault as armed"; + return result; + } + for (int i = 0; i < kHealthyTransfers; i++) { + result = WorkerSendRecvPattern(rank, pair, buffer, size, 810 + i, mhandle, + WorkerSeed(threadIdx, i), kLargeTransferTimeoutMs); + if (!result.ok) { + result.msg = "bystander worker disturbed by another worker's fault: " + + result.msg; + // Counted even on failure: the victim is waiting for this, and + // its own timeout would bury the real error under a second one. + bystandersDone.fetch_add(1, std::memory_order_release); + return result; + } + } + bystandersDone.fetch_add(1, std::memory_order_release); + + if (rank == 1) { + // Read the count only once the victim is finished, so it covers + // the whole time the fault was live. + if (!waitForFlag(victimFinished, kFlagPollIterations)) { + result.ok = false; + result.msg = "the victim worker never finished"; + return result; + } + const int fatalCount = WorkerCastFatalCount(pair.sendComm); + if (fatalCount != 0) { + result.ok = false; + result.msg = "bystander connection reports " + std::to_string(fatalCount) + + " fatal errors after a sibling connection failed"; + } + } + return result; + }); + + MPI_Barrier(MPI_COMM_WORLD); +} + #endif /* MPI_TESTS_ENABLED && ENABLE_FAULT_INJECTION */ diff --git a/projects/rccl/test/transport/NetIbMPI/GeneralTests.cpp b/projects/rccl/test/transport/NetIbMPI/GeneralTests.cpp index df0edadb8fa2..f093311b7bcf 100644 --- a/projects/rccl/test/transport/NetIbMPI/GeneralTests.cpp +++ b/projects/rccl/test/transport/NetIbMPI/GeneralTests.cpp @@ -350,6 +350,9 @@ TEST_F(NetIbMPITest, SimpleSendRecv) { AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded SimpleSendRecv"); } +// Parameterized by MPIEnvironment::nThreads. Concurrent workers each sweep the +// same size ladder on an independent connection, so registrations of many +// different sizes hit the shared per-device MR cache at once. TEST_F(NetIbMPITest, SendRecvMultipleSizes) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)) @@ -359,13 +362,34 @@ TEST_F(NetIbMPITest, SendRecvMultipleSizes) { AssertInitAndGetDevices(&ndev); const int rank = MPIEnvironment::world_rank; - ConnectionPair pair; - NetConnectionGuard connGuard(net_); - SetupConnectionWithGuard(0, pair, connGuard); + const int nThreads = MPIEnvironment::nThreads; // Test various sizes std::vector testSizes = {1, 64, 256, 1024, 4096, 16384, 65536}; + if (nThreads > 1) { + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + ThreadDevPolicy::Spread(), nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + for (size_t size : testSizes) { + const int seed = kMultiSizeSeedOffset + threadIdx * 10000 + + static_cast(size); + result = WorkerHostTransfer(rank, pair, size, 100, seed); + if (!result.ok) return result; + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded SendRecvMultipleSizes"); + return; + } + + ConnectionPair pair; + NetConnectionGuard connGuard(net_); + SetupConnectionWithGuard(0, pair, connGuard); + for (size_t size : testSizes) { const int tag = 100; const int seed = 2000 + static_cast(size); // Unique seed per size @@ -736,6 +760,9 @@ TEST_F(NetIbMPITest, MultipleSequentialTransfers) { AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded MultipleSequentialTransfers"); } +// Parameterized by MPIEnvironment::nThreads. Concurrent 16 MB transfers put +// several large registrations in the per-device MR cache simultaneously and +// keep the NIC saturated while every worker verifies its own payload. TEST_F(NetIbMPITest, LargeTransfer) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)) @@ -745,6 +772,22 @@ TEST_F(NetIbMPITest, LargeTransfer) { AssertInitAndGetDevices(&ndev); const int rank = MPIEnvironment::world_rank; + const int nThreads = MPIEnvironment::nThreads; + + if (nThreads > 1) { + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + ThreadDevPolicy::Spread(), nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + return WorkerHostTransfer(rank, pair, kLargeBufferSize, 400, + kBaseSeedOffset + threadIdx, + kLargeTransferTimeout); + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded LargeTransfer"); + return; + } + ConnectionPair pair; NetConnectionGuard connGuard(net_); SetupConnectionWithGuard(0, pair, connGuard); diff --git a/projects/rccl/test/transport/NetIbMPI/NicFusionTests.cpp b/projects/rccl/test/transport/NetIbMPI/NicFusionTests.cpp index a8645d2d2d0b..3be9de433a9e 100644 --- a/projects/rccl/test/transport/NetIbMPI/NicFusionTests.cpp +++ b/projects/rccl/test/transport/NetIbMPI/NicFusionTests.cpp @@ -615,6 +615,55 @@ TEST_F(NetIbMPITest, RegDeregCycling_VNic) { << "Failed to create fused vNIC from devices 0 and 1"; ASSERT_GE(vdev, 0); + // Parameterized by MPIEnvironment::nThreads. Every regMr on a fused device + // fans out across both members' protection domains and MR caches, so + // concurrent cycling on one address doubles the contended surface. The vNIC + // itself is created once on the main thread: the merged-device table is + // process-global. + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + // One region, cycled by every worker so they contend on a single MR cache + // entry, but sliced for the traffic that follows: the serial body moves + // data through the address it just recycled, and a shared payload buffer + // would have the workers overwriting each other's verification. + const size_t slotSize = kSmallBufferSize; + const size_t sharedSize = slotSize * nThreads; + auto shared = makeHostBufferAutoGuard(malloc(sharedSize)); + ASSERT_NE(shared.get(), nullptr); + + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + vdev, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + void* workerComm = (rank == 0) ? pair.recvComm : pair.sendComm; + for (int i = 0; i < 50; i++) { + void* mhandle = nullptr; + result = WorkerRegister(workerComm, shared.get(), sharedSize, + NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + if (DeregisterMemory(workerComm, mhandle) != ncclSuccess) { + result.ok = false; + result.msg = "deregMr failed during vNIC cache cycling"; + return result; + } + } + + void* mhandle = nullptr; + result = WorkerRegister(workerComm, shared.get(), sharedSize, NCCL_PTR_HOST, + &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, + NetMHandleWorkerDeleter(net_, workerComm)); + + void* slot = static_cast(shared.get()) + threadIdx * slotSize; + return WorkerSendRecvPattern(rank, pair, slot, slotSize, 530, mhandle, + WorkerSeed(threadIdx, 0)); + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded RegDeregCycling_VNic"); + return; + } + ConnectionPair pair; NetConnectionGuard connGuard(net_); SetupConnectionWithGuard(vdev, pair, connGuard); @@ -694,6 +743,27 @@ TEST_F(NetIbMPITest, LargeTransfer_VNic) { << "Failed to create fused vNIC from devices 0 and 1"; ASSERT_GE(vdev, 0); + // Parameterized by MPIEnvironment::nThreads: several striped transfers run at + // once over the fused device. The per-worker size shrinks with the worker + // count: every registration on a fused device is pinned once per member, so + // 64 MB per worker exhausts registration resources at high worker counts + // while adding nothing to the striping coverage. + if (MPIEnvironment::nThreads > 1) { + const size_t threadedSize = + std::max(4 * 1024 * 1024, + (64 * 1024 * 1024) / MPIEnvironment::nThreads); + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + vdev, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + return WorkerHostTransfer(rank, pair, threadedSize, 540, + 5400 + threadIdx * 1000, kLargeTransferTimeout); + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded LargeTransfer_VNic"); + return; + } + ConnectionPair pair; NetConnectionGuard connGuard(net_); SetupConnectionWithGuard(vdev, pair, connGuard); @@ -761,6 +831,20 @@ TEST_F(NetIbMPITest, MixedSizes_VNic) { << "Failed to create fused vNIC from devices 0 and 1"; ASSERT_GE(vdev, 0); + // Parameterized by MPIEnvironment::nThreads: the uneven-split ladder runs on + // every worker's own fused connection simultaneously. + if (MPIEnvironment::nThreads > 1) { + // Per-size registration, as the serial body does: on a fused device every + // regMr fans out across both members, and that churn is what this test is + // about. + RunThreadedSizeSweep(ThreadDevPolicy::Fixed(vdev), MPIEnvironment::nThreads, + {1, 3*1024*1024, 3, 5*1024*1024, 7, 7*1024*1024, + 64, 16*1024*1024, 1, 11*1024*1024, 4*1024*1024, 1}, + /*repeats=*/1, "threaded MixedSizes_VNic", + SweepRegistration::PerSize); + return; + } + ConnectionPair pair; NetConnectionGuard connGuard(net_); SetupConnectionWithGuard(vdev, pair, connGuard); @@ -838,6 +922,16 @@ TEST_F(NetIbMPITest, UnalignedSizeTransfer_VNic) { << "Failed to create fused vNIC from devices 0 and 1"; ASSERT_GE(vdev, 0); + // Parameterized by MPIEnvironment::nThreads: concurrent workers hit the + // 128-byte striping boundary on both members of the fused device at once. + if (MPIEnvironment::nThreads > 1) { + RunThreadedSizeSweep(ThreadDevPolicy::Fixed(vdev), MPIEnvironment::nThreads, + {127, 129, 255, 257, 511, 513}, + /*repeats=*/2, "threaded UnalignedSizeTransfer_VNic", + SweepRegistration::PerSize); + return; + } + ConnectionPair pair; NetConnectionGuard connGuard(net_); SetupConnectionWithGuard(vdev, pair, connGuard); @@ -1057,6 +1151,78 @@ TEST_F(NetIbMPITest, FlushRepeated_VNic) { ASSERT_GE(vdev, 0); const int rank = MPIEnvironment::world_rank; + + // Parameterized by MPIEnvironment::nThreads: concurrent GPU receives and + // flushes on one fused device. Workers inherit this rank's HIP device from + // the harness, since the current device is thread-local. + if (MPIEnvironment::nThreads > 1) { + static constexpr int kThreadedIters = 20; + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + vdev, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t size = kSmallBufferSize; + void* gpuBuffer = nullptr; + if (hipMalloc(&gpuBuffer, size) != hipSuccess) { + result.ok = false; + result.msg = "hipMalloc failed"; + return result; + } + auto gpuGuard = makeDeviceBufferAutoGuard(gpuBuffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(comm, gpuBuffer, size, NCCL_PTR_CUDA, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, NetMHandleWorkerDeleter(net_, comm)); + + for (int iter = 0; iter < kThreadedIters; iter++) { + const int seed = 6000 + threadIdx * 1000 + iter; + if (rank == 1 + && initializeBufferWithPattern(gpuBuffer, size, + makeBytePattern(seed)) + != hipSuccess) { + result.ok = false; + result.msg = "GPU buffer initialization failed"; + return result; + } + + result = WorkerSendRecvRaw(rank, pair, gpuBuffer, size, 600, mhandle); + if (!result.ok) return result; + + if (rank != 0) continue; + + void* flushBuffers[1] = {gpuBuffer}; + int flushSizes[1] = {static_cast(size)}; + void* flushHandles[1] = {mhandle}; + void* flushRequest = nullptr; + // A null request means the flush was a no-op; an error + // return means the flush itself failed. + if (FlushRecv(pair.recvComm, 1, flushBuffers, flushSizes, flushHandles, + &flushRequest) + != ncclSuccess) { + result.ok = false; + result.msg = "FlushRecv failed at iteration " + std::to_string(iter); + return result; + } + if (flushRequest != nullptr) { + result = WorkerWait(flushRequest, nullptr); + if (!result.ok) return result; + } + if (!verifyBufferData(gpuBuffer, size, makeBytePattern(seed))) { + result.ok = false; + result.msg = "GPU data verification failed after flush"; + return result; + } + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded FlushRepeated_VNic"); + return; + } + ConnectionPair pair; NetConnectionGuard connGuard(net_); SetupConnectionWithGuard(vdev, pair, connGuard); @@ -1104,6 +1270,7 @@ TEST_F(NetIbMPITest, FlushRepeated_VNic) { ncclResult_t flushResult = FlushRecv(pair.recvComm, 1, flushBuffers, flushSizes, flushHandles, &flushRequest); + EXPECT_EQ(flushResult, ncclSuccess) << "Iter " << iter << ": FlushRecv failed"; if (flushResult == ncclSuccess && flushRequest != nullptr) { ASSERT_EQ(WaitForCompletion(flushRequest, nullptr), ncclSuccess); } @@ -1141,6 +1308,43 @@ TEST_F(NetIbMPITest, SequentialTransfers_VNic) { const int rank = MPIEnvironment::world_rank; + // Parameterized by MPIEnvironment::nThreads: every worker reuses one + // registration across its iterations, so the fused device serves many + // long-lived MRs at once. + if (MPIEnvironment::nThreads > 1) { + static constexpr int kThreadedIters = 100; + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + vdev, MPIEnvironment::nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t size = kSmallBufferSize; + void* buffer = malloc(size); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(comm, buffer, size, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, NetMHandleWorkerDeleter(net_, comm)); + + for (int iter = 0; iter < kThreadedIters; iter++) { + result = WorkerSendRecvPattern(rank, pair, buffer, size, 700, mhandle, + 7000 + threadIdx * 1000 + iter); + if (!result.ok) return result; + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded SequentialTransfers_VNic"); + return; + } + // Single connection through the vNIC, reused across all 100 iterations. ConnectionPair pair; NetConnectionGuard connGuard(net_); diff --git a/projects/rccl/test/transport/NetIbMPI/StressTests.cpp b/projects/rccl/test/transport/NetIbMPI/StressTests.cpp index fcb0072a0603..b7b96e1d1279 100644 --- a/projects/rccl/test/transport/NetIbMPI/StressTests.cpp +++ b/projects/rccl/test/transport/NetIbMPI/StressTests.cpp @@ -48,12 +48,65 @@ TEST_F(NetIbMPITest, InvalidRecvCount) { // E1. MrCacheRefCount — registers the same host buffer twice on the same // comm: cache hit bumps refs to 2, two Dereg calls bring it back to 0 // (covers the refs>0 branch in ncclIbDeregMrInternal). +// Parameterized by MPIEnvironment::nThreads: every worker registers the +// SAME address from its own communicator, so all of them contend on one +// cache entry's refcount rather than merely on the cache mutex. TEST_F(NetIbMPITest, MrCacheRefCount) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)); int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + const size_t sharedSz = 4096; + auto shared = makeHostBufferAutoGuard(malloc(sharedSz)); + ASSERT_NE(shared.get(), nullptr); + + // Every worker must be holding both of its registrations before any of + // them starts releasing. The body gate only synchronizes entry, so + // without this a worker could register twice and release both before the + // next one ran, and the refcount would never climb past the two that the + // serial body already covers. + std::atomic bothHeld{0}; + static constexpr int kRendezvousPolls = 3000; // 3000 * 10ms = 30s + + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + (void)threadIdx; + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mh1 = nullptr; + void* mh2 = nullptr; + ThreadResult result = WorkerRegister(comm, shared.get(), sharedSz, + NCCL_PTR_HOST, &mh1); + if (!result.ok) return result; + result = WorkerRegister(comm, shared.get(), sharedSz, NCCL_PTR_HOST, &mh2); + if (!result.ok) { + DeregisterMemory(comm, mh1); + return result; + } + if (!WorkerRendezvous(bothHeld, nThreads, kRendezvousPolls)) { + DeregisterMemory(comm, mh1); + DeregisterMemory(comm, mh2); + result.ok = false; + result.msg = "only " + std::to_string(bothHeld.load()) + " of " + + std::to_string(nThreads) + + " workers reached the point where all hold two registrations"; + return result; + } + if (DeregisterMemory(comm, mh1) != ncclSuccess + || DeregisterMemory(comm, mh2) != ncclSuccess) { + result.ok = false; + result.msg = "deregMr failed while unwinding the cache refcount"; + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded MrCacheRefCount"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(/*dev=*/0, cp, guard); @@ -168,12 +221,50 @@ TEST_F(NetIbMPITest, NullCommClose) { // E4. TagZeroReuse — 50 messages all sent with tag=0. // Verifies FIFO ordering: messages arrive in send order because // the FIFO is a strict ring (slot = fifoHead % MAX_REQUESTS). +// Parameterized by MPIEnvironment::nThreads: with every worker reusing +// tag 0 on its own connection, a completion routed to the wrong comm shows +// up as a data mismatch instead of passing silently. TEST_F(NetIbMPITest, TagZeroReuse) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)); int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + static constexpr int kTagZeroIters = 50; + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + ThreadDevPolicy::Spread(), nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t sz = kSmallBufferSize; + void* buffer = malloc(sz); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mh = nullptr; + result = WorkerRegister(comm, buffer, sz, NCCL_PTR_HOST, &mh); + if (!result.ok) return result; + NetMHandleWorkerGuard mhGuard(mh, NetMHandleWorkerDeleter(net_, comm)); + + for (int i = 0; i < kTagZeroIters; i++) { + result = WorkerSendRecvPattern(rank, pair, buffer, sz, /*tag=*/0, mh, + threadIdx * 10000 + i); + if (!result.ok) return result; + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded TagZeroReuse"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(/*dev=*/0, cp, guard); @@ -275,12 +366,29 @@ TEST_F(NetIbMPITest, InlineSendBoundary) { // E7. MixedSizeBarrage — 25 sizes from 1B to 64MB on a single connection. // Exercises alignment boundaries, AR threshold, inline paths, and // large-MR registration. +// Parameterized by MPIEnvironment::nThreads: concurrent workers sweep the +// whole ladder at once, so alignment and chunking paths run under real +// bandwidth pressure with several 64 MB registrations live per device. TEST_F(NetIbMPITest, MixedSizeBarrage) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)); int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + static const std::vector kBarrageSizes = { + 1, 7, 13, 64, 127, 128, 129, 255, 256, 512, + 1023, 1024, 4095, 4096, 8191, 8192, 8193, + 16384, 32768, 65536, 131072, + 1048576, 4194304, 16777216, 67108864 + }; + + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + RunThreadedSizeSweep(ThreadDevPolicy::Spread(), nThreads, kBarrageSizes, + /*repeats=*/1, "threaded MixedSizeBarrage"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -322,12 +430,83 @@ TEST_F(NetIbMPITest, MixedSizeBarrage) { // A1. FifoPressureSenderFast — receiver deliberately slow, sender in // tight loop. Verifies that isend returns *request==NULL when the // FIFO slot is not ready (backpressure). +// Parameterized by MPIEnvironment::nThreads: N slow receivers against N +// tight senders replace an idle fabric with genuine queue contention, while +// the backpressure assertion stays per-connection. TEST_F(NetIbMPITest, FifoPressureSenderFast) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)); int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + static constexpr int kThreadedMsgs = 20; + static constexpr int kThreadedRecvDelayUs = 50000; // 50ms + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t sz = kSmallBufferSize; + void* buffer = malloc(sz); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mh = nullptr; + result = WorkerRegister(comm, buffer, sz, NCCL_PTR_HOST, &mh); + if (!result.ok) return result; + NetMHandleWorkerGuard mhGuard(mh, NetMHandleWorkerDeleter(net_, comm)); + + int nullCount = 0; + for (int i = 0; i < kThreadedMsgs; i++) { + void* request = nullptr; + if (rank == 0) { + if (i > 0) usleep(kThreadedRecvDelayUs); + result = WorkerPostRecv(pair.recvComm, buffer, sz, i, mh, &request); + } else { + fillHostBufferWithPattern(buffer, sz, makeBytePattern(i)); + int attempts = 0; + do { + if (PostSend(pair.sendComm, buffer, sz, i, mh, &request) + != ncclSuccess) { + result.ok = false; + result.msg = "isend failed"; + return result; + } + if (request) break; + nullCount++; + usleep(1000); + if (++attempts > kMaxRetryAttempts) { + result.ok = false; + result.msg = "isend never got a FIFO slot"; + return result; + } + } while (!request); + } + if (!result.ok) return result; + + int sizes[1] = {0}; + result = WorkerWait(request, sizes, kStressTimeoutMs); + if (!result.ok) return result; + } + + // The slow receiver guarantees the sender saw backpressure. + if (rank != 0 && nullCount == 0) { + result.ok = false; + result.msg = "expected at least one NULL request (FIFO backpressure)"; + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded FifoPressureSenderFast"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -392,12 +571,91 @@ TEST_F(NetIbMPITest, FifoPressureSenderFast) { // A2. RequestSlotExhaustion — post NCCL_NET_MAX_REQUESTS (32) irecvs, // then drain all via matching sends, then verify slots are recycled. +// Parameterized by MPIEnvironment::nThreads: request pools are per-comm, so +// the point is N x 32 requests in flight on one NIC and the confirmation +// that completion routing never crosses communicators. TEST_F(NetIbMPITest, RequestSlotExhaustion) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)); int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + static constexpr int kMaxReqsPerComm = 32; // NCCL_NET_MAX_REQUESTS + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t sz = 256; + void* buffer = malloc(sz * kMaxReqsPerComm); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mh = nullptr; + result = WorkerRegister(comm, buffer, sz * kMaxReqsPerComm, NCCL_PTR_HOST, &mh); + if (!result.ok) return result; + NetMHandleWorkerGuard mhGuard(mh, NetMHandleWorkerDeleter(net_, comm)); + + // No rank handshake here: the sender's retry loop waits for the + // receiver's FIFO slots instead of a barrier the worker cannot use. + // Per-worker seeds, checked on arrival: with the slot index alone + // every worker would send the same bytes for a slot, and a + // completion delivered to the wrong worker's request would satisfy + // the wait unnoticed. + std::vector requests(kMaxReqsPerComm, nullptr); + for (int i = 0; i < kMaxReqsPerComm; i++) { + char* slot = static_cast(buffer) + i * sz; + if (rank == 0) { + memset(slot, 0, sz); + result = WorkerPostRecv(pair.recvComm, slot, sz, i, mh, &requests[i]); + } else { + fillHostBufferWithPattern( + slot, sz, makeBytePattern(WorkerSeed(threadIdx, i))); + result = WorkerPostSend(pair.sendComm, slot, sz, i, mh, &requests[i]); + } + if (!result.ok) return result; + } + for (int i = 0; i < kMaxReqsPerComm; i++) { + int sizes[1] = {0}; + result = WorkerWait(requests[i], sizes, kStressTimeoutMs); + if (!result.ok) return result; + if (rank != 0) continue; + const char* slot = static_cast(buffer) + i * sz; + if (sizes[0] != (int)sz) { + result.ok = false; + result.msg = "slot " + std::to_string(i) + " completed with size " + + std::to_string(sizes[0]) + " instead of " + + std::to_string(sz); + return result; + } + if (!verifyHostBufferData(slot, sz, + makeBytePattern(WorkerSeed(threadIdx, i)))) { + result.ok = false; + result.msg = "slot " + std::to_string(i) + + " holds another worker's payload"; + return result; + } + } + + // Slots must be recycled: another round on the same comm. + for (int i = 0; i < 4; i++) { + result = WorkerSendRecvPattern(rank, pair, buffer, sz, 100 + i, mh, + threadIdx * 10000 + i, kStressTimeoutMs); + if (!result.ok) return result; + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded RequestSlotExhaustion"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -455,12 +713,89 @@ TEST_F(NetIbMPITest, RequestSlotExhaustion) { // A3. MemoryRegistrationStorm — register/deregister 512 unique buffers // in various patterns to stress the MR cache. +// Parameterized by MPIEnvironment::nThreads: the MR cache is per physical +// device behind a single mutex, so concurrent storms are the only way to +// race its insert, lookup and eviction paths against each other. TEST_F(NetIbMPITest, MemoryRegistrationStorm) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)); int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + // Keep the aggregate registration count near the serial test's 512 rather + // than multiplying it by the worker count. What is under test is + // contention on the one per-device cache, and a per-worker storm that + // grows with N only spreads the workers apart in time until the + // verification transfer outlives any sane timeout. + const int kThreadedBufs = std::max(16, 512 / nThreads); + static constexpr size_t kThreadedBufSz = 4096; + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + + std::vector bufs(kThreadedBufs, nullptr); + auto bufCleanup = makeScopeGuard([&]() { + for (auto* p : bufs) free(p); + }); + for (int i = 0; i < kThreadedBufs; i++) { + bufs[i] = malloc(kThreadedBufSz); + if (!bufs[i]) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + } + + std::vector handles(kThreadedBufs, nullptr); + auto registerAll = [&]() { + for (int i = 0; i < kThreadedBufs; i++) { + result = WorkerRegister(comm, bufs[i], kThreadedBufSz, NCCL_PTR_HOST, + &handles[i]); + if (!result.ok) return false; + } + return true; + }; + auto deregisterOrder = [&](const std::vector& order) { + for (int idx : order) { + if (DeregisterMemory(comm, handles[idx]) != ncclSuccess) { + result.ok = false; + result.msg = "deregMr failed"; + return false; + } + handles[idx] = nullptr; + } + return true; + }; + + std::vector reverse(kThreadedBufs); + for (int i = 0; i < kThreadedBufs; i++) reverse[i] = kThreadedBufs - 1 - i; + std::vector shuffled(kThreadedBufs); + for (int i = 0; i < kThreadedBufs; i++) shuffled[i] = i; + for (int i = kThreadedBufs - 1; i > 0; i--) { + int j = (i * 42 + 17 + threadIdx) % (i + 1); + std::swap(shuffled[i], shuffled[j]); + } + + if (!registerAll()) return result; + if (!deregisterOrder(reverse)) return result; + if (!registerAll()) return result; + if (!deregisterOrder(shuffled)) return result; + + // The connection must still work after the storm. The peer's + // worker may be several hundred registrations behind, so this + // transfer waits on the stress budget rather than the default. + return WorkerHostTransfer(rank, pair, kThreadedBufSz, 0, + 999 + threadIdx * 10000, kStressTimeoutMs); + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded MemoryRegistrationStorm"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -888,6 +1223,16 @@ TEST_F(NetIbMPITest, MultiQpSplitDataStress) { int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads: concurrent workers multiply the + // QPs in flight per device while each one checks its own split payloads. + if (MPIEnvironment::nThreads > 1) { + RunThreadedSizeSweep(ThreadDevPolicy::Spread(), MPIEnvironment::nThreads, + {127, 128, 129, 255, 256, 257, 1023, 1024, 1025, + 4095, 4096, 4097, 65535, 65536, 65537}, + /*repeats=*/2, "threaded MultiQpSplitDataStress"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -933,6 +1278,15 @@ TEST_F(NetIbMPITest, MultiQpNoSplitStress) { int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + // Parameterized by MPIEnvironment::nThreads: same sweep on the nDataQps path. + if (MPIEnvironment::nThreads > 1) { + RunThreadedSizeSweep(ThreadDevPolicy::Spread(), MPIEnvironment::nThreads, + {127, 128, 129, 255, 256, 257, 1023, 1024, 1025, + 4095, 4096, 4097, 65535, 65536, 65537}, + /*repeats=*/2, "threaded MultiQpNoSplitStress"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -1404,12 +1758,50 @@ TEST_F(NetIbMPITest, BidirectionalSaturation) { // F1. LongRunningEndurance — 10000 transfers on a single connection, // tag cycling, RDMA checkpoints. +// Parameterized by MPIEnvironment::nThreads. The leak checkpoints move to +// the main thread around the whole run: an in-loop snapshot would see the +// other workers' live QPs and report a false leak. TEST_F(NetIbMPITest, LongRunningEndurance) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)); int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + static constexpr int kThreadedIters = 2000; + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + ThreadDevPolicy::Spread(), nThreads, + [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t sz = kSmallBufferSize; + void* buffer = malloc(sz); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mh = nullptr; + result = WorkerRegister(comm, buffer, sz, NCCL_PTR_HOST, &mh); + if (!result.ok) return result; + NetMHandleWorkerGuard mhGuard(mh, NetMHandleWorkerDeleter(net_, comm)); + + for (int i = 0; i < kThreadedIters; i++) { + result = WorkerSendRecvPattern(rank, pair, buffer, sz, i % 1000, mh, + WorkerSeed(threadIdx, i), kStressTimeoutMs); + if (!result.ok) return result; + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded LongRunningEndurance"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -1458,6 +1850,80 @@ TEST_F(NetIbMPITest, GpuMemoryTransferStress) { GTEST_SKIP() << "No GPU Direct RDMA support"; } + // Parameterized by MPIEnvironment::nThreads: concurrent GPU registrations + // drive the peer-memory path into the same per-device MR cache. Workers + // inherit this rank's HIP device from the harness, since the current device + // is thread-local. + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + static constexpr int kThreadedCycles = 3; + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t cycleSizes[kThreadedCycles] = {512 * 1024, 1024 * 1024, + 2 * 1024 * 1024}; + for (int cycle = 0; cycle < kThreadedCycles; cycle++) { + const size_t sz = cycleSizes[cycle]; + const int seed = threadIdx * 100 + cycle; + + void* devBuf = nullptr; + if (hipMalloc(&devBuf, sz) != hipSuccess) { + result.ok = false; + result.msg = "hipMalloc failed"; + return result; + } + auto devGuard = makeDeviceBufferAutoGuard(devBuf); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mh = nullptr; + result = WorkerRegister(comm, devBuf, sz, NCCL_PTR_CUDA, &mh); + if (!result.ok) return result; + NetMHandleWorkerGuard mhGuard(mh, NetMHandleWorkerDeleter(net_, comm)); + + if (rank == 1 + && initializeBufferWithPattern(devBuf, sz, + makeBytePattern(seed)) + != hipSuccess) { + result.ok = false; + result.msg = "GPU buffer initialization failed"; + return result; + } + + result = WorkerSendRecvRaw(rank, pair, devBuf, sz, cycle, mh, + kLargeTransferTimeoutMs); + if (!result.ok) return result; + + if (rank == 0) { + void* flushRequest = nullptr; + int flushSize = static_cast(sz); + // A null request means the flush was a no-op; an error + // return means the flush itself failed. + if (FlushRecv(pair.recvComm, 1, &devBuf, &flushSize, &mh, &flushRequest) + != ncclSuccess) { + result.ok = false; + result.msg = "FlushRecv failed at cycle " + std::to_string(cycle); + return result; + } + if (flushRequest) { + int flushed[1] = {0}; + result = WorkerWait(flushRequest, flushed, kLargeTransferTimeoutMs); + if (!result.ok) return result; + } + if (!verifyBufferData(devBuf, sz, makeBytePattern(seed))) { + result.ok = false; + result.msg = "GPU data mismatch"; + return result; + } + } + } + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded GpuMemoryTransferStress"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -1499,6 +1965,7 @@ TEST_F(NetIbMPITest, GpuMemoryTransferStress) { void* flushReq = nullptr; int flushSz = static_cast(sz); ncclResult_t fr = FlushRecv(cp.recvComm, 1, &devBuf, &flushSz, &mh, &flushReq); + EXPECT_EQ(fr, ncclSuccess) << "FlushRecv failed at cycle " << cycle; if (fr == ncclSuccess && flushReq) { int fsz = 0; EXPECT_EQ(WaitForCompletion(flushReq, &fsz, kLargeTransferTimeoutMs), ncclSuccess); @@ -1516,12 +1983,65 @@ TEST_F(NetIbMPITest, GpuMemoryTransferStress) { // ===================================================================== // J1. RapidRecvPostDrain — 100 cycles of post-32-recv → send-32 → drain. +// Parameterized by MPIEnvironment::nThreads: batches from all workers are +// in flight on one NIC at once, so recycling has to hold up while other +// communicators are draining their own batches. TEST_F(NetIbMPITest, RapidRecvPostDrain) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, false, kMinGpusPerNode, kNoNodeLimit)); int rank = MPIEnvironment::world_rank; AssertInitAndGetDevices(nullptr); + const int nThreads = MPIEnvironment::nThreads; + if (nThreads > 1) { + static constexpr int kThreadedCycles = 25; + static constexpr int kThreadedBatch = 32; + const RdmaResourceCounts before = CaptureRdmaResources(); + RunMultiThreadedIndependent( + 0, nThreads, [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + const size_t sz = 256; + void* buffer = malloc(sz * kThreadedBatch); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mh = nullptr; + result = WorkerRegister(comm, buffer, sz * kThreadedBatch, NCCL_PTR_HOST, &mh); + if (!result.ok) return result; + NetMHandleWorkerGuard mhGuard(mh, NetMHandleWorkerDeleter(net_, comm)); + + for (int cycle = 0; cycle < kThreadedCycles; cycle++) { + std::vector requests(kThreadedBatch, nullptr); + for (int i = 0; i < kThreadedBatch; i++) { + char* slot = static_cast(buffer) + i * sz; + if (rank == 0) { + result = WorkerPostRecv(pair.recvComm, slot, sz, i, mh, &requests[i]); + } else { + fillHostBufferWithPattern( + slot, sz, makeBytePattern(cycle * kThreadedBatch + i)); + result = WorkerPostSend(pair.sendComm, slot, sz, i, mh, &requests[i]); + } + if (!result.ok) return result; + } + for (int i = 0; i < kThreadedBatch; i++) { + int sizes[1] = {0}; + result = WorkerWait(requests[i], sizes, kStressTimeoutMs); + if (!result.ok) return result; + } + } + (void)threadIdx; + return result; + }); + MPI_Barrier(MPI_COMM_WORLD); + AssertNoRdmaLeaks(before, CaptureRdmaResources(), "threaded RapidRecvPostDrain"); + return; + } + ConnectionPair cp; NetConnectionGuard guard(net_); SetupConnectionWithGuard(0, cp, guard); @@ -1586,5 +2106,169 @@ TEST_F(NetIbMPITest, SetNetAttrNoOp) { MPI_Barrier(MPI_COMM_WORLD); } +// ============================================================================= +// Test: ThreadedProgressDoesNotSerialize +// +// Multithread-only gate on the data path itself: independent workers must make +// progress at the same time, not one after another. It exists because a single +// process-global lock anywhere under isend/irecv/test is enough to turn every +// concurrent communicator into a queue while every functional test still passes. +// The known instance is the libibverbs ops shims, whose poll_cq takes one +// process-global mutex per call, but nothing about this test is specific to +// them: any future global lock on the progress path shows up here. +// +// Method. The same body runs in one process, on the same devices: once with a +// single worker, then with all of them, and each worker times only its own loop, +// so connection setup is outside the measurement. Every worker does the same +// fixed number of transfers, so the parallel phase moves N times the data. +// Perfect overlap therefore lands near the single-worker time, and full +// serialization near N times it. Calibration runs before and after the parallel +// phase and is averaged, so a machine that drifts mid-test is visible rather +// than silently shifting the verdict. A fourth phase holds a mutex across every +// transfer, and the same measurement has to notice it -- otherwise the verdict +// on the plugin means nothing. +// +// What it does and does not see. A lock held across a progress call serializes +// the call itself and shows up here immediately. A lock taken and released +// around one, like the two acquisitions shimPollCq adds per poll_cq, does not: +// measured at four and sixteen workers, with the ops shims installed and +// without, the factor was 1.00 either way, because the critical section is a map +// lookup next to a driver call. That leak is caught deterministically instead, by +// FaultInjectionShimsAbsentUnlessRequested. +// ============================================================================= +TEST_F(NetIbMPITest, ThreadedProgressDoesNotSerialize) { + ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, + false, kMinGpusPerNode, kNoNodeLimit)) + << "Test requires exactly " << kExactTwoProcesses << " processes"; + + const int rank = MPIEnvironment::world_rank; + const int nThreads = MPIEnvironment::nThreads; + // Below four workers the ratio is too close to one to separate contention + // from ordinary noise. + if (nThreads < 4) { + GTEST_SKIP() << "requires --net_ib_nthreads of at least 4 to separate contention " + "from noise"; + } + + AssertInitAndGetDevices(nullptr); + + // Small messages: the point is how many transfers overlap, not bandwidth. + static constexpr size_t kMsgSize = 256; + static constexpr int kIters = 100; + static constexpr int kTimeoutMs = 60000; + // Full serialization costs N times the single-worker time. Passing means + // staying below 60% of that, which on healthy hardware measures near 1. + static constexpr double kSerializedFraction = 0.6; + + std::vector workerMs(nThreads, 0.0); + // Phase 3 holds this across each transfer to imitate a global lock on the + // progress path, which is how the measurement proves it can still see one. + std::mutex serializeAll; + bool serialize = false; + + auto body = [&](int threadIdx, ConnectionPair& pair) -> ThreadResult { + ThreadResult result; + void* buffer = malloc(kMsgSize); + if (!buffer) { + result.ok = false; + result.msg = "malloc failed"; + return result; + } + auto bufferGuard = makeHostBufferAutoGuard(buffer); + memset(buffer, 0x5A, kMsgSize); + + void* comm = (rank == 0) ? pair.recvComm : pair.sendComm; + void* mhandle = nullptr; + result = WorkerRegister(comm, buffer, kMsgSize, NCCL_PTR_HOST, &mhandle); + if (!result.ok) return result; + NetMHandleWorkerGuard mhandleGuard(mhandle, NetMHandleWorkerDeleter(net_, comm)); + + const auto start = std::chrono::steady_clock::now(); + for (int i = 0; i < kIters; i++) { + // Only the sending rank serializes. Locking both would let the two + // ranks admit different workers at the same time, and each would then + // wait for a peer that its own lock is keeping out -- a deadlock, + // which is its own argument against a global lock on this path. + if (serialize && rank == 1) { + std::lock_guard lock(serializeAll); + result = WorkerSendRecvRaw(rank, pair, buffer, kMsgSize, 950 + (i % 32), mhandle, + kTimeoutMs); + } else { + result = WorkerSendRecvRaw(rank, pair, buffer, kMsgSize, 950 + (i % 32), mhandle, + kTimeoutMs); + } + if (!result.ok) return result; + } + const auto end = std::chrono::steady_clock::now(); + workerMs[threadIdx] = + std::chrono::duration(end - start).count(); + return result; + }; + + // The slowest worker bounds the phase: the others waited for it. + auto measure = [&](int workers) -> double { + std::fill(workerMs.begin(), workerMs.end(), 0.0); + RunMultiThreadedIndependent(0, workers, body); + double worst = 0.0; + for (int t = 0; t < workers; t++) worst = std::max(worst, workerMs[t]); + return worst; + }; + + const double singleBefore = measure(1); + const double parallel = measure(nThreads); + const double singleAfter = measure(1); + serialize = true; + const double serialized = measure(nThreads); + serialize = false; + + // A worker can fail on one rank only, and the harness reports that on both + // ranks already. What must not happen here is a fatal assertion: this rank + // would leave the test while the peer, whose own timings are fine, waits in + // the barrier below forever. So a missing measurement is reported and the + // verdict skipped, and both ranks reach the barrier either way. + const bool measured = singleBefore > 0.0 && singleAfter > 0.0 && parallel > 0.0 + && serialized > 0.0; + EXPECT_TRUE(measured) + << "a phase produced no timing on this rank (single " << singleBefore << "/" + << singleAfter << " ms, parallel " << parallel << " ms, serialized " << serialized + << " ms), so the scaling verdict is skipped"; + + if (measured) { + const double single = (singleBefore + singleAfter) / 2.0; + const double factor = parallel / single; + const double budget = kSerializedFraction * nThreads; + const double serializedFactor = serialized / single; + + TEST_INFO("progress scaling: %d workers, single %.1f/%.1f ms, parallel %.1f ms, " + "factor %.2f; deliberately serialized %.1f ms, factor %.2f; budget %.2f", + nThreads, singleBefore, singleAfter, parallel, factor, serialized, + serializedFactor, budget); + + // Without this the gate could pass by being blind: a wait that sleeps, or a + // workload that spends its time off the progress path, would report a flat + // factor no matter what. The same measurement has to flag a lock it knows is + // there before its verdict on the plugin means anything. + EXPECT_GT(serializedFactor, budget) + << "the measurement did not notice a mutex held across every transfer (factor " + << serializedFactor << ", budget " << budget + << "), so it cannot be trusted to notice one inside the plugin either"; + + EXPECT_LT(factor, budget) + << nThreads << " workers each moved " << kIters << " messages in " << parallel + << " ms while one worker needed " << single << " ms, a factor of " << factor + << " where full serialization is " << nThreads << ". Concurrent progress is being " + << "serialized somewhere under isend/irecv/test: a lock held across a progress " + << "call would do exactly this."; + + // A machine that changed speed under us would invalidate the ratio above. + const double drift = + std::max(singleBefore, singleAfter) / std::min(singleBefore, singleAfter); + EXPECT_LT(drift, 1.5) << "single-worker calibration drifted between " << singleBefore + << " ms and " << singleAfter + << " ms, so the scaling factor cannot be trusted"; + } + + MPI_Barrier(MPI_COMM_WORLD); +} #endif /* MPI_TESTS_ENABLED */ From f8a3cfdb92e0a2c074f0ee6f2d4a67982dd6fd45 Mon Sep 17 00:00:00 2001 From: Ilya Kosarev Date: Thu, 20 Aug 2026 17:49:32 +0000 Subject: [PATCH 3/7] net-ib: wire the threaded matrices into the runner, and unbreak what they land in The matrix grows from 352 to 445 executions and gains IB-CAST and fault-injection matrices, wired into CI through ainic.json A17-A19 and a four-worker suite in each of the six per-cluster functional configs. Every threaded suite hangs off a base that declares no tests. extends merges test lists rather than replacing them (lib/test_config.py, _merge_tests), so a suite that extended one carrying tests would also run those at its own worker count -- and where the parent was fault_inject_cast, would run the six ops-overload tests with the fault injection they need switched off. Two settings need justification. RCCL_IB_QP_SCHED_UPDATE_INTERVAL of 10 s, which four CAST toggle tests require so the RTT timer cannot rewrite tokens mid-assertion; every shipped config supplied 50 us, so those four had always skipped and never run. And RCCL_IB_FAULT_INJECTION off in the threaded fault suites, which use the per-communicator hooks: the shims add process-global mutex traffic to every poll_cq and their registry is per device, so one worker's clear disarms every other worker's faults. Two pre-existing problems in the files these suites land in, both of which made them inert or dishonest: - The schema allowed num_gpus of 1 or more, while nccl4py_build_smoke declares 0 because it is a CPU-only build smoke. Validation is per file and throws, so that one value stopped five of the six per-cluster configs from loading at all. Zero is allowed now, and the executor rejects it where it cannot place ranks. - 54 gtest filters spelled fixtures from RegistrationMPITests.cpp as AlltoAll where the tests are AllToAll. A filter that matches nothing runs zero tests and exits successfully, so those suites reported success having run nothing. Only these were touched; AlltoAll.*, CeMPI_AlltoAll.* and the Dda and Ce eligibility fixtures are real names elsewhere. Co-authored-by: Cursor --- .../scripts/test_runner/configs/ainic.json | 73 ++++ .../configs/mi300_mellanox_ib_func.json | 123 +++++- .../configs/mi300x_mellanox_ib.json | 123 +++++- .../configs/mi355x_thor2_roce.json | 123 +++++- .../configs/mi355x_thor2_roce_func.json | 123 +++++- .../test_runner/configs/mi455_ainic_roce.json | 123 +++++- .../configs/mi455_ainic_roce_func.json | 123 +++++- .../test_runner/configs/net_ib_transport.json | 391 +++++++++++++++++- .../tools/scripts/test_runner/lib/schema.json | 4 +- .../scripts/test_runner/lib/test_executor.py | 15 + 10 files changed, 1157 insertions(+), 64 deletions(-) diff --git a/projects/rccl/tools/scripts/test_runner/configs/ainic.json b/projects/rccl/tools/scripts/test_runner/configs/ainic.json index 005f47d46c1a..632b2c053fa8 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/ainic.json +++ b/projects/rccl/tools/scripts/test_runner/configs/ainic.json @@ -155,6 +155,76 @@ } ] }, + "a_netib_multithread": { + "extends": "a_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "description": "A17: four worker threads per rank, each on its own NET/IB context and connection. Mirrors how a framework process with several communicators drives the plugin from several proxy threads.", + "tests": [ + { + "name": "A17_NetIB_MultiThread_SimpleSendRecv", + "description": "Concurrent workers each transfer on an independent connection.", + "test_filter": "NetIbMPITest.SimpleSendRecv" + }, + { + "name": "A17_NetIB_MultiThread_SequentialTransfers", + "description": "Concurrent workers each run repeated transfers on an independent connection.", + "test_filter": "NetIbMPITest.MultipleSequentialTransfers" + }, + { + "name": "A17_NetIB_MultiThread_MemoryRegistrationStorm", + "description": "Concurrent register/deregister storms contend on the shared per-device MR cache.", + "test_filter": "NetIbMPITest.MemoryRegistrationStorm" + }, + { + "name": "A17_NetIB_MultiThread_MrCacheRefCount", + "description": "All workers register the same address, contending on one cache entry's refcount.", + "test_filter": "NetIbMPITest.MrCacheRefCount" + }, + { + "name": "A17_NetIB_MultiThread_ProgressDoesNotSerialize", + "description": "Gate: concurrent workers must progress at the same time. Catches a lock held across a progress call, which functional tests pass straight through.", + "test_filter": "NetIbMPITest.ThreadedProgressDoesNotSerialize" + } + ] + }, + "a_cast_multithread": { + "extends": "a_cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "description": "A18: four worker threads per rank on independent IB-CAST connections, so several WRR schedulers advance at once.", + "tests": [ + { + "name": "A18_Cast_MultiThread_StressMultiRound", + "description": "Concurrent CAST connections, each auditing its own WRR token ledger.", + "test_filter": "NetIbMPITest.CastStressMultiRoundTwoConns" + } + ] + }, + "a_fault_inject_multithread": { + "extends": "a_cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "RCCL_IB_FAULT_INJECTION": "0" + }, + "description": "A19: fault isolation between workers. One worker breaks its own connection while the others must keep transferring cleanly. Fault injection stays off: the tests use the per-comm hooks, and the libibverbs shims add process-global mutex traffic to every poll_cq while their per-device registry lets one worker's clear disarm the others.", + "tests": [ + { + "name": "A19_Cast_FaultIsolationAcrossWorkers", + "description": "Injected fault state stays inside its communicator; bystander connections stay clean.", + "test_filter": "NetIbMPITest.FaultIsolationAcrossWorkers" + }, + { + "name": "A19_Cast_FaultShimsAbsentUnlessRequested", + "description": "The libibverbs ops shims must not be installed while RCCL_IB_FAULT_INJECTION is off: their poll_cq serializes the whole process on one mutex.", + "test_filter": "NetIbMPITest.FaultInjectionShimsAbsentUnlessRequested" + } + ] + }, "a_gin": { "extends": "a_base", @@ -448,6 +518,9 @@ { "name": "A14. GIN enable gate (ROCM-27070)", "config": "a_gin_enable_gate", "enabled": true }, { "name": "A15. GIN AINIC backend (ROCM-27070)", "config": "a_gin_ainic_backend", "enabled": true }, { "name": "A16. Component - Bootstrap socket poll", "config": "a_socket_poll_bootstrap", "enabled": true }, + { "name": "A17. Component - NetIb multithread", "config": "a_netib_multithread", "enabled": true }, + { "name": "A18. Component - CAST multithread", "config": "a_cast_multithread", "enabled": true }, + { "name": "A19. Component - Fault isolation across workers", "config": "a_fault_inject_multithread", "enabled": true }, { "name": "B1. IB baseline", "config": "b1_ib", "enabled": true }, { "name": "B2. CAST alltoall full (1K-16G)", "config": "b_alltoall_cast_full", "enabled": true }, diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300_mellanox_ib_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi300_mellanox_ib_func.json index 655b50d940c0..31e59e4c9264 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi300_mellanox_ib_func.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi300_mellanox_ib_func.json @@ -190,6 +190,91 @@ "NCCL_DMABUF_ENABLE": "1" } }, + "net_ib_multithread_independent": { + "extends": "net_ib_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "tests": [ + { + "name": "NetIB_MultiThread_IndependentSimpleSendRecv", + "description": "Concurrent workers each transfer on an independent NET/IB context and connection", + "test_filter": "NetIbMPITest.SimpleSendRecv" + }, + { + "name": "NetIB_MultiThread_IndependentSequentialTransfers", + "description": "Concurrent workers each run repeated transfers on an independent connection", + "test_filter": "NetIbMPITest.MultipleSequentialTransfers" + }, + { + "name": "NetIB_MultiThread_MultipleSizes", + "description": "Concurrent workers sweep the size ladder into the shared per-device MR cache", + "test_filter": "NetIbMPITest.SendRecvMultipleSizes" + }, + { + "name": "NetIB_MultiThread_MemoryRegistrationStorm", + "description": "Concurrent register/deregister storms race the per-device MR cache", + "test_filter": "NetIbMPITest.MemoryRegistrationStorm" + }, + { + "name": "NetIB_MultiThread_MrCacheRefCount", + "description": "All workers register the same address, contending on one cache entry's refcount", + "test_filter": "NetIbMPITest.MrCacheRefCount" + }, + { + "name": "NetIB_MultiThread_LongRunningEndurance", + "description": "Sustained concurrent traffic with the RDMA leak check around the whole run", + "test_filter": "NetIbMPITest.LongRunningEndurance" + } + ] + }, + "cast_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_MultiThread_StressMultiRound", + "description": "Concurrent IB-CAST connections, each auditing its own WRR token ledger", + "test_filter": "NetIbMPITest.CastStressMultiRoundTwoConns" + }, + { + "name": "Cast_MultiThread_SendRecvMultipleSizes", + "description": "Concurrent schedulers sweep both sides of the WRR/split boundary", + "test_filter": "NetIbMPITest.CastSendRecvMultipleSizes" + } + ] + }, + "fault_inject_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "NCCL_IB_QPS_PER_CONNECTION": "4", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "100", + "RCCL_IB_QP_SCHED_RESET_INTERVAL": "1000", + "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN": "2000", + "NCCL_IB_RETURN_ASYNC_EVENTS": "1", + "RCCL_IB_FAULT_INJECTION": "0" + }, + "tests": [ + { + "name": "FaultInj_MultiThread_IsolationAcrossWorkers", + "description": "One worker breaks its own connection while the others keep transferring cleanly", + "test_filter": "NetIbMPITest.FaultIsolationAcrossWorkers" + }, + { + "name": "FaultInj_MultiThread_QpErrorClearRecovers", + "description": "Each worker breaks one connection, clears the fault and must move data on a second one", + "test_filter": "NetIbMPITest.FaultInjCastQpErrorClearRecovers" + } + ] + }, "net_ib_gdr_flush": { "extends": "net_ib_base", "timeout": 120, @@ -902,7 +987,7 @@ { "name": "NetIB_Stress_AlltoAll", "description": "Full mesh, 50 iterations", - "test_filter": "NetIbMPITest.AlltoAllStress" + "test_filter": "NetIbMPITest.AllToAllStress" }, { "name": "NetIB_Stress_BidirectionalMultiRank", @@ -1975,7 +2060,7 @@ { "name": "UBR_AlltoAll_OutOfPlace_MultiNode", "description": "AlltoAll out-of-place with UBR on multi-node", - "test_filter": "UBR_AlltoAll.OutOfPlace_MultiNode" + "test_filter": "UBR_AllToAll.OutOfPlace_MultiNode" }, { "name": "UBR_SendRecv_RingPattern_MultiNode", @@ -1985,19 +2070,19 @@ { "name": "UBR_AlltoAllStress_InterleavedMultiSize", "description": "Stress test with interleaved AlltoAll/AlltoAllv across small/medium/large sizes", - "test_filter": "UBR_AlltoAllStress.InterleavedMultiSize_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedMultiSize_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_SplitComm", "description": "Stress test with interleaved AlltoAll/AlltoAllv using split communicators", - "test_filter": "UBR_AlltoAllStress.InterleavedWithSplitComm_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedWithSplitComm_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_ConcurrentMultiComm", "description": "Concurrent AlltoAll on multiple communicators with separate streams", - "test_filter": "UBR_AlltoAllStress.ConcurrentMultiComm_MultiNode", + "test_filter": "UBR_AllToAllStress.ConcurrentMultiComm_MultiNode", "timeout": 300 }, { @@ -2008,12 +2093,12 @@ { "name": "UBR_ConcurrentRegHang_AlltoAll_Contiguous", "description": "ncclAlltoAll with contiguous registered buffers (should pass)", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_ConcurrentRegHang_AlltoAll_Stress", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2059,12 +2144,12 @@ { "name": "UBR_AlltoAll_Contiguous_Success", "description": "ncclAlltoAll with contiguous registered buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_AlltoAll_Stress_Success", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2105,7 +2190,7 @@ { "name": "GraphCapture_AlltoAll_MultiNode", "description": "AlltoAll graph capture with IPC registration on multi-node", - "test_filter": "GraphCapture_AlltoAll.MultiNode" + "test_filter": "GraphCapture_AllToAll.MultiNode" }, { "name": "GraphCapture_AllReduce_MultiNode", @@ -2751,6 +2836,24 @@ "config": "net_transport_eth_multinode", "enabled": true }, + { + "name": "NET IB - Multithread Independent Connections (4)", + "description": "4 worker threads/rank, each on an independent NET/IB context and connection", + "config": "net_ib_multithread_independent", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (4)", + "description": "4 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread", + "enabled": true + }, + { + "name": "NET IB - Fault Injection Multithread (4)", + "description": "Per-communicator fault injection across concurrent workers, including cross-worker isolation", + "config": "fault_inject_multithread", + "enabled": true + }, { "name": "GDR Flush - Whitebox ncclIbIflush (Single-Node, 2-Rank)", "description": "Whitebox coverage of the NET/IB GDR flush fix: dma-buf and peermem RO=0 scratchpad reads, feature-disabled fallback, repeated-flush burst, and the fault-injection regression guard", 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 8404a65af396..791f32246bdd 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 @@ -446,6 +446,91 @@ "NCCL_DMABUF_ENABLE": "1" } }, + "net_ib_multithread_independent": { + "extends": "net_ib_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "tests": [ + { + "name": "NetIB_MultiThread_IndependentSimpleSendRecv", + "description": "Concurrent workers each transfer on an independent NET/IB context and connection", + "test_filter": "NetIbMPITest.SimpleSendRecv" + }, + { + "name": "NetIB_MultiThread_IndependentSequentialTransfers", + "description": "Concurrent workers each run repeated transfers on an independent connection", + "test_filter": "NetIbMPITest.MultipleSequentialTransfers" + }, + { + "name": "NetIB_MultiThread_MultipleSizes", + "description": "Concurrent workers sweep the size ladder into the shared per-device MR cache", + "test_filter": "NetIbMPITest.SendRecvMultipleSizes" + }, + { + "name": "NetIB_MultiThread_MemoryRegistrationStorm", + "description": "Concurrent register/deregister storms race the per-device MR cache", + "test_filter": "NetIbMPITest.MemoryRegistrationStorm" + }, + { + "name": "NetIB_MultiThread_MrCacheRefCount", + "description": "All workers register the same address, contending on one cache entry's refcount", + "test_filter": "NetIbMPITest.MrCacheRefCount" + }, + { + "name": "NetIB_MultiThread_LongRunningEndurance", + "description": "Sustained concurrent traffic with the RDMA leak check around the whole run", + "test_filter": "NetIbMPITest.LongRunningEndurance" + } + ] + }, + "cast_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_MultiThread_StressMultiRound", + "description": "Concurrent IB-CAST connections, each auditing its own WRR token ledger", + "test_filter": "NetIbMPITest.CastStressMultiRoundTwoConns" + }, + { + "name": "Cast_MultiThread_SendRecvMultipleSizes", + "description": "Concurrent schedulers sweep both sides of the WRR/split boundary", + "test_filter": "NetIbMPITest.CastSendRecvMultipleSizes" + } + ] + }, + "fault_inject_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "NCCL_IB_QPS_PER_CONNECTION": "4", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "100", + "RCCL_IB_QP_SCHED_RESET_INTERVAL": "1000", + "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN": "2000", + "NCCL_IB_RETURN_ASYNC_EVENTS": "1", + "RCCL_IB_FAULT_INJECTION": "0" + }, + "tests": [ + { + "name": "FaultInj_MultiThread_IsolationAcrossWorkers", + "description": "One worker breaks its own connection while the others keep transferring cleanly", + "test_filter": "NetIbMPITest.FaultIsolationAcrossWorkers" + }, + { + "name": "FaultInj_MultiThread_QpErrorClearRecovers", + "description": "Each worker breaks one connection, clears the fault and must move data on a second one", + "test_filter": "NetIbMPITest.FaultInjCastQpErrorClearRecovers" + } + ] + }, "net_ib_initialization": { "extends": "net_ib_base", "timeout": 60, @@ -1334,7 +1419,7 @@ { "name": "UBR_AlltoAll_OutOfPlace_MultiNode", "description": "AlltoAll out-of-place with UBR on multi-node", - "test_filter": "UBR_AlltoAll.OutOfPlace_MultiNode" + "test_filter": "UBR_AllToAll.OutOfPlace_MultiNode" }, { "name": "UBR_SendRecv_RingPattern_MultiNode", @@ -1344,19 +1429,19 @@ { "name": "UBR_AlltoAllStress_InterleavedMultiSize", "description": "Stress test with interleaved AlltoAll/AlltoAllv across small/medium/large sizes", - "test_filter": "UBR_AlltoAllStress.InterleavedMultiSize_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedMultiSize_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_SplitComm", "description": "Stress test with interleaved AlltoAll/AlltoAllv using split communicators", - "test_filter": "UBR_AlltoAllStress.InterleavedWithSplitComm_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedWithSplitComm_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_ConcurrentMultiComm", "description": "Concurrent AlltoAll on multiple communicators with separate streams", - "test_filter": "UBR_AlltoAllStress.ConcurrentMultiComm_MultiNode", + "test_filter": "UBR_AllToAllStress.ConcurrentMultiComm_MultiNode", "timeout": 300 }, { @@ -1367,12 +1452,12 @@ { "name": "UBR_ConcurrentRegHang_AlltoAll_Contiguous", "description": "ncclAlltoAll with contiguous registered buffers (should pass)", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_ConcurrentRegHang_AlltoAll_Stress", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -1455,12 +1540,12 @@ { "name": "UBR_AlltoAll_Contiguous_Success", "description": "ncclAlltoAll with contiguous registered buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_AlltoAll_Stress_Success", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -1855,7 +1940,7 @@ { "name": "GraphCapture_AlltoAll_MultiNode", "description": "AlltoAll graph capture with IPC registration on multi-node", - "test_filter": "GraphCapture_AlltoAll.MultiNode" + "test_filter": "GraphCapture_AllToAll.MultiNode" }, { "name": "GraphCapture_AllReduce_MultiNode", @@ -1988,7 +2073,7 @@ "tests": [ { "name": "NetIB_Stress_FanIn", "description": "Fan-in 3->1, 100 iterations", "test_filter": "NetIbMPITest.FanInStress" }, { "name": "NetIB_Stress_FanOut", "description": "Fan-out 1->3, 100 iterations", "test_filter": "NetIbMPITest.FanOutStress" }, - { "name": "NetIB_Stress_AlltoAll", "description": "Full mesh, 50 iterations", "test_filter": "NetIbMPITest.AlltoAllStress" }, + { "name": "NetIB_Stress_AlltoAll", "description": "Full mesh, 50 iterations", "test_filter": "NetIbMPITest.AllToAllStress" }, { "name": "NetIB_Stress_BidirectionalMultiRank","description": "All pairs bidirectional, 50 iterations", "test_filter": "NetIbMPITest.BidirectionalMultiRank" }, { "name": "NetIB_Stress_LongRunningMultiRank", "description": "Fan-in 3->1, 1000 iterations", "test_filter": "NetIbMPITest.LongRunningMultiRank" } ] @@ -3879,6 +3964,24 @@ "config": "net_transport_eth_multinode", "enabled": true }, + { + "name": "NET IB - Multithread Independent Connections (4)", + "description": "4 worker threads/rank, each on an independent NET/IB context and connection", + "config": "net_ib_multithread_independent", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (4)", + "description": "4 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread", + "enabled": true + }, + { + "name": "NET IB - Fault Injection Multithread (4)", + "description": "Per-communicator fault injection across concurrent workers, including cross-worker isolation", + "config": "fault_inject_multithread", + "enabled": true + }, { "name": "GDR Flush - Whitebox ncclIbIflush (Single-Node, 2-Rank)", "description": "Whitebox coverage of the NET/IB GDR flush fix: dma-buf and peermem RO=0 scratchpad reads, feature-disabled fallback, repeated-flush burst, and the fault-injection regression guard", 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 845885c1919b..82936859bde8 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 @@ -209,6 +209,91 @@ "NCCL_DMABUF_ENABLE": "1" } }, + "net_ib_multithread_independent": { + "extends": "net_ib_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "tests": [ + { + "name": "NetIB_MultiThread_IndependentSimpleSendRecv", + "description": "Concurrent workers each transfer on an independent NET/IB context and connection", + "test_filter": "NetIbMPITest.SimpleSendRecv" + }, + { + "name": "NetIB_MultiThread_IndependentSequentialTransfers", + "description": "Concurrent workers each run repeated transfers on an independent connection", + "test_filter": "NetIbMPITest.MultipleSequentialTransfers" + }, + { + "name": "NetIB_MultiThread_MultipleSizes", + "description": "Concurrent workers sweep the size ladder into the shared per-device MR cache", + "test_filter": "NetIbMPITest.SendRecvMultipleSizes" + }, + { + "name": "NetIB_MultiThread_MemoryRegistrationStorm", + "description": "Concurrent register/deregister storms race the per-device MR cache", + "test_filter": "NetIbMPITest.MemoryRegistrationStorm" + }, + { + "name": "NetIB_MultiThread_MrCacheRefCount", + "description": "All workers register the same address, contending on one cache entry's refcount", + "test_filter": "NetIbMPITest.MrCacheRefCount" + }, + { + "name": "NetIB_MultiThread_LongRunningEndurance", + "description": "Sustained concurrent traffic with the RDMA leak check around the whole run", + "test_filter": "NetIbMPITest.LongRunningEndurance" + } + ] + }, + "cast_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_MultiThread_StressMultiRound", + "description": "Concurrent IB-CAST connections, each auditing its own WRR token ledger", + "test_filter": "NetIbMPITest.CastStressMultiRoundTwoConns" + }, + { + "name": "Cast_MultiThread_SendRecvMultipleSizes", + "description": "Concurrent schedulers sweep both sides of the WRR/split boundary", + "test_filter": "NetIbMPITest.CastSendRecvMultipleSizes" + } + ] + }, + "fault_inject_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "NCCL_IB_QPS_PER_CONNECTION": "4", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "100", + "RCCL_IB_QP_SCHED_RESET_INTERVAL": "1000", + "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN": "2000", + "NCCL_IB_RETURN_ASYNC_EVENTS": "1", + "RCCL_IB_FAULT_INJECTION": "0" + }, + "tests": [ + { + "name": "FaultInj_MultiThread_IsolationAcrossWorkers", + "description": "One worker breaks its own connection while the others keep transferring cleanly", + "test_filter": "NetIbMPITest.FaultIsolationAcrossWorkers" + }, + { + "name": "FaultInj_MultiThread_QpErrorClearRecovers", + "description": "Each worker breaks one connection, clears the fault and must move data on a second one", + "test_filter": "NetIbMPITest.FaultInjCastQpErrorClearRecovers" + } + ] + }, "net_ib_gdr_flush": { "extends": "net_ib_base", "timeout": 120, @@ -1619,7 +1704,7 @@ { "name": "UBR_AlltoAll_OutOfPlace_MultiNode", "description": "AlltoAll out-of-place with UBR on multi-node", - "test_filter": "UBR_AlltoAll.OutOfPlace_MultiNode" + "test_filter": "UBR_AllToAll.OutOfPlace_MultiNode" }, { "name": "UBR_SendRecv_RingPattern_MultiNode", @@ -1629,19 +1714,19 @@ { "name": "UBR_AlltoAllStress_InterleavedMultiSize", "description": "Stress test with interleaved AlltoAll/AlltoAllv across small/medium/large sizes", - "test_filter": "UBR_AlltoAllStress.InterleavedMultiSize_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedMultiSize_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_SplitComm", "description": "Stress test with interleaved AlltoAll/AlltoAllv using split communicators", - "test_filter": "UBR_AlltoAllStress.InterleavedWithSplitComm_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedWithSplitComm_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_ConcurrentMultiComm", "description": "Concurrent AlltoAll on multiple communicators with separate streams", - "test_filter": "UBR_AlltoAllStress.ConcurrentMultiComm_MultiNode", + "test_filter": "UBR_AllToAllStress.ConcurrentMultiComm_MultiNode", "timeout": 300 }, { @@ -1652,12 +1737,12 @@ { "name": "UBR_ConcurrentRegHang_AlltoAll_Contiguous", "description": "ncclAlltoAll with contiguous registered buffers (should pass)", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_ConcurrentRegHang_AlltoAll_Stress", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -1702,12 +1787,12 @@ { "name": "UBR_AlltoAll_Contiguous_Success", "description": "ncclAlltoAll with contiguous registered buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_AlltoAll_Stress_Success", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2090,7 +2175,7 @@ { "name": "GraphCapture_AlltoAll_MultiNode", "description": "AlltoAll graph capture with IPC registration on multi-node", - "test_filter": "GraphCapture_AlltoAll.MultiNode" + "test_filter": "GraphCapture_AllToAll.MultiNode" }, { "name": "GraphCapture_AllReduce_MultiNode", @@ -2299,7 +2384,7 @@ { "name": "NetIB_Stress_AlltoAll", "description": "Full mesh, 50 iterations", - "test_filter": "NetIbMPITest.AlltoAllStress" + "test_filter": "NetIbMPITest.AllToAllStress" }, { "name": "NetIB_Stress_BidirectionalMultiRank", @@ -4249,6 +4334,24 @@ "config": "net_transport_eth_multinode", "enabled": true }, + { + "name": "NET IB - Multithread Independent Connections (4)", + "description": "4 worker threads/rank, each on an independent NET/IB context and connection", + "config": "net_ib_multithread_independent", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (4)", + "description": "4 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread", + "enabled": true + }, + { + "name": "NET IB - Fault Injection Multithread (4)", + "description": "Per-communicator fault injection across concurrent workers, including cross-worker isolation", + "config": "fault_inject_multithread", + "enabled": true + }, { "name": "GDR Flush - Whitebox ncclIbIflush (Single-Node, 2-Rank)", "description": "Whitebox coverage of the NET/IB GDR flush fix: dma-buf and peermem RO=0 scratchpad reads, feature-disabled fallback, repeated-flush burst, and the fault-injection regression guard", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce_func.json index 8ecea09693d3..64726306e887 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce_func.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce_func.json @@ -196,6 +196,91 @@ "NCCL_DMABUF_ENABLE": "1" } }, + "net_ib_multithread_independent": { + "extends": "net_ib_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "tests": [ + { + "name": "NetIB_MultiThread_IndependentSimpleSendRecv", + "description": "Concurrent workers each transfer on an independent NET/IB context and connection", + "test_filter": "NetIbMPITest.SimpleSendRecv" + }, + { + "name": "NetIB_MultiThread_IndependentSequentialTransfers", + "description": "Concurrent workers each run repeated transfers on an independent connection", + "test_filter": "NetIbMPITest.MultipleSequentialTransfers" + }, + { + "name": "NetIB_MultiThread_MultipleSizes", + "description": "Concurrent workers sweep the size ladder into the shared per-device MR cache", + "test_filter": "NetIbMPITest.SendRecvMultipleSizes" + }, + { + "name": "NetIB_MultiThread_MemoryRegistrationStorm", + "description": "Concurrent register/deregister storms race the per-device MR cache", + "test_filter": "NetIbMPITest.MemoryRegistrationStorm" + }, + { + "name": "NetIB_MultiThread_MrCacheRefCount", + "description": "All workers register the same address, contending on one cache entry's refcount", + "test_filter": "NetIbMPITest.MrCacheRefCount" + }, + { + "name": "NetIB_MultiThread_LongRunningEndurance", + "description": "Sustained concurrent traffic with the RDMA leak check around the whole run", + "test_filter": "NetIbMPITest.LongRunningEndurance" + } + ] + }, + "cast_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_MultiThread_StressMultiRound", + "description": "Concurrent IB-CAST connections, each auditing its own WRR token ledger", + "test_filter": "NetIbMPITest.CastStressMultiRoundTwoConns" + }, + { + "name": "Cast_MultiThread_SendRecvMultipleSizes", + "description": "Concurrent schedulers sweep both sides of the WRR/split boundary", + "test_filter": "NetIbMPITest.CastSendRecvMultipleSizes" + } + ] + }, + "fault_inject_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "NCCL_IB_QPS_PER_CONNECTION": "4", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "100", + "RCCL_IB_QP_SCHED_RESET_INTERVAL": "1000", + "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN": "2000", + "NCCL_IB_RETURN_ASYNC_EVENTS": "1", + "RCCL_IB_FAULT_INJECTION": "0" + }, + "tests": [ + { + "name": "FaultInj_MultiThread_IsolationAcrossWorkers", + "description": "One worker breaks its own connection while the others keep transferring cleanly", + "test_filter": "NetIbMPITest.FaultIsolationAcrossWorkers" + }, + { + "name": "FaultInj_MultiThread_QpErrorClearRecovers", + "description": "Each worker breaks one connection, clears the fault and must move data on a second one", + "test_filter": "NetIbMPITest.FaultInjCastQpErrorClearRecovers" + } + ] + }, "net_ib_gdr_flush": { "extends": "net_ib_base", "timeout": 120, @@ -908,7 +993,7 @@ { "name": "NetIB_Stress_AlltoAll", "description": "Full mesh, 50 iterations", - "test_filter": "NetIbMPITest.AlltoAllStress" + "test_filter": "NetIbMPITest.AllToAllStress" }, { "name": "NetIB_Stress_BidirectionalMultiRank", @@ -2017,7 +2102,7 @@ { "name": "UBR_AlltoAll_OutOfPlace_MultiNode", "description": "AlltoAll out-of-place with UBR on multi-node", - "test_filter": "UBR_AlltoAll.OutOfPlace_MultiNode" + "test_filter": "UBR_AllToAll.OutOfPlace_MultiNode" }, { "name": "UBR_SendRecv_RingPattern_MultiNode", @@ -2027,19 +2112,19 @@ { "name": "UBR_AlltoAllStress_InterleavedMultiSize", "description": "Stress test with interleaved AlltoAll/AlltoAllv across small/medium/large sizes", - "test_filter": "UBR_AlltoAllStress.InterleavedMultiSize_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedMultiSize_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_SplitComm", "description": "Stress test with interleaved AlltoAll/AlltoAllv using split communicators", - "test_filter": "UBR_AlltoAllStress.InterleavedWithSplitComm_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedWithSplitComm_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_ConcurrentMultiComm", "description": "Concurrent AlltoAll on multiple communicators with separate streams", - "test_filter": "UBR_AlltoAllStress.ConcurrentMultiComm_MultiNode", + "test_filter": "UBR_AllToAllStress.ConcurrentMultiComm_MultiNode", "timeout": 300 }, { @@ -2050,12 +2135,12 @@ { "name": "UBR_ConcurrentRegHang_AlltoAll_Contiguous", "description": "ncclAlltoAll with contiguous registered buffers (should pass)", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_ConcurrentRegHang_AlltoAll_Stress", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2101,12 +2186,12 @@ { "name": "UBR_AlltoAll_Contiguous_Success", "description": "ncclAlltoAll with contiguous registered buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_AlltoAll_Stress_Success", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2147,7 +2232,7 @@ { "name": "GraphCapture_AlltoAll_MultiNode", "description": "AlltoAll graph capture with IPC registration on multi-node", - "test_filter": "GraphCapture_AlltoAll.MultiNode" + "test_filter": "GraphCapture_AllToAll.MultiNode" }, { "name": "GraphCapture_AllReduce_MultiNode", @@ -2793,6 +2878,24 @@ "config": "net_transport_eth_multinode", "enabled": true }, + { + "name": "NET IB - Multithread Independent Connections (4)", + "description": "4 worker threads/rank, each on an independent NET/IB context and connection", + "config": "net_ib_multithread_independent", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (4)", + "description": "4 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread", + "enabled": true + }, + { + "name": "NET IB - Fault Injection Multithread (4)", + "description": "Per-communicator fault injection across concurrent workers, including cross-worker isolation", + "config": "fault_inject_multithread", + "enabled": true + }, { "name": "GDR Flush - Whitebox ncclIbIflush (Single-Node, 2-Rank)", "description": "Whitebox coverage of the NET/IB GDR flush fix: dma-buf and peermem RO=0 scratchpad reads, feature-disabled fallback, repeated-flush burst, and the fault-injection regression guard", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json index b71b13ff2374..24bc2e34c36c 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json @@ -205,6 +205,91 @@ "NCCL_DMABUF_ENABLE": "1" } }, + "net_ib_multithread_independent": { + "extends": "net_ib_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "tests": [ + { + "name": "NetIB_MultiThread_IndependentSimpleSendRecv", + "description": "Concurrent workers each transfer on an independent NET/IB context and connection", + "test_filter": "NetIbMPITest.SimpleSendRecv" + }, + { + "name": "NetIB_MultiThread_IndependentSequentialTransfers", + "description": "Concurrent workers each run repeated transfers on an independent connection", + "test_filter": "NetIbMPITest.MultipleSequentialTransfers" + }, + { + "name": "NetIB_MultiThread_MultipleSizes", + "description": "Concurrent workers sweep the size ladder into the shared per-device MR cache", + "test_filter": "NetIbMPITest.SendRecvMultipleSizes" + }, + { + "name": "NetIB_MultiThread_MemoryRegistrationStorm", + "description": "Concurrent register/deregister storms race the per-device MR cache", + "test_filter": "NetIbMPITest.MemoryRegistrationStorm" + }, + { + "name": "NetIB_MultiThread_MrCacheRefCount", + "description": "All workers register the same address, contending on one cache entry's refcount", + "test_filter": "NetIbMPITest.MrCacheRefCount" + }, + { + "name": "NetIB_MultiThread_LongRunningEndurance", + "description": "Sustained concurrent traffic with the RDMA leak check around the whole run", + "test_filter": "NetIbMPITest.LongRunningEndurance" + } + ] + }, + "cast_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_MultiThread_StressMultiRound", + "description": "Concurrent IB-CAST connections, each auditing its own WRR token ledger", + "test_filter": "NetIbMPITest.CastStressMultiRoundTwoConns" + }, + { + "name": "Cast_MultiThread_SendRecvMultipleSizes", + "description": "Concurrent schedulers sweep both sides of the WRR/split boundary", + "test_filter": "NetIbMPITest.CastSendRecvMultipleSizes" + } + ] + }, + "fault_inject_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "NCCL_IB_QPS_PER_CONNECTION": "4", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "100", + "RCCL_IB_QP_SCHED_RESET_INTERVAL": "1000", + "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN": "2000", + "NCCL_IB_RETURN_ASYNC_EVENTS": "1", + "RCCL_IB_FAULT_INJECTION": "0" + }, + "tests": [ + { + "name": "FaultInj_MultiThread_IsolationAcrossWorkers", + "description": "One worker breaks its own connection while the others keep transferring cleanly", + "test_filter": "NetIbMPITest.FaultIsolationAcrossWorkers" + }, + { + "name": "FaultInj_MultiThread_QpErrorClearRecovers", + "description": "Each worker breaks one connection, clears the fault and must move data on a second one", + "test_filter": "NetIbMPITest.FaultInjCastQpErrorClearRecovers" + } + ] + }, "net_ib_gdr_flush": { "extends": "net_ib_base", "timeout": 120, @@ -917,7 +1002,7 @@ { "name": "NetIB_Stress_AlltoAll", "description": "Full mesh, 50 iterations", - "test_filter": "NetIbMPITest.AlltoAllStress" + "test_filter": "NetIbMPITest.AllToAllStress" }, { "name": "NetIB_Stress_BidirectionalMultiRank", @@ -1985,7 +2070,7 @@ { "name": "UBR_AlltoAll_OutOfPlace_MultiNode", "description": "AlltoAll out-of-place with UBR on multi-node", - "test_filter": "UBR_AlltoAll.OutOfPlace_MultiNode" + "test_filter": "UBR_AllToAll.OutOfPlace_MultiNode" }, { "name": "UBR_SendRecv_RingPattern_MultiNode", @@ -1995,19 +2080,19 @@ { "name": "UBR_AlltoAllStress_InterleavedMultiSize", "description": "Stress test with interleaved AlltoAll/AlltoAllv across small/medium/large sizes", - "test_filter": "UBR_AlltoAllStress.InterleavedMultiSize_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedMultiSize_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_SplitComm", "description": "Stress test with interleaved AlltoAll/AlltoAllv using split communicators", - "test_filter": "UBR_AlltoAllStress.InterleavedWithSplitComm_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedWithSplitComm_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_ConcurrentMultiComm", "description": "Concurrent AlltoAll on multiple communicators with separate streams", - "test_filter": "UBR_AlltoAllStress.ConcurrentMultiComm_MultiNode", + "test_filter": "UBR_AllToAllStress.ConcurrentMultiComm_MultiNode", "timeout": 300 }, { @@ -2018,12 +2103,12 @@ { "name": "UBR_ConcurrentRegHang_AlltoAll_Contiguous", "description": "ncclAlltoAll with contiguous registered buffers (should pass)", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_ConcurrentRegHang_AlltoAll_Stress", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2069,12 +2154,12 @@ { "name": "UBR_AlltoAll_Contiguous_Success", "description": "ncclAlltoAll with contiguous registered buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_AlltoAll_Stress_Success", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2115,7 +2200,7 @@ { "name": "GraphCapture_AlltoAll_MultiNode", "description": "AlltoAll graph capture with IPC registration on multi-node", - "test_filter": "GraphCapture_AlltoAll.MultiNode" + "test_filter": "GraphCapture_AllToAll.MultiNode" }, { "name": "GraphCapture_AllReduce_MultiNode", @@ -3581,6 +3666,24 @@ "config": "net_transport_eth_multinode", "enabled": true }, + { + "name": "NET IB - Multithread Independent Connections (4)", + "description": "4 worker threads/rank, each on an independent NET/IB context and connection", + "config": "net_ib_multithread_independent", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (4)", + "description": "4 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread", + "enabled": true + }, + { + "name": "NET IB - Fault Injection Multithread (4)", + "description": "Per-communicator fault injection across concurrent workers, including cross-worker isolation", + "config": "fault_inject_multithread", + "enabled": true + }, { "name": "GDR Flush - Whitebox ncclIbIflush (Single-Node, 2-Rank)", "description": "Whitebox coverage of the NET/IB GDR flush fix: dma-buf and peermem RO=0 scratchpad reads, feature-disabled fallback, repeated-flush burst, and the fault-injection regression guard", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json index b49d9e81c74d..fed3ed264951 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json @@ -190,6 +190,91 @@ "NCCL_DMABUF_ENABLE": "1" } }, + "net_ib_multithread_independent": { + "extends": "net_ib_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "tests": [ + { + "name": "NetIB_MultiThread_IndependentSimpleSendRecv", + "description": "Concurrent workers each transfer on an independent NET/IB context and connection", + "test_filter": "NetIbMPITest.SimpleSendRecv" + }, + { + "name": "NetIB_MultiThread_IndependentSequentialTransfers", + "description": "Concurrent workers each run repeated transfers on an independent connection", + "test_filter": "NetIbMPITest.MultipleSequentialTransfers" + }, + { + "name": "NetIB_MultiThread_MultipleSizes", + "description": "Concurrent workers sweep the size ladder into the shared per-device MR cache", + "test_filter": "NetIbMPITest.SendRecvMultipleSizes" + }, + { + "name": "NetIB_MultiThread_MemoryRegistrationStorm", + "description": "Concurrent register/deregister storms race the per-device MR cache", + "test_filter": "NetIbMPITest.MemoryRegistrationStorm" + }, + { + "name": "NetIB_MultiThread_MrCacheRefCount", + "description": "All workers register the same address, contending on one cache entry's refcount", + "test_filter": "NetIbMPITest.MrCacheRefCount" + }, + { + "name": "NetIB_MultiThread_LongRunningEndurance", + "description": "Sustained concurrent traffic with the RDMA leak check around the whole run", + "test_filter": "NetIbMPITest.LongRunningEndurance" + } + ] + }, + "cast_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_MultiThread_StressMultiRound", + "description": "Concurrent IB-CAST connections, each auditing its own WRR token ledger", + "test_filter": "NetIbMPITest.CastStressMultiRoundTwoConns" + }, + { + "name": "Cast_MultiThread_SendRecvMultipleSizes", + "description": "Concurrent schedulers sweep both sides of the WRR/split boundary", + "test_filter": "NetIbMPITest.CastSendRecvMultipleSizes" + } + ] + }, + "fault_inject_multithread": { + "extends": "cast_base", + "timeout": 600, + "command_args": "--net_ib_nthreads=4", + "env_variables": { + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "NCCL_IB_QPS_PER_CONNECTION": "4", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "100", + "RCCL_IB_QP_SCHED_RESET_INTERVAL": "1000", + "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN": "2000", + "NCCL_IB_RETURN_ASYNC_EVENTS": "1", + "RCCL_IB_FAULT_INJECTION": "0" + }, + "tests": [ + { + "name": "FaultInj_MultiThread_IsolationAcrossWorkers", + "description": "One worker breaks its own connection while the others keep transferring cleanly", + "test_filter": "NetIbMPITest.FaultIsolationAcrossWorkers" + }, + { + "name": "FaultInj_MultiThread_QpErrorClearRecovers", + "description": "Each worker breaks one connection, clears the fault and must move data on a second one", + "test_filter": "NetIbMPITest.FaultInjCastQpErrorClearRecovers" + } + ] + }, "net_ib_gdr_flush": { "extends": "net_ib_base", "timeout": 120, @@ -902,7 +987,7 @@ { "name": "NetIB_Stress_AlltoAll", "description": "Full mesh, 50 iterations", - "test_filter": "NetIbMPITest.AlltoAllStress" + "test_filter": "NetIbMPITest.AllToAllStress" }, { "name": "NetIB_Stress_BidirectionalMultiRank", @@ -1970,7 +2055,7 @@ { "name": "UBR_AlltoAll_OutOfPlace_MultiNode", "description": "AlltoAll out-of-place with UBR on multi-node", - "test_filter": "UBR_AlltoAll.OutOfPlace_MultiNode" + "test_filter": "UBR_AllToAll.OutOfPlace_MultiNode" }, { "name": "UBR_SendRecv_RingPattern_MultiNode", @@ -1980,19 +2065,19 @@ { "name": "UBR_AlltoAllStress_InterleavedMultiSize", "description": "Stress test with interleaved AlltoAll/AlltoAllv across small/medium/large sizes", - "test_filter": "UBR_AlltoAllStress.InterleavedMultiSize_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedMultiSize_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_SplitComm", "description": "Stress test with interleaved AlltoAll/AlltoAllv using split communicators", - "test_filter": "UBR_AlltoAllStress.InterleavedWithSplitComm_MultiNode", + "test_filter": "UBR_AllToAllStress.InterleavedWithSplitComm_MultiNode", "timeout": 300 }, { "name": "UBR_AlltoAllStress_ConcurrentMultiComm", "description": "Concurrent AlltoAll on multiple communicators with separate streams", - "test_filter": "UBR_AlltoAllStress.ConcurrentMultiComm_MultiNode", + "test_filter": "UBR_AllToAllStress.ConcurrentMultiComm_MultiNode", "timeout": 300 }, { @@ -2003,12 +2088,12 @@ { "name": "UBR_ConcurrentRegHang_AlltoAll_Contiguous", "description": "ncclAlltoAll with contiguous registered buffers (should pass)", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_ConcurrentRegHang_AlltoAll_Stress", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2054,12 +2139,12 @@ { "name": "UBR_AlltoAll_Contiguous_Success", "description": "ncclAlltoAll with contiguous registered buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_SHOULD_WORK" + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_SHOULD_WORK" }, { "name": "UBR_AlltoAll_Stress_Success", "description": "Stress test: 50 iterations of ncclAlltoAll with contiguous buffers", - "test_filter": "UBR_ConcurrentRegHang.AlltoAll_ContiguousBuffer_Stress_SHOULD_WORK", + "test_filter": "UBR_ConcurrentRegHang.AllToAll_ContiguousBuffer_Stress_SHOULD_WORK", "timeout": 300 }, { @@ -2100,7 +2185,7 @@ { "name": "GraphCapture_AlltoAll_MultiNode", "description": "AlltoAll graph capture with IPC registration on multi-node", - "test_filter": "GraphCapture_AlltoAll.MultiNode" + "test_filter": "GraphCapture_AllToAll.MultiNode" }, { "name": "GraphCapture_AllReduce_MultiNode", @@ -2746,6 +2831,24 @@ "config": "net_transport_eth_multinode", "enabled": true }, + { + "name": "NET IB - Multithread Independent Connections (4)", + "description": "4 worker threads/rank, each on an independent NET/IB context and connection", + "config": "net_ib_multithread_independent", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (4)", + "description": "4 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread", + "enabled": true + }, + { + "name": "NET IB - Fault Injection Multithread (4)", + "description": "Per-communicator fault injection across concurrent workers, including cross-worker isolation", + "config": "fault_inject_multithread", + "enabled": true + }, { "name": "GDR Flush - Whitebox ncclIbIflush (Single-Node, 2-Rank)", "description": "Whitebox coverage of the NET/IB GDR flush fix: dma-buf and peermem RO=0 scratchpad reads, feature-disabled fallback, repeated-flush burst, and the fault-injection regression guard", diff --git a/projects/rccl/tools/scripts/test_runner/configs/net_ib_transport.json b/projects/rccl/tools/scripts/test_runner/configs/net_ib_transport.json index 8e8abcab214e..40984e4987a5 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/net_ib_transport.json +++ b/projects/rccl/tools/scripts/test_runner/configs/net_ib_transport.json @@ -595,6 +595,22 @@ ] }, + "cast_shim_leak_guard": { + "extends": "cast_base", + "timeout": 120, + "description": "Runs the shim-leak probe in a plain CAST configuration, where RCCL_IB_FAULT_INJECTION is simply absent. The damage a leak does is process-wide and has nothing to do with threading, so the guard belongs in an ordinary run and not only in the threaded fault suites.", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1" + }, + "tests": [ + { + "name": "Cast_FaultShimsAbsentUnlessRequested", + "description": "Arming an ops-level fault must be refused while RCCL_IB_FAULT_INJECTION is off, which is only true while the libibverbs shims are absent. If they are present, every poll_cq in the process pays for a global mutex twice", + "test_filter": "NetIbMPITest.FaultInjectionShimsAbsentUnlessRequested" + } + ] + }, "cast_stress": { "extends": "cast_base", "timeout": 600, @@ -976,13 +992,17 @@ "NCCL_NET": "ib-cast" } }, - "net_ib_multithread_independent": { + "net_ib_multithread_base": { "extends": "net_ib_base", "timeout": 300, + "description": "Shared environment for the threaded suites. Deliberately carries no tests: extends merges test lists rather than replacing them, so a threaded suite that extended a suite with tests would also run those tests, at whatever worker count it sets.", "env_variables": { "NCCL_DEBUG": "INFO", "NCCL_DEBUG_SUBSYS": "NET,INIT" - }, + } + }, + "net_ib_multithread_independent": { + "extends": "net_ib_multithread_base", "tests": [ { "name": "NetIB_MultiThread_IndependentSimpleSendRecv", @@ -993,6 +1013,124 @@ "name": "NetIB_MultiThread_IndependentSequentialTransfers", "description": "Independent worker connections concurrently execute sequential host-memory transfers", "test_filter": "NetIbMPITest.MultipleSequentialTransfers" + }, + { + "name": "NetIB_MultiThread_MultipleSizes", + "description": "Concurrent workers sweep the size ladder, filling the shared per-device MR cache with registrations of many sizes", + "test_filter": "NetIbMPITest.SendRecvMultipleSizes" + }, + { + "name": "NetIB_MultiThread_ProgressDoesNotSerialize", + "description": "Gate: the same transfer loop at one worker and at all of them, so a lock held across a progress call shows up as time that scales with the worker count. Skips below four workers", + "test_filter": "NetIbMPITest.ThreadedProgressDoesNotSerialize" + }, + { + "name": "NetIB_MultiThread_LargeTransfer", + "description": "Concurrent 16 MB transfers keep several large registrations live on one device", + "test_filter": "NetIbMPITest.LargeTransfer" + }, + { + "name": "NetIB_MultiThread_MemoryRegistrationStorm", + "description": "Concurrent register/deregister storms race the insert, lookup and eviction paths of the per-device MR cache", + "test_filter": "NetIbMPITest.MemoryRegistrationStorm" + }, + { + "name": "NetIB_MultiThread_MrCacheRefCount", + "description": "All workers register the same address, contending on one cache entry's refcount", + "test_filter": "NetIbMPITest.MrCacheRefCount" + }, + { + "name": "NetIB_MultiThread_MixedSizeBarrage", + "description": "Concurrent workers walk 1 B to 64 MB, exercising alignment and chunking under bandwidth pressure", + "test_filter": "NetIbMPITest.MixedSizeBarrage" + }, + { + "name": "NetIB_MultiThread_TagZeroReuse", + "description": "Every worker reuses tag 0 on its own connection, so a misrouted completion shows up as a data mismatch", + "test_filter": "NetIbMPITest.TagZeroReuse" + }, + { + "name": "NetIB_MultiThread_RequestSlotExhaustion", + "description": "Concurrent workers hold 32 requests each in flight on one NIC and verify slot recycling", + "test_filter": "NetIbMPITest.RequestSlotExhaustion" + }, + { + "name": "NetIB_MultiThread_RapidRecvPostDrain", + "description": "Concurrent post-and-drain batches, so recycling holds up while other communicators drain", + "test_filter": "NetIbMPITest.RapidRecvPostDrain" + }, + { + "name": "NetIB_MultiThread_FifoPressure", + "description": "N slow receivers against N tight senders, with the backpressure assertion kept per connection", + "test_filter": "NetIbMPITest.FifoPressureSenderFast" + }, + { + "name": "NetIB_MultiThread_LongRunningEndurance", + "description": "Sustained concurrent traffic with the RDMA leak check taken around the whole run", + "test_filter": "NetIbMPITest.LongRunningEndurance" + }, + { + "name": "NetIB_MultiThread_GpuMemoryTransfer", + "description": "Concurrent GPU-buffer registrations drive the peer-memory path into the shared MR cache", + "test_filter": "NetIbMPITest.GpuMemoryTransferStress" + }, + { + "name": "NetIB_MultiThread_VNicRegDeregCycling", + "description": "Concurrent register/deregister cycling on a fused device, fanning out across both members' protection domains", + "test_filter": "NetIbMPITest.RegDeregCycling_VNic" + }, + { + "name": "NetIB_MultiThread_VNicSequentialTransfers", + "description": "Concurrent long-lived registrations on one fused device", + "test_filter": "NetIbMPITest.SequentialTransfers_VNic" + }, + { + "name": "NetIB_MultiThread_VNicMixedSizes", + "description": "Concurrent uneven QP splits across a fused device", + "test_filter": "NetIbMPITest.MixedSizes_VNic" + }, + { + "name": "NetIB_MultiThread_VNicLargeTransfer", + "description": "Concurrent 64 MB striped transfers over a fused device", + "test_filter": "NetIbMPITest.LargeTransfer_VNic" + }, + { + "name": "NetIB_MultiThread_VNicUnalignedSizes", + "description": "Concurrent transfers at the 128-byte striping boundary of a fused device", + "test_filter": "NetIbMPITest.UnalignedSizeTransfer_VNic" + }, + { + "name": "NetIB_MultiThread_VNicFlushRepeated", + "description": "Concurrent GPU receives and flushes on one fused device", + "test_filter": "NetIbMPITest.FlushRepeated_VNic" + } + ] + }, + "net_ib_multithread_multiqp": { + "extends": "net_ib_multithread_base", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "4", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1" + }, + "tests": [ + { + "name": "NetIB_MultiThread_MultiQpSplitData", + "description": "Concurrent workers multiply the QPs in flight per device while each checks its own split payloads", + "test_filter": "NetIbMPITest.MultiQpSplitDataStress" + } + ] + }, + "net_ib_multithread_multiqp_nosplit": { + "extends": "net_ib_multithread_base", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "4", + "NCCL_IB_SPLIT_DATA_ON_QPS": "0" + }, + "tests": [ + { + "name": "NetIB_MultiThread_MultiQpNoSplit", + "description": "Same concurrent sweep on the nDataQps path", + "test_filter": "NetIbMPITest.MultiQpNoSplitStress" } ] }, @@ -1007,6 +1145,189 @@ "net_ib_multithread_independent_16": { "extends": "net_ib_multithread_independent", "command_args": "--net_ib_nthreads=16" + }, + "net_ib_multithread_multiqp_4": { + "extends": "net_ib_multithread_multiqp", + "command_args": "--net_ib_nthreads=4" + }, + "net_ib_multithread_multiqp_nosplit_4": { + "extends": "net_ib_multithread_multiqp_nosplit", + "command_args": "--net_ib_nthreads=4" + }, + "cast_multithread": { + "extends": "cast_base", + "timeout": 300, + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_MultiThread_StressMultiRound", + "description": "Every worker drives its own CAST connection, so the WRR schedulers advance simultaneously instead of one after another", + "test_filter": "NetIbMPITest.CastStressMultiRoundTwoConns" + }, + { + "name": "Cast_MultiThread_SendRecvMultipleSizes", + "description": "Concurrent schedulers sweep both sides of the WRR/split boundary, each auditing its own token ledger", + "test_filter": "NetIbMPITest.CastSendRecvMultipleSizes" + }, + { + "name": "Cast_MultiThread_LargeTransfer", + "description": "Concurrent 16 MB split-path transfers, each confirming its scheduler consumed no tokens", + "test_filter": "NetIbMPITest.CastLargeTransfer" + }, + { + "name": "Cast_MultiThread_SendRecvZeroSize", + "description": "Zero-byte sends still take the WRR path on every worker's scheduler", + "test_filter": "NetIbMPITest.CastSendRecvZeroSize" + }, + { + "name": "Cast_MultiThread_SplitDataThresholdBoundary", + "description": "Concurrent schedulers decide the split boundary at the same time, each auditing its own tokens", + "test_filter": "NetIbMPITest.CastSplitDataThresholdBoundary" + } + ] + }, + "cast_multithread_2": { + "extends": "cast_multithread", + "command_args": "--net_ib_nthreads=2" + }, + "cast_multithread_4": { + "extends": "cast_multithread", + "command_args": "--net_ib_nthreads=4" + }, + "cast_multithread_16": { + "extends": "cast_multithread", + "command_args": "--net_ib_nthreads=16" + }, + "cast_wrr_split_long_interval": { + "extends": "cast_base", + "timeout": 120, + "description": "Extends cast_base rather than cast_wrr_split so it does not also inherit that suite's tests, several of which need the RTT timer to fire and would fail at this interval.", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_EnableDisableSched_LongInterval", + "description": "Toggle schedEnable mid-connection. Needs an update interval above 10 s so the RTT timer cannot rewrite tokens mid-assertion; the default 50 us makes this test skip", + "test_filter": "NetIbMPITest.CastEnableDisableSched" + } + ] + }, + "cast_wrr_no_split_long_interval": { + "extends": "cast_base", + "timeout": 120, + "description": "Extends cast_base rather than cast_wrr_no_split for the same reason as its split counterpart.", + "env_variables": { + "NCCL_IB_QPS_PER_CONNECTION": "2", + "NCCL_IB_SPLIT_DATA_ON_QPS": "0", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "10000000" + }, + "tests": [ + { + "name": "Cast_WeightsDistributionOneRound_LongInterval", + "description": "Two-phase WRR distribution with the RTT timer suspended", + "test_filter": "NetIbMPITest.CastWeightsDistributionOneRound" + }, + { + "name": "Cast_AlternatingWrrNonWrr_LongInterval", + "description": "Toggle doWrr mid-connection with the RTT timer suspended", + "test_filter": "NetIbMPITest.CastAlternatingWrrNonWrr" + }, + { + "name": "Cast_EnableDisableSplitData_LongInterval", + "description": "Toggle splitData at the threshold with the RTT timer suspended", + "test_filter": "NetIbMPITest.CastEnableDisableSplitData" + } + ] + }, + "fault_multithread_base": { + "extends": "cast_base", + "timeout": 600, + "description": "IB-CAST fault environment for the threaded suites, carrying no tests of its own for the reason given in net_ib_multithread_base. RCCL_IB_FAULT_INJECTION stays off: these tests use the per-communicator pre-call hooks and have no use for the libibverbs ops shims, which add two acquisitions of one process-global mutex to every poll_cq and whose registry is per device, so one worker's clear disarms every other worker's faults.", + "env_variables": { + "NCCL_IB_SPLIT_DATA_ON_QPS": "1", + "NCCL_IB_QPS_PER_CONNECTION": "4", + "RCCL_IB_QP_SCHED_UPDATE_INTERVAL": "100", + "RCCL_IB_QP_SCHED_RESET_INTERVAL": "1000", + "RCCL_IB_QP_SCHED_SPLIT_DATA_MIN": "2000", + "NCCL_IB_RETURN_ASYNC_EVENTS": "1", + "NCCL_IB_RESILIENCY_PORT_FAILOVER": "1", + "RCCL_IB_FAULT_INJECTION": "0" + } + }, + "fault_inject_multithread": { + "extends": "fault_multithread_base", + "tests": [ + { + "name": "FaultInj_ShimsAbsentUnlessRequested", + "description": "With RCCL_IB_FAULT_INJECTION off, arming an ops-level fault must be refused: if it succeeds the libibverbs shims are installed, and their poll_cq takes one process-global mutex per call, serializing progress for the whole process", + "test_filter": "NetIbMPITest.FaultInjectionShimsAbsentUnlessRequested" + }, + { + "name": "FaultInj_MultiThread_IsolationAcrossWorkers", + "description": "One worker breaks its own connection while the others keep transferring: proves injected fault state stays inside its communicator", + "test_filter": "NetIbMPITest.FaultIsolationAcrossWorkers" + }, + { + "name": "FaultInj_MultiThread_QpErrorIsFatal", + "description": "Every worker arms error injection on its own QPs and must observe the failure itself", + "test_filter": "NetIbMPITest.FaultInjCastQpErrorIsFatal" + }, + { + "name": "FaultInj_MultiThread_DelayDataIntegrity", + "description": "Every worker slows its own QP 0 and still delivers intact data while the others load the device", + "test_filter": "NetIbMPITest.FaultInjCastDelayDataIntegrity" + }, + { + "name": "FaultInj_MultiThread_QpErrorClearRecovers", + "description": "Each worker breaks one connection, clears the fault and must move data cleanly on a second connection", + "test_filter": "NetIbMPITest.FaultInjCastQpErrorClearRecovers" + }, + { + "name": "FaultInj_MultiThread_FailoverCqeErrorRecovered", + "description": "N links fail at once, so the process-global recovery thread serves several communicators", + "test_filter": "NetIbMPITest.FailoverCqeErrorRecovered" + }, + { + "name": "FaultInj_MultiThread_FailoverLargeMessage", + "description": "Concurrent 64 KB messages survive their own link failure byte for byte", + "test_filter": "NetIbMPITest.FailoverLargeMessageDataIntegrity" + }, + { + "name": "FaultInj_MultiThread_FailoverMultiRequest", + "description": "Several messages per worker cross a link failure, exercising concurrent failed requests", + "test_filter": "NetIbMPITest.FailoverMultiRequestInFlight" + } + ] + }, + "fault_inject_multithread_2": { + "extends": "fault_inject_multithread", + "command_args": "--net_ib_nthreads=2" + }, + "fault_inject_multithread_4": { + "extends": "fault_inject_multithread", + "command_args": "--net_ib_nthreads=4" + }, + "fault_recovery_multithread_2": { + "extends": "fault_multithread_base", + "timeout": 900, + "command_args": "--net_ib_nthreads=2", + "env_variables": { + "NCCL_IB_RESILIENCY_PORT_RECOVERY": "1" + }, + "tests": [ + { + "name": "FaultInj_MultiThread_RecoverySuccessRestoresTraffic", + "description": "Concurrent recoveries share one global recovery thread, then every worker's traffic must resume", + "test_filter": "NetIbMPITest.RecoverySuccessRestoresTraffic" + } + ] } }, "test_suites": [ @@ -1124,6 +1445,12 @@ "config": "cast_max_qp", "enabled": true }, + { + "name": "NET IB - CAST shim leak guard", + "description": "The libibverbs ops shims must not be installed in a plain run", + "config": "cast_shim_leak_guard", + "enabled": true + }, { "name": "NET IB - CAST Stress (QPS=2)", "description": "5000 sends on 2 concurrent connections + large-size split-path stress (1MB/4MB/16MB × 20): RTT timer fires mid-run, consistency invariants verified", @@ -1285,6 +1612,66 @@ "description": "16 worker threads/rank, each using an independent NET/IB context and connection", "config": "net_ib_multithread_independent_16", "enabled": true + }, + { + "name": "NET IB - Multithread Multi-QP Split (4)", + "description": "4 worker threads/rank with 4 QPs per connection and data split across QPs", + "config": "net_ib_multithread_multiqp_4", + "enabled": true + }, + { + "name": "NET IB - Multithread Multi-QP NoSplit (4)", + "description": "4 worker threads/rank with 4 QPs per connection on the nDataQps path", + "config": "net_ib_multithread_multiqp_nosplit_4", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (2)", + "description": "2 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread_2", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (4)", + "description": "4 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread_4", + "enabled": true + }, + { + "name": "NET IB - CAST Multithread (16)", + "description": "16 worker threads/rank driving independent IB-CAST connections and WRR schedulers", + "config": "cast_multithread_16", + "enabled": true + }, + { + "name": "NET IB - CAST WRR+Split, long update interval", + "description": "Scheduler toggle tests that need RCCL_IB_QP_SCHED_UPDATE_INTERVAL >= 10 s; they skip in every other config", + "config": "cast_wrr_split_long_interval", + "enabled": true + }, + { + "name": "NET IB - CAST WRR only, long update interval", + "description": "Scheduler toggle tests that need RCCL_IB_QP_SCHED_UPDATE_INTERVAL >= 10 s; they skip in every other config", + "config": "cast_wrr_no_split_long_interval", + "enabled": true + }, + { + "name": "NET IB - Fault Injection Multithread (2)", + "description": "2 worker threads/rank injecting per-communicator faults, including cross-worker isolation", + "config": "fault_inject_multithread_2", + "enabled": true + }, + { + "name": "NET IB - Fault Injection Multithread (4)", + "description": "4 worker threads/rank injecting per-communicator faults, including cross-worker isolation", + "config": "fault_inject_multithread_4", + "enabled": true + }, + { + "name": "NET IB - Fault Recovery Multithread (2)", + "description": "Concurrent link recoveries through the single process-global recovery thread", + "config": "fault_recovery_multithread_2", + "enabled": true } ] } diff --git a/projects/rccl/tools/scripts/test_runner/lib/schema.json b/projects/rccl/tools/scripts/test_runner/lib/schema.json index f2642de10d2d..4450eb5def0e 100644 --- a/projects/rccl/tools/scripts/test_runner/lib/schema.json +++ b/projects/rccl/tools/scripts/test_runner/lib/schema.json @@ -194,9 +194,9 @@ "description": "Number of nodes across which ranks are distributed" }, "num_gpus": { - "description": "GPUs per node; controls --map-by ppr:{num_gpus}:node for multi-node runs. The string \"auto\" (also the default when omitted) resolves to the detected GPUs per node.", + "description": "GPUs per node; controls --map-by ppr:{num_gpus}:node for multi-node runs. The string \"auto\" (also the default when omitted) resolves to the detected GPUs per node. Zero declares a test that needs no GPU, such as a CPU-only build smoke; the executor only uses this value for multi-node placement, so zero is inert for single-node runs.", "oneOf": [ - { "type": "integer", "minimum": 1 }, + { "type": "integer", "minimum": 0 }, { "type": "string", "enum": ["auto"] } ] }, diff --git a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py index f7660fe6bd70..051e51ba5b4a 100644 --- a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py +++ b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py @@ -1041,6 +1041,21 @@ def run_test(self, test_config, suite_config): "duration": 0, "error": msg, } + + # num_gpus of 0 declares a test that needs no GPU, which only makes sense + # on one node: multi-node placement builds "host:{num_gpus}" and + # "--map-by ppr:{num_gpus}:node" from it, and zero there would launch + # nothing. Fail this test rather than emit that. + if num_nodes > 1 and num_gpus < 1: + msg = (f"Invalid num_gpus for test '{test_name}': {num_gpus} GPUs/node cannot place " + f"ranks across {num_nodes} nodes") + print(f"ERROR: {msg}") + return { + "name": test_name, + "result": TestResult.RESULT_FAILED.value, + "duration": 0, + "error": msg, + } timeout = test_config.get("timeout", 0) env_vars = test_config.get("env_variables", {}) From f2cd5174b43d077a48fe6c3d103d0f4e7252a1e3 Mon Sep 17 00:00:00 2001 From: Ilya Kosarev Date: Thu, 20 Aug 2026 17:49:33 +0000 Subject: [PATCH 4/7] net-ib: document how the threaded tests were chosen Records the production threading model the option mirrors, why concurrent use of one communicator and concurrent plugin init stay out of scope, and the two-stage gate pattern the worker bodies use. Co-authored-by: Cursor --- .../multithread-validation.md | 35 +++++++++++++++++++ 1 file changed, 35 insertions(+) diff --git a/projects/rccl/.claude/skills/feature-unit-testing/multithread-validation.md b/projects/rccl/.claude/skills/feature-unit-testing/multithread-validation.md index a7f1d79b1998..acf6ac5df8b1 100644 --- a/projects/rccl/.claude/skills/feature-unit-testing/multithread-validation.md +++ b/projects/rccl/.claude/skills/feature-unit-testing/multithread-validation.md @@ -51,6 +51,41 @@ turns a multithread configuration into a single-threaded run: - reduce the value across all ranks and fail before the suite starts unless every rank resolved the same one. +## Choose which tests get a threaded variant + +Pick by the state concurrency would actually stress, not by how easy the body is +to convert. In descending order of value: + +- **Registration paths.** The MR cache and protection domain are per physical + device behind one mutex, so registration storms, refcount churn on a single + shared address, and fused-device cycling (where one registration fans out + across both members) are the only way to race insert, lookup and eviction. +- **Fault injection whose state is per communicator.** One worker breaking its + own connection while the others keep transferring is an assertion no serial + test can make. Concurrent link failures are also the only way to load a + process-global recovery thread. +- **Per-communicator scheduler state**, such as WRR tokens and cursor: every + worker arms and audits its own ledger, so N schedulers advance at once while + sharing nothing. +- **Data paths under pressure** — size ladders, multi-QP splits, request-pool + churn, GPU buffers. There is no software race to find in per-comm state; the + value is real queue contention plus proof that completion routing never + crosses communicators. + +Give every worker a distinct payload seed. A transfer delivered on the wrong +connection then fails verification instead of passing silently, which is the one +cheap defence against the whole model being wrong. + +Leave alone anything that asserts on process-global state: device and merged-NIC +tables, counters snapshotted process-wide, and timing claims about a shared +background thread. Resource-leak checks still work, but only on the main thread +around the entire run — an in-loop snapshot sees other workers' live resources +and reports a leak that is not there. + +Spread workers across devices when the suite has several NICs. Pinning every +worker to device 0 leaves the other devices' protection domains and caches +untouched. + ## Harness structure - The main thread creates every context and connection, exchanges handles over From 0bb81025facf5b77f56db914ffa0b1631051d37b Mon Sep 17 00:00:00 2001 From: Ilya Kosarev Date: Thu, 20 Aug 2026 22:58:45 +0000 Subject: [PATCH 5/7] net-ib: keep the CAST helpers buildable without fault injection, and reject empty placements WorkerCastPrepareTokens is an ordinary IB-CAST helper, but it was calling WorkerCastLiveNqps, which sat inside the ENABLE_FAULT_INJECTION block. Asking the scheduler how many QPs a connection uses has nothing to do with fault injection, and -DFAULT_INJECTION=OFF is a documented option, so that build could not compile the CAST tests at all. The helper moves out; only the fault APIs stay guarded. The placement guard also missed a case: num_ranks of "auto" derives the rank count from num_gpus, so zero GPUs resolved to zero ranks, which the single-node path accepted and mpirun would have been asked to run with -np 0. Any non-positive resolved rank or node count is now rejected along with a negative GPU count. Co-authored-by: Cursor --- .../transport/NetIbMPI/NetIbMPITestBase.hpp | 38 +++++++++---------- .../scripts/test_runner/lib/test_executor.py | 10 +++-- 2 files changed, 25 insertions(+), 23 deletions(-) diff --git a/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp b/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp index f76652e48d0f..6865887ceaff 100644 --- a/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp +++ b/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp @@ -1175,6 +1175,25 @@ class NetIbMPITest : public MPITestBase { return result; } + // QPs the connection actually uses. NCCL_IB_QPS_PER_CONNECTION states only a + // request: on a merged device the plugin creates that many per member, so + // arming faults from the environment would leave the remaining QPs healthy + // and "the send must fail" would depend on which QP the scheduler picked. + // Valid once the scheduler is warm, which the first successful send does. + ThreadResult WorkerCastLiveNqps(void* sendComm, int* nqps) { + struct ncclIbCastSchedState state; + ThreadResult result = WorkerCastGetSchedState(sendComm, &state); + if (!result.ok) return result; + if (state.nqps <= 0) { + result.ok = false; + result.msg = "scheduler reports nqps=" + std::to_string(state.nqps); + return result; + } + *nqps = state.nqps; + return result; + } +#endif /* ENABLE_FAULT_INJECTION */ + // Warm the scheduler up (it initializes on the first send) and arm equal // weights. nqps comes from the environment so both ranks agree without a // collective; the sender confirms the connection really uses that many QPs. @@ -1533,25 +1552,6 @@ class NetIbMPITest : public MPITestBase { return WorkerSendRecvPattern(rank, pair, buffer, size, tag, mhandle, seed, timeoutMs); } - // QPs the connection actually uses. NCCL_IB_QPS_PER_CONNECTION states only a - // request: on a merged device the plugin creates that many per member, so - // arming faults from the environment would leave the remaining QPs healthy - // and "the send must fail" would depend on which QP the scheduler picked. - // Valid once the scheduler is warm, which the first successful send does. - ThreadResult WorkerCastLiveNqps(void* sendComm, int* nqps) { - struct ncclIbCastSchedState state; - ThreadResult result = WorkerCastGetSchedState(sendComm, &state); - if (!result.ok) return result; - if (state.nqps <= 0) { - result.ok = false; - result.msg = "scheduler reports nqps=" + std::to_string(state.nqps); - return result; - } - *nqps = state.nqps; - return result; - } -#endif /* ENABLE_FAULT_INJECTION */ - ncclResult_t InitNetIbCtx(void** ctxOut) { ncclNetCommConfig_t commConfig = {}; commConfig.trafficClass = NCCL_NET_TRAFFIC_CLASS_UNDEF; diff --git a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py index 051e51ba5b4a..af8f1d9585b2 100644 --- a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py +++ b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py @@ -1045,10 +1045,12 @@ def run_test(self, test_config, suite_config): # num_gpus of 0 declares a test that needs no GPU, which only makes sense # on one node: multi-node placement builds "host:{num_gpus}" and # "--map-by ppr:{num_gpus}:node" from it, and zero there would launch - # nothing. Fail this test rather than emit that. - if num_nodes > 1 and num_gpus < 1: - msg = (f"Invalid num_gpus for test '{test_name}': {num_gpus} GPUs/node cannot place " - f"ranks across {num_nodes} nodes") + # nothing. The same goes for the resolved rank count, which "auto" derives + # from num_gpus, so zero GPUs would ask mpirun for -np 0. Fail the test + # rather than launch nothing and call it a pass. + if num_ranks < 1 or num_nodes < 1 or num_gpus < 0 or (num_nodes > 1 and num_gpus < 1): + msg = (f"Invalid placement for test '{test_name}': {num_ranks} rank(s) over " + f"{num_nodes} node(s) at {num_gpus} GPU(s)/node") print(f"ERROR: {msg}") return { "name": test_name, From ec07f054d20bc0b8f3e01be737b7c7ecdcb6fa34 Mon Sep 17 00:00:00 2001 From: Ilya Kosarev Date: Fri, 21 Aug 2026 10:06:37 +0000 Subject: [PATCH 6/7] net-ib: keep the fault-injection guard closing its own block Moving the scheduler query out of the fault-injection section took the block's #endif with it, so that directive closed MPI_TESTS_ENABLED and the fault-injection #if swallowed the rest of the class including its closing brace. Builds with fault injection enabled compiled anyway, which is every configuration on the cluster; a -DFAULT_INJECTION=OFF build saw an unterminated class. Co-authored-by: Cursor --- projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp b/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp index 6865887ceaff..c5d7a437da95 100644 --- a/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp +++ b/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp @@ -1192,7 +1192,6 @@ class NetIbMPITest : public MPITestBase { *nqps = state.nqps; return result; } -#endif /* ENABLE_FAULT_INJECTION */ // Warm the scheduler up (it initializes on the first send) and arm equal // weights. nqps comes from the environment so both ranks agree without a @@ -1551,6 +1550,7 @@ class NetIbMPITest : public MPITestBase { } return WorkerSendRecvPattern(rank, pair, buffer, size, tag, mhandle, seed, timeoutMs); } +#endif /* ENABLE_FAULT_INJECTION */ ncclResult_t InitNetIbCtx(void** ctxOut) { ncclNetCommConfig_t commConfig = {}; From ff81447513cfac519142a61d7b95bfbe881fa39c Mon Sep 17 00:00:00 2001 From: Ilya Kosarev Date: Fri, 21 Aug 2026 11:33:28 +0000 Subject: [PATCH 7/7] net-ib: stop the serialization gate from measuring its own poll interval ThreadedProgressDoesNotSerialize used the harness waits, which back off 10 ms between polls -- both the completion wait and the isend retry that waits for the receiver's FIFO slot. A 256-byte transfer needs neither: every transfer cost one full interval, so the single-worker phase measured 10.08 ms per transfer and the phase timings reported the interval instead of the plugin. Nothing shorter than 10 ms of held lock could be seen, and the self-check phase passed only because the test's own mutex spanned a sleeping wait. Polling without sleeping puts a transfer at about 7 us, and the measurement changes character: a healthy build reports 1.02 to 1.16 at four workers and 1.65 to 2.11 at sixteen, while the same build with RCCL_IB_FAULT_INJECTION set, whose ops shims take a process-global mutex per poll_cq, reports 1.74 to 1.94 and 5.39 to 6.20. Those ranges do not overlap, so the earlier conclusion that the shims cannot be seen by timing was an artifact of the backoff. The budget stays at 0.6 x N so a loaded node cannot fail the suite, and the whole test now takes two seconds instead of four and a half minutes. Only this gate polls without sleeping; the functional tests keep the backoff, which costs them nothing and does not burn a core per waiting worker. Co-authored-by: Cursor --- .../transport/NetIbMPI/NetIbMPITestBase.hpp | 59 +++++++++++++++++-- .../test/transport/NetIbMPI/StressTests.cpp | 34 +++++++---- 2 files changed, 77 insertions(+), 16 deletions(-) diff --git a/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp b/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp index c5d7a437da95..def0a76ab73c 100644 --- a/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp +++ b/projects/rccl/test/transport/NetIbMPI/NetIbMPITestBase.hpp @@ -1034,10 +1034,18 @@ class NetIbMPITest : public MPITestBase { return result; } + // isend returns success with a null request until the receiver's FIFO slot is + // available, so the caller has to retry. busyPoll retries without sleeping, + // for timing tests: with the backoff, waiting for the slot costs one poll + // interval and that interval, not the plugin, is what a measurement then + // reports. ThreadResult WorkerPostSend(void* sendComm, void* data, size_t size, int tag, - void* mhandle, void** request) { + void* mhandle, void** request, bool busyPoll = false, + int timeoutMs = kDefaultTimeoutMs) { ThreadResult result; int attempts = 0; + const auto deadline = + std::chrono::steady_clock::now() + std::chrono::milliseconds(timeoutMs); do { if (PostSend(sendComm, data, size, tag, mhandle, request) != ncclSuccess) { result.ok = false; @@ -1045,6 +1053,14 @@ class NetIbMPITest : public MPITestBase { return result; } if (*request != nullptr) break; + if (busyPoll) { + if (std::chrono::steady_clock::now() >= deadline) { + result.ok = false; + result.msg = "isend kept returning a NULL request until the timeout"; + return result; + } + continue; + } if (++attempts >= kMaxRetryAttempts) { result.ok = false; result.msg = "isend returned a NULL request after retries"; @@ -1079,19 +1095,50 @@ class NetIbMPITest : public MPITestBase { return result; } + // Same wait without the 10 ms backoff between polls. A small transfer + // completes in tens of microseconds, so the sleeping wait spends one full + // interval on every one of them: measurements built on it report the poll + // interval rather than anything the plugin did, and contention inside a + // plugin call disappears under it. Only timing tests should use this -- it + // burns a core per waiting worker, which is why the functional tests keep + // the sleeping version. + ThreadResult WorkerWaitBusy(void* request, int* sizes, int timeoutMs = kDefaultTimeoutMs) { + ThreadResult result; + const auto deadline = + std::chrono::steady_clock::now() + std::chrono::milliseconds(timeoutMs); + int done = 0; + while (!done) { + if (TestRequest(request, &done, sizes) != ncclSuccess) { + result.ok = false; + result.msg = "test failed while polling for completion"; + return result; + } + if (done) break; + if (std::chrono::steady_clock::now() >= deadline) { + result.ok = false; + result.msg = "request did not complete within the timeout"; + return result; + } + } + return result; + } + // One transfer on the worker's own connection, without touching buffer // contents. For GPU buffers, or when the caller verifies the payload itself. ThreadResult WorkerSendRecvRaw(int rank, ConnectionPair& pair, void* buffer, size_t size, int tag, void* mhandle, int timeoutMs = kDefaultTimeoutMs, - int* receivedSize = nullptr) { + int* receivedSize = nullptr, bool busyPoll = false) { void* request = nullptr; - ThreadResult result = (rank == 0) - ? WorkerPostRecv(pair.recvComm, buffer, size, tag, mhandle, &request) - : WorkerPostSend(pair.sendComm, buffer, size, tag, mhandle, &request); + ThreadResult result = + (rank == 0) + ? WorkerPostRecv(pair.recvComm, buffer, size, tag, mhandle, &request) + : WorkerPostSend(pair.sendComm, buffer, size, tag, mhandle, &request, busyPoll, + timeoutMs); if (!result.ok) return result; int sizes[1] = {0}; - result = WorkerWait(request, sizes, timeoutMs); + result = busyPoll ? WorkerWaitBusy(request, sizes, timeoutMs) + : WorkerWait(request, sizes, timeoutMs); if (!result.ok) return result; if (receivedSize) *receivedSize = sizes[0]; return result; diff --git a/projects/rccl/test/transport/NetIbMPI/StressTests.cpp b/projects/rccl/test/transport/NetIbMPI/StressTests.cpp index b7b96e1d1279..d74f0c31e950 100644 --- a/projects/rccl/test/transport/NetIbMPI/StressTests.cpp +++ b/projects/rccl/test/transport/NetIbMPI/StressTests.cpp @@ -2128,13 +2128,24 @@ TEST_F(NetIbMPITest, SetNetAttrNoOp) { // transfer, and the same measurement has to notice it -- otherwise the verdict // on the plugin means nothing. // -// What it does and does not see. A lock held across a progress call serializes -// the call itself and shows up here immediately. A lock taken and released -// around one, like the two acquisitions shimPollCq adds per poll_cq, does not: -// measured at four and sixteen workers, with the ops shims installed and -// without, the factor was 1.00 either way, because the critical section is a map -// lookup next to a driver call. That leak is caught deterministically instead, by -// FaultInjectionShimsAbsentUnlessRequested. +// Both waits poll without sleeping, deliberately: the completion wait, and the +// isend retry that waits for the receiver's FIFO slot. The harness backs off +// 10 ms in either place, which is longer than everything being measured -- with +// the backoff a transfer here costs 10 ms and the phase reports the poll interval +// rather than the plugin, and nothing shorter than the interval can be seen at +// all. Without it a transfer costs about 7 us, so contention has somewhere to +// show up. The price is a core per waiting worker for the couple of seconds this +// test runs, which is why the functional tests keep the sleeping waits. +// +// What it sees, measured on mia1 at two ranks: a healthy build reports 1.02 to +// 1.16 at four workers and 1.65 to 2.11 at sixteen, the growth being the NIC and +// the cores rather than a lock. With RCCL_IB_FAULT_INJECTION set, whose ops shims +// take one process-global mutex per poll_cq, the same measurement reports 1.74 to +// 1.94 and 5.39 to 6.20: separated from the healthy range with no overlap. The +// budget stays at 0.6 x N so a loaded node cannot fail the suite, which means the +// shims pass it -- they are gated deterministically by +// FaultInjectionShimsAbsentUnlessRequested instead, and this test's own numbers +// are the evidence it would see a lock that mattered. // ============================================================================= TEST_F(NetIbMPITest, ThreadedProgressDoesNotSerialize) { ASSERT_TRUE(validateTestPrerequisites(kExactTwoProcesses, kExactTwoProcesses, @@ -2153,8 +2164,11 @@ TEST_F(NetIbMPITest, ThreadedProgressDoesNotSerialize) { AssertInitAndGetDevices(nullptr); // Small messages: the point is how many transfers overlap, not bandwidth. + // The loop polls for completion without sleeping, so a transfer costs what it + // costs -- tens of microseconds -- and enough of them are needed for a phase + // to be measurable against scheduling noise. static constexpr size_t kMsgSize = 256; - static constexpr int kIters = 100; + static constexpr int kIters = 5000; static constexpr int kTimeoutMs = 60000; // Full serialization costs N times the single-worker time. Passing means // staying below 60% of that, which on healthy hardware measures near 1. @@ -2192,10 +2206,10 @@ TEST_F(NetIbMPITest, ThreadedProgressDoesNotSerialize) { if (serialize && rank == 1) { std::lock_guard lock(serializeAll); result = WorkerSendRecvRaw(rank, pair, buffer, kMsgSize, 950 + (i % 32), mhandle, - kTimeoutMs); + kTimeoutMs, nullptr, /*busyPoll=*/true); } else { result = WorkerSendRecvRaw(rank, pair, buffer, kMsgSize, 950 + (i % 32), mhandle, - kTimeoutMs); + kTimeoutMs, nullptr, /*busyPoll=*/true); } if (!result.ok) return result; }