abouteiller · Oct 20, 2026 · Sep 25, 2026 · Oct 13, 2026 · Oct 13, 2026 · Oct 14, 2026
@@ -114,7 +114,7 @@ static int test_qp_creation(int rank, int nranks, MPI_Comm comm) {
   void **gpu_qps = nullptr;
 
   int rc = rocshmem_gin_create_qps(nranks, rank, mpi_allgather, &comm,
-                                    &qp_set, &gpu_qps);
+                                    /*traffic_class=*/0, &qp_set, &gpu_qps);
   if (rc != 0) {
     fprintf(stderr, "[rank %d] FAIL: rocshmem_gin_create_qps returned %d\n", rank, rc);
     return -1;
@@ -135,7 +135,7 @@ static int test_put(int rank, int nranks, MPI_Comm comm) {
   void **gpu_qps = nullptr;
 
   int rc = rocshmem_gin_create_qps(nranks, rank, mpi_allgather, &comm,
-                                    &qp_set, &gpu_qps);
+                                    /*traffic_class=*/0, &qp_set, &gpu_qps);
   if (rc != 0) {
     fprintf(stderr, "[rank %d] FAIL: QP creation failed\n", rank);
     return -1;
@@ -226,7 +226,7 @@ static int test_atomic_signal(int rank, int nranks, MPI_Comm comm) {
   void **gpu_qps = nullptr;
 
   int rc = rocshmem_gin_create_qps(nranks, rank, mpi_allgather, &comm,
-                                    &qp_set, &gpu_qps);
+                                    /*traffic_class=*/0, &qp_set, &gpu_qps);
   if (rc != 0) {
     fprintf(stderr, "[rank %d] FAIL: QP creation failed\n", rank);
     return -1;
@@ -287,7 +287,7 @@ static int test_put_signal(int rank, int nranks, MPI_Comm comm) {
   void **gpu_qps = nullptr;
 
   int rc = rocshmem_gin_create_qps(nranks, rank, mpi_allgather, &comm,
-                                    &qp_set, &gpu_qps);
+                                    /*traffic_class=*/0, &qp_set, &gpu_qps);
   if (rc != 0) {
     fprintf(stderr, "[rank %d] FAIL: QP creation failed\n", rank);
     return -1;

@@ -42,6 +42,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
   * `ncclAllGather`: DDA LL up to nRanks x 16 KiB, DDA LL128 up to 256 MiB from 16 ranks, and Ring LL128 up to 64 MiB per rank.
   * `ncclReduceScatter`: DDA LL up to nRanks x 256 KiB, DDA LL128 up to nRanks x 512 KiB, then DDA VMM up to nRanks x 2 MiB; Ring LL up to 512 KiB per rank and LL128 up to 8 MiB per rank (32 MiB for 1-byte types); Ring Simple on at most 128 channels between 8 and 32 MiB per rank unless `NCCL_MIN_NCHANNELS` or `NCCL_MAX_NCHANNELS` is set.
   * DDA VMM `ncclAllReduce` and `ncclReduceScatter` launch 2048 / nRanks blocks, clamped to 32 to 128 and never more than before.
+* The rocSHMEM GDA GIN backend (`NCCL_GIN_TYPE=6`) now sets the RoCE traffic class of its queue pairs from `NCCL_GIN_IB_TC`, then `NCCL_IB_TC`, then the communicator traffic class, matching the IB proxy backend. It previously read `ROCSHMEM_GDA_TRAFFIC_CLASS`, which no longer affects these queue pairs. The default remains traffic class 0 when none of these is set.
 
 ### Removed
 * `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it.

@@ -254,10 +254,13 @@ indexed GIN resources.
 ``NCCL_GIN_IB_TC``, added in upstream NCCL 2.30.7, applies to the connections
 of the IB proxy backend (``NCCL_GIN_TYPE=2``). The one-sided host RMA
 operations run on the same backend and use the same traffic class. The
-rocSHMEM GDA backend (``NCCL_GIN_TYPE=6``) doesn't read ``NCCL_GIN_IB_TC``,
-``NCCL_IB_TC``, or ``ginTrafficClass``; its queue pairs take their traffic
-class from ``ROCSHMEM_GDA_TRAFFIC_CLASS``. Anvil SDMA (``NCCL_GIN_TYPE=7``)
-moves data within a node without a NIC, so traffic class doesn't apply to it.
+rocSHMEM GDA backend (``NCCL_GIN_TYPE=6``) also honors it: its RoCE queue
+pairs take their traffic class from ``NCCL_GIN_IB_TC`` when set, then from
+``NCCL_IB_TC``, and only then from the host communicator traffic class.
+Unlike the IB proxy backend, the GDA backend does not consult the
+device-communicator ``ginTrafficClass``, because its queue pairs are created
+before that value is available. The ``ROCSHMEM_GDA_TRAFFIC_CLASS`` environment
+variable no longer affects RCCL-driven GIN-GDA queue pairs.
 
 The 128-byte, versioned GIN proxy descriptor and per-context proxy progress are
 internal implementation details and require no application configuration.

@@ -438,13 +438,15 @@ NCCL_GIN_IB_TC
 (since 2.30.7)
 
 Defines the InfiniBand traffic class for GPU-initiated networking (GIN)
-connections, independently of NCCL_IB_TC. GIN traffic can then use a different
+connections on both the IB proxy and rocSHMEM GDA backends, independently of
+NCCL_IB_TC. GIN traffic can then use a different
 RoCE traffic class from the collective and point-to-point connections, which
 keep using NCCL_IB_TC. The one-sided host RMA operations run on the same IB
 proxy backend and therefore also use this traffic class.
 
-In RCCL, only the IB proxy GIN backend (``NCCL_GIN_TYPE=2``) reads this variable.
-The RCCL device API and GIN how-to describes the other backends.
+In RCCL, both the IB proxy GIN backend (``NCCL_GIN_TYPE=2``) and the rocSHMEM
+GDA backend (``NCCL_GIN_TYPE=6``) read this variable. The RCCL device API and
+GIN how-to describes the backends and their precedence.
 
 Values accepted
 ^^^^^^^^^^^^^^^

@@ -22,6 +22,10 @@
 #include "gin/gin_rocshmem_gda_factory.h"
 #include <hip/hip_runtime.h>
 
+// RCCL env-var parameters for traffic class (defined in transport/net_ib/)
+extern int64_t ncclParamGinIbTc();
+extern int64_t ncclParamIbTc();
+
 // collComm: per-connection state, returned by connect()
 struct ginRocshmemGdaCollCtx {
   int nranks;
@@ -120,9 +124,22 @@ static ncclResult_t ginRocshmemGdaConnect(void* ctx, void* handles[], int nranks
   cctx->rank = rank;
   cctx->comm = ictx->comm;
 
+  // Resolve traffic class: NCCL_GIN_IB_TC > NCCL_IB_TC > host comm config > default 0.
+  // The proxy path also has a device-communicator ginTrafficClass rung, but the GDA
+  // backend creates its QPs here in connect(), before ginConfig.trafficClass is known
+  // (it only reaches createContext()), so that rung cannot be honored on this path.
+  int trafficClass = 0;
+  if (ncclParamGinIbTc() != -1) {
+    trafficClass = static_cast<int>(ncclParamGinIbTc());
+  } else if (ncclParamIbTc() != -1) {
+    trafficClass = static_cast<int>(ncclParamIbTc());
+  } else if (cctx->comm->config.trafficClass != NCCL_CONFIG_UNDEF_INT) {
+    trafficClass = cctx->comm->config.trafficClass;
+  }
+
   // Create QPs during connect (like GDAKI creates IB connections in connect)
-  int rc = rocshmem_gin_create_qps(nranks, rank, ginGdaBootstrapAllgather, cctx->comm->bootstrap, &cctx->qpSet,
-                                   &cctx->gpu_qp_ptrs);
+  int rc = rocshmem_gin_create_qps(nranks, rank, ginGdaBootstrapAllgather, cctx->comm->bootstrap, trafficClass,
+                                   &cctx->qpSet, &cctx->gpu_qp_ptrs);
   if (rc != 0) {
     WARN("GIN rocshmem-gda: failed to create QPs");
     delete cctx;
@@ -132,7 +149,7 @@ static ncclResult_t ginRocshmemGdaConnect(void* ctx, void* handles[], int nranks
   // Initialize rocshmem __constant__ device memory
   rocshmem_gin_init_constmem(rocshmem_gin_get_provider(cctx->qpSet), rank);
 
-  INFO(NCCL_INIT, "GIN rocshmem-gda: QPs created (%d ranks)", nranks);
+  INFO(NCCL_INIT, "GIN rocshmem-gda: QPs created (%d ranks, tc=%d)", nranks, trafficClass);
   *collComm = cctx;
   return ncclSuccess;
 }

@@ -157,6 +157,7 @@ struct rocshmem_gin_qp_set {
   struct ibv_mr* fetch_atomic_dummy_mr = nullptr;
 
   uint32_t sq_size = envvar::gda::sq_size;
+  uint8_t traffic_class = 0;
 };
 
 ///////////////////////////////////////////////////////////////////////////////
@@ -867,7 +868,7 @@ static int gin_modify_qps_init_to_rtr(rocshmem_gin_qp_set* set, struct gin_dest_
       attr.ah_attr.is_global = 1;
       attr.ah_attr.grh.hop_limit = 255;
       attr.ah_attr.sl = 1;
-      attr.ah_attr.grh.traffic_class = envvar::gda::traffic_class;
+      attr.ah_attr.grh.traffic_class = set->traffic_class;
       memcpy(&attr.ah_attr.grh.dgid, &remote_info[i].gid, 16);
     } else {
       attr.ah_attr.is_global = 0;
@@ -1083,14 +1084,16 @@ int rocshmem_gin_qp_set::initialize_host_gpu_qp(QueuePair* host_gpu_qp, int idx)
 ///////////////////////////////////////////////////////////////////////////////
 
 int rocshmem_gin_create_qps(int nRanks, int myRank, int (*allgather)(void* ctx, void* buf, size_t size),
-                            void* allgather_ctx, rocshmem_gin_qp_set_t* out_qp_set, void*** out_gpu_qps) {
+                            void* allgather_ctx, int traffic_class, rocshmem_gin_qp_set_t* out_qp_set,
+                            void*** out_gpu_qps) {
   // Note: log_pe_number and device-side logd_constants must be initialized
   // by the consumer binary (e.g. rccl-tests) which links rocshmem's device
   // bitcode and can resolve HIP_SYMBOL(logd_constants).
 
   auto* set = new rocshmem_gin_qp_set();
   set->nRanks = nRanks;
   set->myRank = myRank;
+  set->traffic_class = static_cast<uint8_t>(traffic_class);
 
   // 1. Open IB device and detect provider
   if (gin_open_ib_device(set) != 0) {
@@ -1236,7 +1239,8 @@ int rocshmem_gin_create_qps(int nRanks, int myRank, int (*allgather)(void* ctx,
     *out_gpu_qps = gpu_ptr_array;
   }
 
-  LOG_TRACE("GIN QP factory: %d QPs ready on %s (rank %d/%d)", nRanks, set->nic.nic_name.c_str(), myRank, nRanks);
+  LOG_TRACE("GIN QP factory: %d QPs ready on %s (rank %d/%d, tc=%d)", nRanks, set->nic.nic_name.c_str(), myRank,
+            nRanks, set->traffic_class);
   *out_qp_set = set;
   return 0;
 

@@ -71,13 +71,16 @@ typedef struct rocshmem_gin_qp_set* rocshmem_gin_qp_set_t;
  *                          Signature: int (*)(void* ctx, void* buf, size_t size)
  *                          Performs in-place allgather of size bytes per rank.
  * @param[in]  allgather_ctx Opaque context passed to allgather.
+ * @param[in]  traffic_class IB traffic class (0-255) for QP INIT->RTR.
+ *                           Value 0 means no specific TC (default behavior).
  * @param[out] out_qp_set   Opaque handle to the created QP set.
  * @param[out] out_gpu_qps  GPU-accessible array of nRanks QueuePair pointers.
  *
  * @return 0 on success, non-zero on failure.
  */
 GIN_QP_API int rocshmem_gin_create_qps(int nRanks, int myRank, int (*allgather)(void* ctx, void* buf, size_t size),
-                                       void* allgather_ctx, rocshmem_gin_qp_set_t* out_qp_set, void*** out_gpu_qps);
+                                       void* allgather_ctx, int traffic_class, rocshmem_gin_qp_set_t* out_qp_set,
+                                       void*** out_gpu_qps);
 
 /**
  * @brief Destroy a set of GIN QPs and release all IB resources.