Revert "[GRAPH] Increase default nChannels to 112 for gfx950 (#1596)" (#1620)

* Revert "[GRAPH] Increase default nChannels to 112 for gfx950 (#1596)"

This reverts commit 1df73e209e.

* [DOC] Update Changelog

* [DOC] Update CHANGELOG

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

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();