GPU基础与集群架构：从计算核心到互联拓扑

GPU工作原理：大规模并行计算的硬件基础

GPU（图形处理器）最初为图形渲染设计，其架构天然适合大量简单计算单元的并行执行。与CPU追求单线程低延迟不同，GPU将更多晶体管用于计算单元（ALU）而非控制逻辑与缓存，从而在相同芯片面积下获得远超CPU的浮点运算能力。现代数据中心GPU（如NVIDIA A100/H100、AMD MI300）本质上是高度并行的流式多处理器阵列。

理解GPU工作原理需抓住三个关键层级：

  • 流式多处理器（SM）：GPU的基本调度单元，每个SM包含多个CUDA核心、共享内存、寄存器文件和调度器。A100每个SM含64个FP32 CUDA核心，共108个SM。
  • 线程层次结构：GPU以线程束（Warp）为最小调度单位（通常32线程），同一Warp内的线程必须执行相同指令（SIMT模型）。若分支发散，各路径需串行执行，造成效率损失。
  • 内存层次：寄存器（每线程私有）→ 共享内存（每个SM内，低延迟）→ 全局显存（HBM，高带宽但延迟高）→ 主机内存（PCIe/NVLink访问）。数据局部性直接决定访存效率。
关键指标解读： 峰值算力（TFLOPS）受限于SM频率与核心数；显存带宽（GB/s）决定数据搬运上限；Occupancy（占用率）反映SM上活跃Warp占比，过低会导致延迟无法隐藏。调度器需关注的是计算密度（算术强度）与访存密度的匹配，而非单纯峰值。

CUDA编程模型：从线程到任务映射

CUDA（Compute Unified Device Architecture）是NVIDIA提供的并行计算平台，其编程模型将GPU视为协处理器：Host（CPU）负责串行逻辑与调度，Device（GPU）执行大规模并行内核函数（Kernel）。程序员通过定义线程层次来组织并行度，而硬件负责将线程映射到SM上执行。

CUDA的线程组织为三层结构，与硬件调度直接对应：

  1. Grid（网格）：一个Kernel启动对应的全部线程集合，由多个Block组成。
  2. Block（线程块）：可被调度到单个SM上的线程集合，Block内线程可通过共享内存通信并同步。
  3. Thread（线程）：最小执行单元，每个线程有独立寄存器与局部变量。

一个典型的CUDA Kernel启动与执行流程如下：

// 核函数：向量加法
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}

