[nvbugs/6062416][fix] Cache per-comm NCCL symmetric-memory unavailability

After the collective ncclAllReduce(min) inside allocateAndRegisterBuffer agrees that
at least one rank could not ncclMemAlloc, all subsequent requestBuffer() calls for the
same comm short-circuit to an empty NCCLWindowBuffer instead of re-running the failing
ncclMemAlloc + rank-sync allreduce on every autotuner trial. This is cluster-safe
because the unavailable decision is driven by a collective allreduce, so all ranks
agree and remain in sync without further communication. The cache is invalidated when
the comm is torn down.

Signed-off-by: tensorrt-cicd <90828364+tensorrt-cicd@users.noreply.github.com>

Author: tensorrt-cicd

cpp/tensorrt_llm/common/ncclUtils.cpp

@@ -320,6 +320,17 @@ NCCLWindowBuffer NCCLWindowAllocator::requestBuffer(ncclComm_t comm, size_t size
     // This is cheap even if no buffers exist yet - cleanup will just return early
     registerBufferCleanup(comm);
 
+    // If a previous allocateAndRegisterBuffer call collectively concluded that this comm
+    // cannot use NCCL symmetric memory, short-circuit so callers transparently fall back to
+    // regular allreduce. This avoids re-running ncclMemAlloc + the rank-sync allreduce on
+    // every autotuner trial, which would otherwise spam warnings and stress the failing path.
+    // The decision is collective (driven by an ncclAllReduce(min) inside allocateAndRegisterBuffer),
+    // so all ranks reach the same conclusion and stay in sync without further communication.
+    if (mSymmetricUnavailable.find(comm) != mSymmetricUnavailable.end())
+    {
+        return NCCLWindowBuffer();
+    }
+
     // Check if we have an available buffer of at least the requested size for this communicator
     // Use best-fit: find the smallest buffer that's >= requested size
     auto& commBuffers = mBufferPool[comm];
@@ -371,6 +382,13 @@ NCCLWindowBuffer NCCLWindowAllocator::requestBuffer(ncclComm_t comm, size_t size
     {
         commBuffers.push_back({buffer, true});
     }
+    else
+    {
+        // The collective allreduce inside allocateAndRegisterBuffer agreed that at least one
+        // rank could not allocate symmetric memory. Mark this comm so future requests don't
+        // retry the failing path on every autotuner trial.
+        mSymmetricUnavailable.insert(comm);
+    }
 
     return buffer;
 }
@@ -604,6 +622,7 @@ void NCCLWindowAllocator::cleanupBuffersForComm(ncclComm_t comm) noexcept
     {
         // No buffers to clean up, but mark as cleaned
         mRegisteredComms.erase(comm);
+        mSymmetricUnavailable.erase(comm);
         return;
     }
 
@@ -679,6 +698,7 @@ void NCCLWindowAllocator::cleanupBuffersForComm(ncclComm_t comm) noexcept
 
     mBufferPool.erase(commIt);
     mRegisteredComms.erase(comm);
+    mSymmetricUnavailable.erase(comm);
 }
 
 #endif // NCCL_VERSION_CODE >= NCCL_VERSION(2, 28, 0)

cpp/tensorrt_llm/common/ncclUtils.h

@@ -289,6 +289,12 @@ class NCCLWindowAllocator
     mutable std::mutex mMutex;
     std::unordered_map<ncclComm_t, std::vector<BufferEntry>> mBufferPool;
     std::unordered_set<ncclComm_t> mRegisteredComms;
+    // Comms whose symmetric memory path is known to fail collectively (e.g. H100 PCIe without
+    // NVLink fabric where ncclMemAlloc returns ncclUnhandledCudaError on at least one rank).
+    // Once recorded, subsequent requestBuffer() calls short-circuit to NCCLWindowBuffer{} so we
+    // don't repeatedly trigger the warning, the rank-sync allreduce, and the sticky-error drain
+    // for every autotuner trial.
+    std::unordered_set<ncclComm_t> mSymmetricUnavailable;
 };
 
 // RAII wrapper for NCCL window buffers