plugin: finalize contexts on init fallback failure

src/plugin/net.cc

@@ -9,6 +9,7 @@
 #include "bootstrap.h"
 #include "checks.h"
 #include "plugin.h"
+#include "plugin_cleanup.h"
 #include "nccl_net.h"
 
 #include <string.h>
@@ -165,6 +166,7 @@ ncclResult_t ncclNetCheckDeviceVersion(struct ncclComm* comm, ncclNet_t* net, in
 static ncclResult_t ncclNetPluginInit(struct ncclComm* comm, netPluginLib_t* pluginLib) {
   int ndev;
   bool initCompleted = false;
+  bool collNetInitCompleted = false;
   // Init must be called for each new comm to set the right context
   if (pluginLib->ncclNetPluginState >= ncclNetPluginStateInitReady && pluginLib->ncclNet) {
     ncclNetCommConfig_t commConfig = {};
@@ -189,11 +191,19 @@ static ncclResult_t ncclNetPluginInit(struct ncclComm* comm, netPluginLib_t* plu
   if (pluginLib->ncclCollNetPluginState >= ncclNetPluginStateInitReady && pluginLib->ncclCollNet) {
     if (pluginLib->ncclCollNet->init(&comm->collNetContext, comm->commHash, ncclDebugLog) != ncclSuccess) {
       pluginLib->ncclCollNetPluginState = ncclNetPluginStateDisabled;
+    } else {
+      collNetInitCompleted = true;
     }
   }
   // Detection of the devices is only done when the plugin is being initialized the first time
   if (pluginLib->ncclCollNetPluginState == ncclNetPluginStateInitReady && pluginLib->ncclCollNet) {
     if (pluginLib->ncclCollNet->devices(&ndev) != ncclSuccess || ndev <= 0) {
+      if (collNetInitCompleted) {
+        ncclResult_t finalizeRet = ncclPluginFinalizeContext(pluginLib->ncclCollNet->finalize, &comm->collNetContext);
+        if (finalizeRet != ncclSuccess) {
+          WARN("NET/Plugin: CollNet finalize failed during init fallback cleanup: %d", finalizeRet);
+        }
+      }
       pluginLib->ncclCollNetPluginState = ncclNetPluginStateDisabled;
     } else {
       pluginLib->collNetPhysDevs = ndev;

src/plugin/plugin_cleanup.h

@@ -0,0 +1,20 @@
+/*************************************************************************
+ * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * See LICENSE.txt for more license information
+ *************************************************************************/
+
+#ifndef NCCL_PLUGIN_CLEANUP_H_
+#define NCCL_PLUGIN_CLEANUP_H_
+
+typedef ncclResult_t (*ncclPluginFinalizeFn_t)(void* ctx);
+
+static inline ncclResult_t ncclPluginFinalizeContext(ncclPluginFinalizeFn_t finalize, void** ctx) {
+  if (*ctx == nullptr) return ncclSuccess;
+  ncclResult_t ret = finalize(*ctx);
+  *ctx = nullptr;
+  return ret;
+}
+
+#endif

src/plugin/rma.cc

@@ -9,6 +9,7 @@
 #include "bootstrap.h"
 #include "checks.h"
 #include "plugin.h"
+#include "plugin_cleanup.h"
 #include "nccl_rma.h"
 #include "gin/gin_host_proxy.h"
 
@@ -105,14 +106,23 @@ static ncclResult_t ncclRmaPluginLoad(rmaPluginLib_t* pluginLib) {
 
 static ncclResult_t ncclRmaPluginInit(struct ncclComm* comm, rmaPluginLib_t* pluginLib) {
   int ndev;
+  bool rmaInitCompleted = false;
   // Init must be called for each new comm to set the right context
   if (pluginLib->state >= ncclRmaPluginStateInitReady && pluginLib->ncclRma) {
     if (pluginLib->ncclRma->init(&comm->rmaContext, comm->commHash, ncclDebugLog) != ncclSuccess) {
       pluginLib->state = ncclRmaPluginStateDisabled;
+    } else {
+      rmaInitCompleted = true;
     }
   }
   if (pluginLib->state == ncclRmaPluginStateInitReady && pluginLib->ncclRma) {
     if (pluginLib->ncclRma->devices(&ndev) != ncclSuccess || ndev <= 0) {
+      if (rmaInitCompleted) {
+        ncclResult_t finalizeRet = ncclPluginFinalizeContext(pluginLib->ncclRma->finalize, &comm->rmaContext);
+        if (finalizeRet != ncclSuccess) {
+          WARN("RMA/Plugin: finalize failed during init fallback cleanup: %d", finalizeRet);
+        }
+      }
       pluginLib->state = ncclRmaPluginStateDisabled;
     } else {
       pluginLib->physDevs = ndev;

src/plugin/rma/rma_v13.cc

@@ -9,22 +9,37 @@
 #include "nccl_rma.h"
 #include "checks.h"
 #include "os.h"
+#include "plugin_cleanup.h"
 #include <string.h>
 
 static ncclRma_v13_t* ncclRma_v13;
 static ncclRma_t ncclRma;
 
 static ncclResult_t ncclRma_init(void** ctx, uint64_t commId, ncclDebugLogger_t logFunction) {
-  NCCLCHECK(ncclRma_v13->init(ctx, commId, logFunction));
+  ncclResult_t ret = ncclSuccess;
+  void* tmpCtx = nullptr;
+  *ctx = nullptr;
+  NCCLCHECK(ncclRma_v13->init(&tmpCtx, commId, logFunction));
 
   // RMA plugin must report GIN proxy type
   ncclNetProperties_t props;
-  NCCLCHECK(ncclRma_v13->getProperties(0, &props));
+  NCCLCHECKGOTO(ncclRma_v13->getProperties(0, &props), ret, fail);
   if (props.netDeviceType != NCCL_NET_DEVICE_GIN_PROXY) {
     WARN("RMA v13 (%s) requires GIN PROXY type, got netDeviceType %d", ncclRma_v13->name, props.netDeviceType);
-    return ncclInternalError;
+    ret = ncclInternalError;
+    goto fail;
   }
+  *ctx = tmpCtx;
   return ncclSuccess;
+fail:
+  if (tmpCtx) {
+    ncclResult_t finalizeRet = ncclPluginFinalizeContext(ncclRma_v13->finalize, &tmpCtx);
+    if (finalizeRet != ncclSuccess) {
+      WARN("RMA v13 (%s) finalize failed during init fallback cleanup: %d", ncclRma_v13->name, finalizeRet);
+    }
+    if (ret == ncclSuccess) ret = finalizeRet;
+  }
+  return ret;
 }
 
 static ncclResult_t ncclRma_createContext(void* collComm, ncclRmaConfig_v14_t* config, void** rmaCtx) {

tests/plugin_lifecycle/Makefile

@@ -0,0 +1,22 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+CXX ?= g++
+CXXFLAGS ?= -std=c++17 -O2 -Wall -Wextra -Werror
+
+BUILD_DIR ?= ../../build/test/plugin_lifecycle
+TEST_BIN := $(BUILD_DIR)/plugin_lifecycle_test
+
+.PHONY: all run clean
+
+all: $(TEST_BIN)
+
+$(TEST_BIN): plugin_lifecycle_test.cc
+	@mkdir -p $(BUILD_DIR)
+	$(CXX) $(CXXFLAGS) $< -o $@
+
+run: $(TEST_BIN)
+	$(TEST_BIN)
+
+clean:
+	rm -rf $(BUILD_DIR)