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/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 08b90a7a5039..810e6d7bc73d 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -50,6 +50,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * `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. * `fp8_e4m3` and `fp8_e5m2` sum reductions on gfx1250 convert eight elements per instruction using the packed conversion instructions, instead of two. Results are unchanged. Other architectures and reduction operators are unaffected. +* 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. diff --git a/projects/rccl/docs/how-to/device-api-gin.rst b/projects/rccl/docs/how-to/device-api-gin.rst index ac4d3e7f83d5..f7e0e312eb37 100644 --- a/projects/rccl/docs/how-to/device-api-gin.rst +++ b/projects/rccl/docs/how-to/device-api-gin.rst @@ -258,10 +258,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..1ddec65290b9 100644 --- a/projects/rccl/docs/userguide/source/env.rst +++ b/projects/rccl/docs/userguide/source/env.rst @@ -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 ^^^^^^^^^^^^^^^ diff --git a/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc b/projects/rccl/src/gin/gin_plugin_rocshmem_gda.cc index c4ef74019907..b1a86cd6d906 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,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(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; @@ -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; } diff --git a/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc b/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc index eda33dd789f4..4e32b2ebeaab 100644 --- a/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc +++ b/projects/rccl/src/gin/gin_rocshmem_gda_factory.cc @@ -158,6 +158,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; }; /////////////////////////////////////////////////////////////////////////////// @@ -868,7 +869,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; @@ -1084,7 +1085,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). @@ -1092,6 +1094,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) { @@ -1237,7 +1240,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; 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.