From b17338d164fc8c6bcf8f559b3a6e1ee12480ebea Mon Sep 17 00:00:00 2001 From: Nilesh M Negi Date: Fri, 28 Mar 2025 17:57:06 -0500 Subject: [PATCH] Revert "[GRAPH] Increase default nChannels to 112 for gfx950 (#1596)" (#1620) * Revert "[GRAPH] Increase default nChannels to 112 for gfx950 (#1596)" This reverts commit 1df73e209e536397b53e0a1c57560119dcd2abe1. * [DOC] Update Changelog * [DOC] Update CHANGELOG --- CHANGELOG.md | 11 +++++++++++ src/init.cc | 5 +---- 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 876c61948a..abb5f75832 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -2,6 +2,17 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https://rccl.readthedocs.io) +## Unreleased - RCCL 2.22.3 for ROCm 6.5.0 + +### Added + +* Added new GPU target `gfx950` +* Added support for `unroll=1` in device-code generation to improve performance + +### Known issue + +* Using more than 64 channels can cause a segmentation fault when multiple different collectives are used in the same `ncclGroup()` call + ## Unreleased - RCCL 2.22.3 for ROCm 6.4.0 ### Added diff --git a/src/init.cc b/src/init.cc index 12eabd4364..45259a01db 100644 --- a/src/init.cc +++ b/src/init.cc @@ -1396,7 +1396,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p allGather3Data[rank].nc = std::max(allGather3Data[rank].nc, 4/ringGraph->nChannels); if (ringGraph->nChannels > MAXCHANNELS/2) allGather3Data[rank].nc = 1; - if (IsArchMatch(comm->topo->nodes[GPU].nodes[idx].gpu.gcn, "gfx942")) { + if (IsArchMatch(comm->topo->nodes[GPU].nodes[idx].gpu.gcn, "gfx942") || IsArchMatch(comm->topo->nodes[GPU].nodes[idx].gpu.gcn, "gfx950")) { // Multi-node MI300A int managed = 0; CUDACHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); @@ -1413,9 +1413,6 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p allGather3Data[rank].nc = 4; } } - if (IsArchMatch(comm->topo->nodes[GPU].nodes[idx].gpu.gcn, "gfx950")) { - allGather3Data[rank].nc = 4; - } allGather3Data[rank].pivotA2AEnabled = comm->topo->pivotA2AEnabled && rcclParamPivotAlltoallEnable(); comm->topo->ll128Enabled = comm->topo->ll128Enabled || rcclParamLL128ForceEnable();