From 99bdd357f59484ef55df7e3ed7abe47c36d34673 Mon Sep 17 00:00:00 2001 From: Aurelien Bouteiller Date: Fri, 25 Sep 2026 15:46:03 -0400 Subject: [PATCH 1/5] Add traffic-class control to gin-gda --- .../src/gin_rocshmem_gda_factory_test.cu | 8 ++++---- .../rccl/src/gin/gin_plugin_rocshmem_gda.cc | 18 ++++++++++++++++-- .../rccl/src/gin/gin_rocshmem_gda_factory.cc | 9 ++++++--- .../src/include/gin/gin_rocshmem_gda_factory.h | 5 ++++- 4 files changed, 30 insertions(+), 10 deletions(-) diff --git a/projects/rccl-tests/src/gin_rocshmem_gda_factory_test.cu b/projects/rccl-tests/src/gin_rocshmem_gda_factory_test.cu index 1b7cc7531f96..662f064dd445 100644 --- a/projects/rccl-tests/src/gin_rocshmem_gda_factory_test.cu +++ b/projects/rccl-tests/src/gin_rocshmem_gda_factory_test.cu @@ -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; diff --git a/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc b/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc index c4ef74019907..b0609601d022 100644 --- a/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc +++ b/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc @@ -22,6 +22,10 @@ #include "gin/gin_rocshmem_gda_factory.h" #include +// 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,19 @@ 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 > comm config > default 0 + int trafficClass = 0; + if (ncclParamGinIbTc() != -1) { + trafficClass = static_cast(ncclParamGinIbTc()); + } else if (ncclParamIbTc() != -1) { + trafficClass = static_cast(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; diff --git a/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc b/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc index 6c0fe87c2a7b..9d0ed00fe363 100644 --- a/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc +++ b/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc @@ -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,7 +1084,8 @@ 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). @@ -1091,6 +1093,7 @@ int rocshmem_gin_create_qps(int nRanks, int myRank, int (*allgather)(void* ctx, auto* set = new rocshmem_gin_qp_set(); set->nRanks = nRanks; set->myRank = myRank; + set->traffic_class = static_cast(traffic_class); // 1. Open IB device and detect provider if (gin_open_ib_device(set) != 0) { @@ -1236,7 +1239,7 @@ 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, traffic_class); *out_qp_set = set; return 0; diff --git a/projects/rccl/src/include/gin/gin_rocshmem_gda_factory.h b/projects/rccl/src/include/gin/gin_rocshmem_gda_factory.h index fa16628d44d8..c59bf0a0eb60 100644 --- a/projects/rccl/src/include/gin/gin_rocshmem_gda_factory.h +++ b/projects/rccl/src/include/gin/gin_rocshmem_gda_factory.h @@ -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. From 3c400a0e581f5ed7ca7bcb715827b129904c4563 Mon Sep 17 00:00:00 2001 From: Aurelien Bouteiller Date: Tue, 13 Oct 2026 16:11:26 -0400 Subject: [PATCH 2/5] Add a log output for the traffic class --- projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc b/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc index b0609601d022..8c406eda0c99 100644 --- a/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc +++ b/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc @@ -146,7 +146,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; } From 7ddac87b9297b606acf9ca2051533babd0b59996 Mon Sep 17 00:00:00 2001 From: Aurelien Bouteiller Date: Wed, 14 Oct 2026 18:00:14 -0400 Subject: [PATCH 3/5] Address PR review: document GIN-GDA traffic class and fix trace log Update the GIN how-to and env reference to reflect that the rocSHMEM GDA backend now honors NCCL_GIN_IB_TC / NCCL_IB_TC / host-comm traffic class, and that ROCSHMEM_GDA_TRAFFIC_CLASS no longer drives these QPs. Note why the device-communicator ginTrafficClass rung is not consulted on this path. Log the narrowed uint8_t traffic class actually programmed into the QP instead of the pre-truncation int. --- projects/rccl/docs/how-to/device-api-gin.rst | 11 +++++++---- projects/rccl/docs/userguide/source/env.rst | 5 +++-- projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc | 5 ++++- projects/rccl/src/gin/gin_rocshmem_gda_factory.cc | 3 ++- 4 files changed, 16 insertions(+), 8 deletions(-) diff --git a/projects/rccl/docs/how-to/device-api-gin.rst b/projects/rccl/docs/how-to/device-api-gin.rst index 0f905f548196..0b2ddb4c687a 100644 --- a/projects/rccl/docs/how-to/device-api-gin.rst +++ b/projects/rccl/docs/how-to/device-api-gin.rst @@ -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. diff --git a/projects/rccl/docs/userguide/source/env.rst b/projects/rccl/docs/userguide/source/env.rst index b9b6846e4f6c..45263d573699 100644 --- a/projects/rccl/docs/userguide/source/env.rst +++ b/projects/rccl/docs/userguide/source/env.rst @@ -443,8 +443,9 @@ 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 ^^^^^^^^^^^^^^^ diff --git a/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc b/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc index 8c406eda0c99..b1a86cd6d906 100644 --- a/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc +++ b/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc @@ -124,7 +124,10 @@ 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 > comm config > default 0 + // 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(ncclParamGinIbTc()); diff --git a/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc b/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc index b2a21da1395c..e60d5082cbd4 100644 --- a/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc +++ b/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc @@ -1239,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, tc=%d)", nRanks, set->nic.nic_name.c_str(), myRank, nRanks, traffic_class); + 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; From ace5e6f0c358f711ede36060bf57d68ae318dd35 Mon Sep 17 00:00:00 2001 From: Aurelien Bouteiller Date: Wed, 14 Oct 2026 18:17:49 -0400 Subject: [PATCH 4/5] Add CHANGELOG entry for GIN-GDA traffic class control --- projects/rccl/CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 3e00bd52c92a..05014c876f50 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -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. From 4c3e7fc980f89aa0afa6726f1930fbd16da5b444 Mon Sep 17 00:00:00 2001 From: Aurelien Bouteiller Date: Wed, 14 Oct 2026 18:19:28 -0400 Subject: [PATCH 5/5] Update NCCL_GIN_IB_TC env entry prose for GDA backend --- projects/rccl/docs/userguide/source/env.rst | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/rccl/docs/userguide/source/env.rst b/projects/rccl/docs/userguide/source/env.rst index 45263d573699..1ddec65290b9 100644 --- a/projects/rccl/docs/userguide/source/env.rst +++ b/projects/rccl/docs/userguide/source/env.rst @@ -438,7 +438,8 @@ 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.