// Host端调用
int main() {
    float *d_A, *d_B, *d_C;
    // 1. 分配显存并拷贝数据
    cudaMalloc(&d_A, N*sizeof(float));
    cudaMemcpy(d_A, h_A, N*sizeof(float), cudaMemcpyHostToDevice);
    // 2. 配置网格与块维度
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    // 3. 启动Kernel（异步）
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
    // 4. 同步并拷贝结果回主机
    cudaMemcpy(h_C, d_C, N*sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    return 0;
}

调度器在GPU集群中分配任务时，需理解Block数量与SM数量的关系：若Block总数远大于SM数，GPU通过时间片轮转提高利用率；若Block过少，则SM空闲。同时，Block内线程数建议为32的倍数（如128/256/512），避免Warp内线程浪费。

GPU集群硬件拓扑：NVLink、PCIe与InfiniBand

单台GPU服务器的算力有限，大规模训练需将多台服务器组成集群。GPU间通信效率是集群性能的瓶颈之一，理解硬件拓扑是调度器做亲和性决策的前提。下图展示了现代8卡GPU服务器内部及跨节点的典型互联结构：

GPU集群硬件拓扑（单机8卡 + 跨节点） 服务器 Node A（8×GPU） GPU 0 SM:108 GPU 1 SM:108 GPU 2 SM:108 GPU 3 SM:108 GPU 4 SM:108 GPU 5 SM:108 GPU 6 SM:108 GPU 7 SM:108 NVLink NVLink NVLink NVLink CPU（Host） PCIe Switch Gen4 x16 PCIe HCA 网卡0 IB 200Gbps HCA 网卡1 IB 200Gbps InfiniBand InfiniBand 服务器 Node B（8×GPU） GPU 0 GPU 1 GPU 2 GPU 3 GPU 4 GPU 5 GPU 6 GPU 7 IB 交换机（集群级） 通信带宽对比 NVLink 4.0：900 GB/s（双向） PCIe Gen5 x16：~64 GB/s（双向） InfiniBand NDR：400 Gb/s = 50 GB/s NVLink延迟 ~1.3μs，PCIe ~2.5μs，IB ~5μs 调度亲和性：优先同机NVLink，其次同交换机IB

上图揭示了调度器必须感知的三层通信拓扑：

  • NVLink（片间互联）：GPU与GPU之间通过NVLink直连，形成高带宽、低延迟的GPU到GPU通信域。A100支持NVLink 3.0（600GB/s），H100支持NVLink 4.0（900GB/s）。NVLink还支持NVSwitch实现全互联，使任意GPU对之间带宽一致。
  • PCIe（板级互联）：GPU与CPU、网卡之间通过PCIe总线通信。PCIe带宽远低于NVLink，且多GPU共享PCIe Switch时存在带宽争抢。GPU Direct DMA允许GPU绕过CPU直接访问网卡或存储，减少数据拷贝。
  • InfiniBand（节点间互联）：跨服务器通信依赖InfiniBand（或RoCEv2）。IB采用RDMA（远程直接内存访问）技术，无需CPU参与即可在GPU显存间直接传输数据（GPUDirect RDMA），延迟约5μs，远低于传统TCP/IP。

调度亲和性设计：拓扑感知的分配策略

集群调度器（如Slurm、Kubernetes + GPU插件）在分配GPU时，若忽略拓扑，可能将同一训练任务的多个GPU分散在不同服务器上，导致通信走IB而非NVLink，训练性能下降30%~50%。亲和性调度的核心目标是将通信频繁的GPU尽量放置在物理邻近位置。

一个典型的拓扑感知调度流程如下：

  1. 拓扑发现：通过NVIDIA DCGM或nvidia-smi读取GPU的NVLink连接矩阵、PCIe Switch归属、NUMA节点等信息，构建集群拓扑图。
  2. 通信需求分析：解析作业配置（如分布式训练框架的rank数量、all-reduce通信模式），估算GPU间通信流量矩阵。
  3. 放置决策：采用图分区算法或贪心策略，将通信量大的GPU对分配到同一NVLink域；将不同作业的GPU尽量隔离到不同NUMA节点，避免PCIe带宽争抢。
  4. 动态调整：运行时监控实际通信带宽与延迟，若发现拥塞则触发迁移或重新调度。
实践建议： 使用Slurm时，通过--gpus-per-node=8配合--ntasks-per-node=8可确保任务占满单节点；若需多节点，使用--gres=ib:2预留IB网卡。Kubernetes中可结合NodeLabeller为节点打上GPU拓扑标签，再用自定义调度器实现拓扑感知分配。

误区与复盘：拓扑感知的常见陷阱

误区 后果 正确做法
只按GPU空闲数量分配，忽略NVLink域 跨节点通信，all-reduce耗时增加2~4倍 优先分配同一NVLink域内的GPU
未考虑PCIe Switch共享带宽 多个GPU同时访问网卡时带宽争抢 将通信密集作业的GPU分散到不同PCIe Switch
忽略NUMA距离 GPU访问远端CPU内存，延迟增加 绑定GPU与同NUMA节点的CPU核心
未启用GPUDirect RDMA 数据经CPU中转，延迟高且CPU占用大 配置IB网卡与GPU的GPUDirect支持

复盘案例：某训练任务申请16块GPU，调度器随机分配至2台服务器各8块。实测发现，节点内NVLink通信带宽约600GB/s，而跨节点IB仅50GB/s，导致数据并行中的梯度同步（all-reduce）耗时占总训练时间的40%。优化后，调度器将任务改为单节点16卡（若节点支持）或采用层次化all-reduce（节点内NVLink聚合后再跨节点），训练吞吐提升1.8倍。此案例说明：调度器必须将硬件拓扑作为一等公民纳入决策，而非仅看GPU数量。

← 上一章 📖 返回目录 下一章 →