alex-breslow-amd · Nov 10, 2026 · Oct 25, 2026 · Oct 25, 2026 · Oct 29, 2026 · Oct 29, 2026
@@ -46,6 +46,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * Hierarchical symmetric `ncclAllGather` for multi-node using GIN and LSA (`AllGather_RailRing_LsaST`): a GIN ring over the rails team moves data between nodes and LSA stores broadcast it to the other GPUs in each node.
 * nccl4py: per-call collective configuration (`NCCLCollConfig`, `VendorOption`) including the launch completion event, communicator properties (`NCCLCommProperties`), and the `GIN_ONLY` window flag. On ROCm, the HIP `Event` shim provides only an event handle; `record()`, `sync()` and `query` are not implemented.
 * `RCCL_TUNER_EMBEDDED_CONFIG` (default `1`): the built-in CSV tuner configurations in `tuner/*.csv` are now compiled into `librccl.so` and used as the stock tuning defaults, so a cluster no longer has to stage the CSV on every node or on NFS. Set `0` to ignore the embedded configurations. `NCCL_TUNER_CONFIG_FILE` and the on-disk tuner directories still take precedence, and `NCCL_TUNER_PLUGIN=none` still disables the tuner entirely.
+* `NCCL_GIN_TYPES` (RCCL-only): a comma-separated list of GIN backend types that may be active together on a communicator, so DevComms on different backends can share it. For example, `6,7` enables rocSHMEM GDA alongside Anvil SDMA, and `6,2` enables it alongside the IB proxy. Types not in the list are not registered. `NCCL_GIN_TYPE` still forces exactly one backend and overrides the list when both are set. When Anvil SDMA and rocSHMEM GDA are both active, Anvil SDMA is the default backend; a DevComm selects rocSHMEM GDA by requesting `ginType = NCCL_GIN_TYPE_ROCSHMEM_GDA`. Without either variable, the active backends are unchanged and rocSHMEM GDA stays off.
 * `RCCL_TUNING_PARTITION_MODE` (`auto`, `SPX` or `DPX`; default `auto`): selects the gfx1250 tuning table for the GPU's compute-partition mode. `auto` uses the DPX table when the GPU reports DPX, and the SPX table otherwise. A communicator uses the DPX table only when every rank selects it. The DPX table is currently a copy of the SPX table, so tuning is unchanged.
 * `RCCL_CE_AR_2SHOT_MAX_BYTES` (default `-1`): overrides the 2-shot AllReduce size cap from the arch table. Set to a positive value to override `ceNonRegMax[AR]`. Renamed from `RCCL_CE_AR_MAX_MSG_BYTES`.
 * `RCCL_CE_AR_2SHOT_MIN_BYTES` (default `-1`): overrides the minimum message size for the CE 2-shot AllReduce path. When `-1`, resolved from `ceNonRegMin[AllReduce]` in the per-arch table.

@@ -772,9 +772,22 @@ in the following table.
         | ``3``: GDAKI
         | ``4``: GPI
         | ``5``: EFA GDA
-        | ``6``: rocSHMEM GDA (required for that backend to initialize)
+        | ``6``: rocSHMEM GDA (this value, or ``6`` in ``NCCL_GIN_TYPES``, is
+          required for that backend to initialize)
         | ``7``: Anvil SDMA (also accepts unset)
 
+    * - | ``NCCL_GIN_TYPES``
+        | Lists the GIN backend types that may be active together, so DevComms
+          on one communicator can use different backends. RCCL-specific.
+          Ignored when ``NCCL_GIN_TYPE`` is set.
+      - | Comma-separated list of ``NCCL_GIN_TYPE`` values, for example
+          ``6,7`` (rocSHMEM GDA and Anvil SDMA) or ``6,2`` (rocSHMEM GDA and
+          the proxy). Types not in the list are not registered.
+        | When Anvil SDMA and rocSHMEM GDA are both active, Anvil SDMA is the
+          default backend. A DevComm uses rocSHMEM GDA by requesting
+          ``NCCL_GIN_TYPE_ROCSHMEM_GDA`` in ``ncclDevCommRequirements::ginType``.
+        | Default: unset (every backend except rocSHMEM GDA may be active)
+
     * - | ``NCCL_RMA_PLUGIN``
         | Selects external one-sided RMA plugins, which are also the backend
           the built-in GIN proxy forwards to.

