【NCCL】9 PAT AllGather 的 nMaxChannels 和 nWarps

comm->nChannels 的初始化

根据 NCCL v2.28.9 代码分析，对于 NCCL_ALGO_PAT 算法，在 AllGather 通信行为，多机通信场景（每个节点1个GPU，即 comm->nNodes == comm->nRanks）下，comm->nChannels 的初始化过程如下：

comm->nChannels 初始化流程

1. 初始值设定（src/init.cc:1090）

comm->nChannels = std::min(treeGraph->nChannels, ringGraph->nChannels);

初始值 = min(treeGraph->nChannels, ringGraph->nChannels)

• ringGraph->nChannels: Ring 算法的通道数（范围 1 到 MAXCHANNELS/2，即 1-16）
• treeGraph->nChannels: Tree 算法的通道数（等于 ringGraph->nChannels）

因此，初始值通常为 ringGraph->nChannels（1-16之间，取决于拓扑）。

2. 通道复制翻倍（src/graph/connect.cc:452）

nChannels = comm->nChannels = std::min(MAXCHANNELS, nChannels*2);

通道数翻倍（但不超过 MAXCHANNELS=32）。

3. 可能的额外增加

根据特定条件，通道数可能进一步增加：

a) CollNet 场景（src/graph/connect.cc:459-460）

如果启用 CollNet 且节点内带宽 > 节点间带宽：

int collNetNchannels = std::min(MAXCHANNELS, nChannels+nChannels/2);

b) 高性能 GPU（src/graph/connect.cc:473-474）

对于计算能力 ≥ 9.0 的 GPU（Hopper 架构），如果多节点且节点内带宽 > 45 GB/s：

nChannels = comm->nChannels = copyChannels(comm, nChannels, 2*nChannels, ...);

再次翻倍。

c) Unpack 网络设备（src/graph/connect.cc:479-480）

如果使用 Unpack 网络设备且多节点：

nChannels = comm->nChannels = copyChannels(comm, nChannels, 2*nChannels, ...);

4. 最终限制（src/graph/connect.cc:487-491）

nChannels = comm->nChannels = std::min(std::min(ncclMaxNchannels(), nChannels), 
                                       comm->config.maxCTAs);
nChannels = comm->nChannels = copyChannels(comm, nChannels, 
                                           std::max(ncclMinNchannels(), comm->config.minCTAs), ...);

应用用户设置的最小/最大通道数限制（NCCL_MIN_NCHANNELS 和 NCCL_MAX_NCHANNELS 环境变量）。

总结

对于 PAT 算法在多机 AllGather 场景（每节点1个GPU）：

1. 基础值：取决于拓扑计算的 ringGraph->nChannels（通常 1-16）
2. 第一次翻倍：基础值 × 2（不超过32）
3. 可能的额外翻倍：
   • 如果是 Hopper GPU（计算能力≥9.0）且带宽>45GB/s：再 × 2
   • 如果使用 Unpack 网络：再 × 2
4. 最终范围：受 NCCL_MIN_NCHANNELS 和 NCCL_MAX_NCHANNELS 环境变量限制

典型值：

• 普通情况：2-32 个通道
• 高性能场景（Hopper GPU + 高带宽）：可能达到 32 个通道（MAXCHANNELS 上限）

重要提示：虽然 PAT 算法会使用这些通道，但在实际执行时（topoGetAlgoInfo），通道数可能会根据数据大小动态减少：

while (nBytes < nc * nt * threadThreshold) {
  if (nc >= 2) nc--;
  else break;
}

因此，实际使用的通道数 = 1 到 comm->nChannels 之间，取决于传输数据大小。

nc、nt、threadhold对最终nMaxChannels和nWraps的影响

在 src/enqueue.cc 的 topoGetAlgoInfo 函数中，这三个变量的取值如下：

1. nc = comm->nChannels

初始值：comm->nChannels

这个值取决于拓扑结构，通常在初始化时根据系统配置确定。典型值范围是 8-32 个 channel。

动态调整：

• 在第 1886-1890 行，会根据数据量进行动态调整：

// Ring/Tree channel tuning
while (nBytes < nc * nt * threadThreshold) {
  if (nc >= 2) nc--;
  else break;
}

2. nt = comm->maxThreads[info->algorithm][info->protocol]

初始值：comm->maxThreads[NCCL_ALGO_PAT][protocol]

PAT 算法的 maxThreads 没有在 src/graph/tuning.cc 中显式设置，因此会使用默认初始化值（通常为 0 或未定义）。

最终值：在第 1907 行 被强制设置为：

if (info->algorithm == NCCL_ALGO_PAT) nt = NCCL_MAX_NTHREADS;

其中 NCCL_MAX_NTHREADS 在 src/include/device.h 中定义为 640。

3. threadThreshold = comm->threadThresholds[info->algorithm][info->protocol]

取值：comm->threadThresholds[NCCL_ALGO_PAT][protocol]

在 src/graph/tuning.cc 的第 542-546 行，所有算法的 threadThresholds 被统一初始化：

for (int a=0; a<NCCL_NUM_ALGORITHMS; a++) {
  comm->threadThresholds[a][NCCL_PROTO_LL] = NCCL_LL_THREAD_THRESHOLD;
  comm->threadThresholds[a][NCCL_PROTO_LL128] = NCCL_LL128_THREAD_THRESHOLD;
  comm->threadThresholds[a][NCCL_PROTO_SIMPLE] = NCCL_SIMPLE_THREAD_THRESHOLD;
}

根据 protocol 的不同：

• NCCL_PROTO_LL: NCCL_LL_THREAD_THRESHOLD
• NCCL_PROTO_LL128: NCCL_LL128_THREAD_THRESHOLD
• NCCL_PROTO_SIMPLE: NCCL_SIMPLE_THREAD_THRESHOLD

这些常量的具体值需要在头文件中查找，但通常 NCCL_SIMPLE_THREAD_THRESHOLD 约为 64 字节。

总结

对于 PAT 算法 + AllGather + 多机场景：

1. nc（channel 数量）：初始为 comm->nChannels（8-32），会根据数据量动态减少
2. nt（线程数）：最终固定为 640（NCCL_MAX_NTHREADS）
3. threadThreshold：根据 protocol 取对应的阈值常量（如 NCCL_SIMPLE_THREAD_THRESHOLD）

这些值共同决定了 PAT 算法在多机 AllGather 场景下的 channel 选择策略。

nc 和nt 计算完后，给info内的成员变量赋值。

  info->nMaxChannels = nc;
  info->nWarps = nt/WARP_SIZE;