@@ -225,6 +225,7 @@ set(SRC_FILES
   include/gin/gin_host_proxy.h
   include/gin/gin_host_rocshmem_gda.h
   include/gin/gin_rocshmem_gda_factory.h
+  include/gin/gin_type_env.h
   include/dv_dlopen.h
   include/gin/gin_host_anvil_sdma.h
   include/gin/gin_anvil_sdma_factory.h

@@ -50,6 +50,7 @@ for the alltoall\_wg offload path.  Does not enable GIN plugins.
 |----------|---------|-------------|
 | `NCCL_GIN_ENABLE` | `1` | Enable/disable GIN (all backends) |
 | `NCCL_GIN_TYPE` | (auto) | Force a specific backend: `2`=IB proxy, `6`=rocSHMEM GDA, `7`=Anvil SDMA. **Not value-compatible with 2.30.7** (see below). |
+| `NCCL_GIN_TYPES` | (unset) | RCCL-only. Comma-separated backend types that may be active together, e.g. `6,7` (GDA and SDMA) or `6,2` (GDA and IB proxy). Ignored when `NCCL_GIN_TYPE` is set. See [Mixing backends](#mixing-backends). |
 | `NCCL_GIN_ANVIL_SDMA_THRESHOLD` | `128` | Minimum message size (bytes) to use SDMA; smaller messages fall back to the IB proxy |
 | `NCCL_GIN_ANVIL_SDMA_FUSED_SIGNAL` | `0` | Enable fused signal mode for SDMA (experimental) |
 | `NCCL_GIN_ANVIL_SDMA_CONN_CHECK` | `1` | Validate LSA signal connectivity on first bind (LSA team only; skipped under `NCCL_GIN_CONNECTION_RAIL`); set exactly `0` to bypass |
@@ -86,6 +87,27 @@ docker run -it --rm --shm-size 64G \
 
 To use the rocSHMEM GDA backend instead, set `NCCL_GIN_TYPE=6`.
 
+### Mixing backends
+
+With `NCCL_GIN_TYPE` unset, Anvil SDMA and the IB proxy are both active by
+default, and rocSHMEM GDA is off. `NCCL_GIN_TYPES` replaces that default set
+with the listed types, so one communicator can host DevComms on different
+backends:
+
+| `NCCL_GIN_TYPES` | Active backends | Default (`backends[0]`) |
+|------------------|-----------------|-------------------------|
+| `6,7` | rocSHMEM GDA, Anvil SDMA | Anvil SDMA |
+| `6,2` | rocSHMEM GDA, IB proxy | rocSHMEM GDA |
+| `6,7,2` | all three | Anvil SDMA |
+
+The default order is Anvil SDMA, then rocSHMEM GDA, then the IB proxy. A
+DevComm that leaves `ncclDevCommRequirements::ginType` as
+`NCCL_GIN_TYPE_NONE` gets the first backend that can serve it; set `ginType`
+to `NCCL_GIN_TYPE_ROCSHMEM_GDA` (or `NCCL_GIN_TYPE_ANVIL_SDMA`,
+`NCCL_GIN_TYPE_PROXY`) to pick one. Each backend still needs its own platform
+requirements below, and a backend that fails to initialize is dropped from the
+set.
+
 ### Platform requirements
 
 rocSHMEM GDA (`NCCL_GIN_TYPE=6`) requires a supported GDA NIC provider (mlx5, bnxt,

@@ -20,6 +20,7 @@
 #include "nccl_device/gin/anvil_sdma/gin_anvil_sdma_device_host_common.h"
 #include "nccl_device/gin/anvil_sdma/gin_anvil_ipc_table.h"
 #include "gin/gin_anvil_sdma_factory.h"
+#include "gin/gin_type_env.h"
 #include <hip/hip_runtime.h>
 #include <algorithm>
 #include <atomic>
@@ -194,8 +195,7 @@ int ncclGinAnvilPluginTestLiveInitCtxCount(void) {
 }
 
 static ncclResult_t ginAnvilInit(void** ctx, uint64_t commId, ncclDebugLogger_t logFunction) {
-  const char* gin_type = getenv("NCCL_GIN_TYPE");
-  if (gin_type && atoi(gin_type) != NCCL_NET_DEVICE_GIN_ANVIL_SDMA) return ncclInternalError;
+  if (!ncclGinTypeEnabled(NCCL_NET_DEVICE_GIN_ANVIL_SDMA, /*enabledByDefault=*/true)) return ncclInternalError;
   if (gin_anvil_sdma_probe() <= 0) return ncclInternalError;
   auto* ictx = new ginAnvilInitCtx{};
   ginAnvilLiveInitCtxs.fetch_add(1, std::memory_order_relaxed);

@@ -20,6 +20,7 @@
 #include "bootstrap.h"
 #include "nccl_device/gin/rocshmem_gda/gin_rocshmem_device_host_common_gda.h"
 #include "gin/gin_rocshmem_gda_factory.h"
+#include "gin/gin_type_env.h"
 #include <hip/hip_runtime.h>
 
 // RCCL env-var parameters for traffic class (defined in transport/net_ib/)
@@ -98,8 +99,7 @@ static int ginGdaBootstrapAllgather(void* ctx, void* buf, size_t perRankSize) {
 ///////////////////////////////////////////////////////////////////////////////
 
 static ncclResult_t ginRocshmemGdaInit(void** ctx, uint64_t commId, ncclDebugLogger_t logFunction) {
-  const char* gin_type = getenv("NCCL_GIN_TYPE");
-  if (!gin_type || atoi(gin_type) != NCCL_NET_DEVICE_GIN_ROCSHMEM_GDA) return ncclInternalError;
+  if (!ncclGinTypeEnabled(NCCL_NET_DEVICE_GIN_ROCSHMEM_GDA, /*enabledByDefault=*/false)) return ncclInternalError;
   if (rocshmem_gin_probe_devices() <= 0) return ncclInternalError;
   // Allocate init context; bootstrap filled by ncclGinRocshmemSetInitContext
   struct ginRocshmemInitCtx* ictx = new ginRocshmemInitCtx{};

@@ -0,0 +1,60 @@
+/*************************************************************************
+ * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#ifndef RCCL_GIN_TYPE_ENV_H_
+#define RCCL_GIN_TYPE_ENV_H_
+
+#include <cerrno>
+#include <cstdint>
+#include <cstdlib>
+
+// NCCL_GIN_TYPE forces one backend; if unset, NCCL_GIN_TYPES lists the types allowed together.
+// Uncached getenv() so tests can setenv() between cases.
+
+// Parses list into a bitmask of GIN types (bit t set for type t). Returns false if
+// any entry is not an integer in [1, 63]; valid entries are still recorded in *mask.
+static inline bool ncclGinParseTypeList(const char* list, uint64_t* mask) {
+  bool valid = true;
+  *mask = 0;
+  const char* p = list;
+  while (*p) {
+    while (*p == ' ' || *p == '\t') p++;
+    char* end = nullptr;
+    errno = 0;
+    long type = strtol(p, &end, 10);
+    while (*end == ' ' || *end == '\t') end++;
+    if (end == p || errno != 0 || type < 1 || type > 63 || (*end != ',' && *end != '\0')) {
+      valid = false;
+      while (*end != ',' && *end != '\0') end++;
+    } else {
+      *mask |= (uint64_t)1 << type;
+    }
+    p = (*end == ',') ? end + 1 : end;
+  }
+  return valid;
+}
+
+// Returns true and fills *mask when NCCL_GIN_TYPE is unset and NCCL_GIN_TYPES names
+// at least one valid type. Otherwise the list does not apply.
+static inline bool ncclGinTypeListMask(uint64_t* mask) {
+  if (getenv("NCCL_GIN_TYPE")) return false;
+  const char* list = getenv("NCCL_GIN_TYPES");
+  if (!list) return false;
+  ncclGinParseTypeList(list, mask);
+  return *mask != 0;
+}
+
+// Whether the backend of the given type may initialize. enabledByDefault is its
+// behavior when neither NCCL_GIN_TYPE nor a usable NCCL_GIN_TYPES is set.
+static inline bool ncclGinTypeEnabled(int type, bool enabledByDefault) {
+  const char* single = getenv("NCCL_GIN_TYPE");
+  if (single) return atoi(single) == type;
+  uint64_t mask;
+  if (ncclGinTypeListMask(&mask)) return (mask >> type) & 1;
+  return enabledByDefault;
+}
+
+#endif  // RCCL_GIN_TYPE_ENV_H_
@@ -14,13 +14,15 @@
 #include "nccl_gin.h"
 #include "graph/topo.h"
 #include "gin/gin_host_proxy.h"
+#include "gin/gin_type_env.h"
 #ifdef ENABLE_ROCSHMEM_GIN
 #include "gin/gin_host_rocshmem_gda.h"
 #include "gin/gin_host_anvil_sdma.h"
 #endif
 
 #include <string.h>
 #include <errno.h>
+#include <stdlib.h>
 #include <mutex>
 
 typedef ncclGin_t* getNcclGin_t(void* ginPluginLib);
@@ -175,6 +177,13 @@ static ncclResult_t ncclGinPluginAssignToComm(struct ncclComm* comm, int pluginI
       return ncclSuccess;
     }
 
+    uint64_t ginTypeMask;
+    if (ginType == -1 && ncclGinTypeListMask(&ginTypeMask) && !(ginTypeMask & BIT(props.netDeviceType))) {
+      INFO(NCCL_INIT | NCCL_NET, "GIN/Plugin: Skipping plugin %s index %d type %d: not in NCCL_GIN_TYPES=%s", gin->name,
+           pluginIndex, props.netDeviceType, getenv("NCCL_GIN_TYPES"));
+      return ncclSuccess;
+    }
+
     for (int idx = 0; idx < pluginIndex; idx++) {
       if (pluginLibs[idx].state >= ncclGinPluginStateInitReady) {
         ncclNetProperties_t prev_props;
@@ -240,6 +249,24 @@ static void initPluginLibsOnceFunc() {
   int pluginCounter = 0;
 
   memset(pluginLibs, 0, NCCL_GIN_MAX_PLUGINS * sizeof(ginPluginLib_t));
+
+  // Same getenv() source as ncclGinTypeListMask / ncclGinTypeEnabled, so an
+  // NCCL_ENV_PLUGIN value cannot be logged here and then ignored at init.
+  const char* envGinTypes = getenv("NCCL_GIN_TYPES");
+  if (envGinTypes) {
+    uint64_t ginTypeMask;
+    bool valid = ncclGinParseTypeList(envGinTypes, &ginTypeMask);
+    if (getenv("NCCL_GIN_TYPE")) {
+      INFO(NCCL_ENV | NCCL_NET, "NCCL_GIN_TYPES=%s ignored: NCCL_GIN_TYPE is set", envGinTypes);
+    } else if (ginTypeMask == 0) {
+      WARN("NCCL_GIN_TYPES=%s names no valid GIN type, ignoring it", envGinTypes);
+    } else if (!valid) {
+      WARN("NCCL_GIN_TYPES=%s contains an invalid entry, ignoring that entry", envGinTypes);
+    } else {
+      INFO(NCCL_ENV | NCCL_NET, "NCCL_GIN_TYPES set by environment to %s", envGinTypes);
+    }
+  }
+
   envGinPlugin = ncclGetEnv("NCCL_GIN_PLUGIN");
   if (envGinPlugin) {
     INFO(NCCL_ENV | NCCL_NET, "NCCL_GIN_PLUGIN set by environment to %s", envGinPlugin);
@@ -313,7 +340,7 @@ static void initPluginLibsOnceFunc() {
   pluginCounter++;
 
 #ifdef ENABLE_ROCSHMEM_GIN
-  // Add internal rocshmem GDA plugin (device-initiated, NCCL_GIN_TYPE=6)
+  // Add internal rocshmem GDA plugin (device-initiated, NCCL_GIN_TYPE=6 or listed in NCCL_GIN_TYPES)
   pluginLibs[pluginCounter].ncclGin = &ncclGinRocshmemGdaPlugin;
   pluginLibs[pluginCounter].state = ncclGinPluginStateInitReady;
   pluginLibs[pluginCounter].version = ncclGinVersion[0];
@@ -407,16 +434,17 @@ static int ncclGetBackendPriority(struct ncclGinBackendState* backend) {
   case NCCL_GIN_TYPE_GDAKI:
     return 0;
 #ifdef ENABLE_ROCSHMEM_GIN
-  // [RCCL] Device-initiated AMD backends rank with GDAKI, ahead of the host proxy.
-  case NCCL_GIN_TYPE_ROCSHMEM_GDA:
+  // [RCCL] Distinct priorities: qsort is unstable, so ties would make the default backend arbitrary.
   case NCCL_GIN_TYPE_ANVIL_SDMA:
     return 0;
+  case NCCL_GIN_TYPE_ROCSHMEM_GDA:
+    return 1;
 #endif
   case NCCL_GIN_TYPE_PROXY:
-    return 1;
+    return 2;
   case NCCL_GIN_TYPE_EFA_GDA:
   case NCCL_GIN_TYPE_GPI:
-    return 2;
+    return 3;
   default:
     return INT_MAX;
   }

@@ -156,6 +156,41 @@ TEST_F(GinAnvilPluginTest, Init_WrongGinType) {
   EXPECT_EQ(ictx, nullptr);
 }
 
+// AICOMRCCL-3418: with NCCL_GIN_TYPE unset, NCCL_GIN_TYPES decides.
+TEST_F(GinAnvilPluginTest, Init_ListedInGinTypes) {
+  ScopedEnv ty("NCCL_GIN_TYPE", nullptr);
+  ScopedEnv list("NCCL_GIN_TYPES", "6,7");
+  void* ictx = nullptr;
+  ASSERT_EQ(plugin_.init(&ictx, 0, nullptr), ncclSuccess);
+  ASSERT_NE(ictx, nullptr);
+  EXPECT_EQ(plugin_.finalize(ictx), ncclSuccess);
+}
+
+TEST_F(GinAnvilPluginTest, Init_NotListedInGinTypes) {
+  ScopedEnv ty("NCCL_GIN_TYPE", nullptr);
+  ScopedEnv list("NCCL_GIN_TYPES", "6,2");
+  void* ictx = nullptr;
+  EXPECT_EQ(plugin_.init(&ictx, 0, nullptr), ncclInternalError);
+  EXPECT_EQ(ictx, nullptr);
+}
+
+TEST_F(GinAnvilPluginTest, Init_GinTypeOverridesGinTypes) {
+  ScopedEnv list("NCCL_GIN_TYPES", "6,2");
+  void* ictx = nullptr;
+  ASSERT_EQ(plugin_.init(&ictx, 0, nullptr), ncclSuccess);
+  ASSERT_NE(ictx, nullptr);
+  EXPECT_EQ(plugin_.finalize(ictx), ncclSuccess);
+}
+
+TEST_F(GinAnvilPluginTest, Init_OnByDefault) {
+  ScopedEnv ty("NCCL_GIN_TYPE", nullptr);
+  ScopedEnv list("NCCL_GIN_TYPES", nullptr);
+  void* ictx = nullptr;
+  ASSERT_EQ(plugin_.init(&ictx, 0, nullptr), ncclSuccess);
+  ASSERT_NE(ictx, nullptr);
+  EXPECT_EQ(plugin_.finalize(ictx), ncclSuccess);
+}
+
 // G2: probe failure.
 TEST_F(GinAnvilPluginTest, Init_ProbeFails) {
   GinAnvilPluginStubs::SetProbeResult(0);

@@ -11,6 +11,8 @@
 
 #include "gin_rocshmem_gda_plugin_test_stubs.h"
 
+#include "../common/ScopedEnv.hpp"
+
 #include "gin/gin_host_rocshmem_gda.h"
 #include "comm.h"
 #include "nccl_device/gin/rocshmem_gda/gin_rocshmem_device_host_common_gda.h"
@@ -77,6 +79,41 @@ class GinRocshmemGdaPluginTest : public ::testing::Test {
   }
 };
 
+// AICOMRCCL-3418: GDA stays off unless NCCL_GIN_TYPE=6, or NCCL_GIN_TYPE is unset
+// and NCCL_GIN_TYPES lists 6.
+TEST_F(GinRocshmemGdaPluginTest, Init_OffByDefault) {
+  ScopedEnv ty("NCCL_GIN_TYPE", nullptr);
+  ScopedEnv list("NCCL_GIN_TYPES", nullptr);
+  EXPECT_EQ(plugin_.init(&ictx_, 0, nullptr), ncclInternalError);
+  EXPECT_EQ(ictx_, nullptr);
+}
+
+TEST_F(GinRocshmemGdaPluginTest, Init_ListedInGinTypes) {
+  ScopedEnv ty("NCCL_GIN_TYPE", nullptr);
+  ScopedEnv list("NCCL_GIN_TYPES", "6,7");
+  ASSERT_EQ(plugin_.init(&ictx_, 0, nullptr), ncclSuccess);
+  EXPECT_NE(ictx_, nullptr);
+}
+
+TEST_F(GinRocshmemGdaPluginTest, Init_NotListedInGinTypes) {
+  ScopedEnv ty("NCCL_GIN_TYPE", nullptr);
+  ScopedEnv list("NCCL_GIN_TYPES", "2,7");
+  EXPECT_EQ(plugin_.init(&ictx_, 0, nullptr), ncclInternalError);
+  EXPECT_EQ(ictx_, nullptr);
+}
+
+TEST_F(GinRocshmemGdaPluginTest, Init_GinTypeOverridesGinTypes) {
+  ScopedEnv list("NCCL_GIN_TYPES", "2,7");
+  ASSERT_EQ(plugin_.init(&ictx_, 0, nullptr), ncclSuccess);
+  EXPECT_NE(ictx_, nullptr);
+
+  ScopedEnv ty("NCCL_GIN_TYPE", "7");
+  ScopedEnv listWithGda("NCCL_GIN_TYPES", "6,7");
+  void* other = nullptr;
+  EXPECT_EQ(plugin_.init(&other, 0, nullptr), ncclInternalError);
+  EXPECT_EQ(other, nullptr);
+}
+
 TEST_F(GinRocshmemGdaPluginTest, CreateContext_StripesSignalsAndCountersPerContext) {
   constexpr int kRanks = 2, kRank = 1, kContexts = 3, kSignals = 4, kCounters = 2;
   connect(kRanks, kRank);

@@ -57,6 +57,7 @@ struct FakeGinVtable {
   // finalize() frees the ctx it is handed. Leave init null in the vtable when
   // initNullFn is set.
   int ndev = 0;
+  int netDeviceType = 0;  // reported by getProperties()
   void* lastCtx = nullptr;
   bool initNullFn = false;
   bool lastError = false;
@@ -109,7 +110,10 @@ struct FakeGinVtable {
   static ncclResult_t GetProperties(int /*dev*/, ncclNetProperties_t* props) {
     FakeGinVtable& self = current();
     ++self.getPropertiesCalls;
-    if (self.getPropertiesResult == ncclSuccess && props) std::memset(props, 0, sizeof(*props));
+    if (self.getPropertiesResult == ncclSuccess && props) {
+      std::memset(props, 0, sizeof(*props));
+      props->netDeviceType = static_cast<decltype(props->netDeviceType)>(self.netDeviceType);
+    }
     return self.getPropertiesResult;
   }