GPU 硬件与资源共享

一句话结论

GPU 硬件要从三类资源理解：SM/Tensor Core 决定计算吞吐，HBM 决定显存容量和带宽，NVLink/NVSwitch/PCIe 决定数据在 GPU、CPU 和网卡之间怎么流动。面试里不要只背型号参数，要能把“算力、显存、互联”对应到训练和推理瓶颈。

核心概念

主流 GPU 对比

关键机制

GPU 性能通常由下面几类资源共同决定：

1. 计算资源：SM 数量、Tensor Core 能力、时钟频率、支持的数据类型。
2. 片上资源：寄存器、shared memory、L1/L2 cache，决定 occupancy 和数据复用能力。
3. 显存资源：HBM 容量决定能放下多大模型和 batch，HBM 带宽决定 memory-bound 算子的上限。
4. 互联资源：PCIe/NVLink/NVSwitch/RDMA 决定多卡和跨机通信效率。
5. 功耗散热：训练集群里功耗墙、降频和散热也会影响实际吞吐。

常见误区

关联模块

• CPU vs GPU：先理解 GPU 为什么为吞吐而设计。
• CUDA 执行模型：把 SM、warp、block 和 kernel 执行关系串起来。
• 性能指标：用 TFLOPS、HBM 带宽、Roofline 判断硬件上限。
• GPU 互联与数据路径：继续看 PCIe、NVLink、NVSwitch、RDMA 的真实数据流。

一句话结论

在 8 卡 H100/A100 训练服务器里，单机内 GPU-GPU 通信优先走 NVLink/NVSwitch，跨机优先走 GPUDirect RDMA，PCIe 则承担 CPU、GPU、NIC、NVMe 之间的通用 I/O。理解这三者不是为了背硬件名词，而是为了回答一个 AI Infra 的核心问题：同样是「8 张 GPU」，为什么放置位置不同，训练吞吐和延迟会差很多？

系统链路

三者定位

简单类比：

• NVLink：GPU 之间的“高速内环路”。
• PCIe：服务器内部设备之间的“通用高速公路”。
• RDMA：服务器之间的“跨城专线”。

带宽差异

NVLink：单机 GPU-GPU 通信的最高优先级

对 H100 来说，NVIDIA 给出的 H100 SXM 互联规格是 NVLink 900 GB/s，H100 NVL 的 PCIe 形态则列出 NVLink 600 GB/s 和 PCIe Gen5 128 GB/s。

对 A100 来说，NVIDIA 给出的 A100 80GB SXM 互联规格是 NVLink 600 GB/s，PCIe Gen4 为 64 GB/s，并且 A100 PCIe 版本通过 NVLink Bridge 主要支持最多两张 GPU 的桥接。

这里要注意：

• NVIDIA 规格表里的 600 GB/s、900 GB/s 通常是 GPU 级别的 NVLink 聚合带宽。
• 这是面向 GPU-GPU 的专用互联，不是 CPU-GPU 通用总线。
• 在 8 卡 SXM 服务器里，NVLink 往往不是简单点对点线缆，而是通过 NVSwitch 形成交换网络，让多张 GPU 之间可以高带宽互联。

PCIe：通用但相对慢

PCIe 的带宽按 代际 × lane 数计算。PCI-SIG 说明 PCIe 4.0 每 lane 每方向支持 16.0 GT/s，PCIe 5.0 每 lane 每方向支持 32.0 GT/s。

常见 GPU 是 x16 链路，可以近似理解为：

所以 PCIe 和 NVLink 的差距很明显：

也就是说，NVLink 的 GPU-GPU 带宽通常是 PCIe 的数倍到十几倍。这也是为什么张量并行、MoE All-to-All、频繁 collective 更偏好单机 NVLink/NVSwitch。

RDMA / InfiniBand：跨机通信的主力

跨服务器时，数据不能直接走 NVLink，通常走 InfiniBand 或 RoCE RDMA。NVIDIA Quantum-2 InfiniBand 平台支持最高 400 Gb/s，ConnectX-7 InfiniBand 适配器也支持单端口或双端口 400 Gb/s。

换算一下：

400 Gb/s ÷ 8 ≈ 50 GB/s

所以一张 400G 网卡的理论线速大约是 50 GB/s。如果一台服务器有多张 400G NIC，例如 8 张 GPU 搭配 4 到 8 张 NIC，就可以通过 multi-rail 并行提升跨机总带宽。

但即便如此，跨机 RDMA 的单链路带宽通常仍低于单机内 NVLink/NVSwitch：

• 单 GPU NVLink 聚合：数百 GB/s 级别。
• 单 400G RDMA NIC：约 50 GB/s 线速级别。
• 多 NIC 可以并行，但会受 PCIe、NUMA、交换机、拥塞控制、通信库调度影响。

延迟差异

延迟比带宽更依赖平台、拓扑、驱动、通信库和消息大小，所以很难给一个绝对固定值。工程上可以按这个顺序理解：

NVLink/NVSwitch 延迟最低
  < PCIe P2P
  < 跨 Socket PCIe / Host staging
  < RDMA 跨机

NVLink / NVSwitch 延迟

NVLink 面向 GPU-GPU 通信，路径短、协议栈轻、带宽高。单机内 GPU 之间做 AllReduce、AllGather、ReduceScatter 时，如果能走 NVLink/NVSwitch，通常是最优路径。

在 8 卡 SXM 服务器里，GPU 间通信通常是：

CPU 不负责搬运这段数据，只负责启动 kernel、提交通信操作、管理进程和页表等控制逻辑。

PCIe 延迟

PCIe 是通用 I/O 协议，层次更多，路径可能经过：

如果两张 GPU 在同一个 PCIe switch 下，PCIe P2P 还算可接受；如果跨 Socket，则可能经过 CPU 间互联，例如 UPI、Infinity Fabric 等，延迟和抖动都会上升。

更糟糕的情况是不能直接 P2P，需要走 host staging：

这会额外消耗 CPU 内存带宽，并增加一次或多次拷贝。

RDMA 延迟

RDMA 的优势是跨机时避免远端 CPU 参与数据搬运。NVIDIA 对 GPUDirect 的描述是，网络适配器和存储设备可以直接读写 GPU memory，从而消除不必要的内存拷贝、降低 CPU 开销和延迟。

但跨机 RDMA 仍然要经过：

所以它通常比单机 NVLink/NVSwitch 延迟更高，并且容易受网络拥塞、交换机层级、路由、ECN/PFC、RoCE/IB 配置影响。

一台 8 卡 H100/A100 服务器内部，数据如何流转？

这里分两种常见形态：SXM/HGX 8 卡服务器和PCIe 8 卡服务器。

8 卡 H100/A100 SXM/HGX：主路径是 NVLink + NVSwitch

高端 8 卡 A100/H100 训练服务器通常使用 SXM 模组和 HGX/DGX 形态。以这种形态为例，单机内 GPU-GPU 通信一般不走 PCIe，而是走：

在这种服务器内：

模型参数、梯度、激活值、KV Cache 等主要在 GPU HBM。H100 SXM 的 GPU memory bandwidth 是 3.35 TB/s，H100 NVL 规格中列出的 memory bandwidth 是 3.9 TB/s。A100 80GB PCIe/SXM 的 GPU memory bandwidth 分别列为 1,935 GB/s 和 2,039 GB/s。

1. GPU 自己的计算数据在 HBM 中。

例如 GPU0 要把梯度 chunk 发给 GPU3，数据从 GPU0 HBM 读出，经 NVLink 进入 NVSwitch，再通过 NVLink 到 GPU3 HBM。CPU 不拷贝这段数据。

1. GPU-GPU 通信走 NVLink/NVSwitch。

CPU 会启动 CUDA kernel、初始化 NCCL communicator、管理进程、内存注册、同步等，但真正的数据平面不应该经过 CPU 内存。

1. CPU 主要负责控制和调度。

PCIe 仍负责 CPU-GPU 控制、Host memory 到 GPU memory 的数据加载、GPU 到 NIC 的路径、GPU 到 NVMe 或存储路径，以及没有 NVLink 可用时的 fallback。

1. PCIe 仍然存在，但不是 GPU-GPU 主通道。

8 卡 PCIe 服务器：GPU-GPU 可能走 PCIe P2P 或 Host staging

如果是 PCIe 形态的 8 卡服务器，拓扑就更复杂。典型路径可能是：

也可能是：

如果 P2P 不可用，可能退化为：

这就是为什么 8 卡 PCIe 服务器做大模型训练时，性能通常不如 8 卡 SXM/NVSwitch 服务器。PCIe 服务器不是不能训练，而是 GPU-GPU 通信密集型任务会更容易被互联瓶颈限制。

单机内 AllReduce 数据如何走？

以 8 卡数据并行训练为例，每张 GPU 都有一份模型副本，每轮 backward 后需要同步梯度。NCCL 通常会根据拓扑选择 ring、tree、CollNet、NVLS 等算法。概念上可以理解为：

GPU0 梯度 chunk A → GPU1 → GPU2 → ... → GPU7
GPU1 梯度 chunk B → GPU2 → GPU3 → ... → GPU0
...

在 SXM/NVSwitch 机器里，这些 chunk 主要沿 NVLink/NVSwitch 交换。

更具体地说：

1. 每张 GPU 把梯度切成多个 chunk。
2. reduce-scatter 阶段，每个 chunk 沿环或树传播，并在目标 GPU 上累加。
3. all-gather 阶段，累加后的结果再分发回所有 GPU。
4. 如果 NVLink/NVSwitch 可用，NCCL 会尽量让通信走 GPU-GPU 高速路径。
5. 如果某些 GPU 间只能走 PCIe，通信图会变慢，尤其是跨 Socket GPU 对。

因此在单机 8 卡服务器中，拓扑好坏直接影响 AllReduce 的有效带宽。

服务器之间，数据如何流转？

跨服务器时，NVLink 通常只管本机内 GPU；跨机要靠 NIC 和网络。理想路径是 GPUDirect RDMA。

有 GPUDirect RDMA 时

假设 Server A 的 GPU0 要发数据给 Server B 的 GPU3，理想路径是：

这个路径的关键点是：

• NIC 可以直接读写 GPU memory。
• 数据不需要先拷贝到 CPU DRAM。
• CPU 主要负责控制面，例如注册内存、提交 work request、处理中断或轮询 completion queue。
• 数据面由 GPU、PCIe、NIC、网络交换机完成。

NVIDIA 对 GPUDirect RDMA 的说明是，它为远程系统中的 NVIDIA GPU 提供直接通信，并消除系统 CPU 和经由系统内存的数据 buffer copy。

没有 GPUDirect RDMA 或拓扑不佳时

如果不能直接从 GPU memory 做 RDMA，可能退化为：

这会多出两侧 host memory staging：

• GPU → CPU 内存。
• CPU 内存 → NIC。
• 远端 NIC → CPU 内存。
• CPU 内存 → GPU。

这类路径会显著增加延迟、占用 CPU 内存带宽，并降低通信吞吐。

8 卡 H100/A100 多机训练的典型通信层级

多机多卡训练通常采用分层通信：

以两台 8 卡服务器做 AllReduce 为例，较优通信流程通常是：

8 张 GPU 先通过 NVLink/NVSwitch 做本机内 reduce-scatter，这样可以把本机内的高带宽利用起来。

1. 每台机器内部先 reduce。

每台服务器的 GPU 数据通过 GPUDirect RDMA 发到对端。如果有多张 NIC，会做 multi-rail 并行。

1. 每台机器通过 NIC 跨机交换部分结果。

跨机同步完成后，再通过本机 NVLink/NVSwitch 分发给本机所有 GPU。

1. 每台机器内部再 all-gather。

这就是为什么大规模训练通信库通常会做 topology-aware 优化：先用单机内最快的 NVLink/NVSwitch，跨机部分才使用 RDMA。

为什么 NUMA、Socket、NIC 亲和性很重要？

即使都有 H100/A100 和 400G RDMA，性能也可能差很多，原因是本机内 GPU 到 NIC 的路径不同。

情况一：GPU 和 NIC 在同一 PCIe switch / 同一 Socket 下

这是比较理想的路径：

• 路径短。
• 不跨 CPU socket。
• PCIe 带宽更可控。
• GPUDirect RDMA 效率较高。

情况二：GPU 和 NIC 跨 Socket

这会带来：

• 更高延迟。
• 更低有效带宽。
• 更强抖动。
• CPU socket 间互联被占用。
• 多任务并发时更容易产生瓶颈。

所以调度 8 卡任务时，不仅要看 GPU-GPU 是否同属 NVLink/NVSwitch 域，还要看 GPU-NIC 路径。跨机训练时，GPU 到 NIC 的亲和性几乎和 GPU 到 GPU 的亲和性一样重要。

通信路径优先级总结

单机内 GPU-GPU

跨机 GPU-GPU

综合排序

注意：这里的排序是工程上的常见优先级，不代表所有消息大小下都绝对成立。小消息更敏感于 latency，大消息更敏感于 bandwidth，实际还要看 NCCL 算法、batch size、通信 overlap、网络拥塞和拓扑绑定。

对调度和性能优化的启示

单机 8 卡任务

如果任务正好需要 8 张 GPU，最优是：

同一台 HGX/DGX 8 卡机器
  + 同一 NVSwitch fabric
  + GPU 到 NIC 拓扑均衡

不建议把一个强通信任务拆成：

4 卡在机器 A + 4 卡在机器 B

除非单机没有 8 卡资源，或者任务本身跨机并行效率很好。因为这样会把本来可以走 NVLink/NVSwitch 的通信变成 RDMA 跨机通信。

4 卡任务

如果一个任务只需要 4 张 GPU，优先级通常是：

1. 同一 NVLink/NVSwitch 域内的 4 张 GPU。
2. 同一 PCIe switch 或同一 Socket 下的 4 张 GPU。
3. 同机但跨 Socket 的 4 张 GPU。
4. 跨机器 2+2 或 1+3。

原因是跨 Socket 和跨机都会增加通信路径长度，并引入 PCIe、CPU interconnect、NIC、交换机等额外瓶颈。

多机 8×N 卡任务

多机训练时，调度器需要同时考虑：

• 每台机器内部是否是完整 8 卡 NVSwitch 拓扑。
• 每张 GPU 到 NIC 的距离。
• NIC 数量和 rail 分布。
• 多台机器是否在同一 leaf switch 或同一网络 pod。
• 是否有 RDMA 拥塞。
• NCCL topology file 或自动探测结果是否正确。

常见排查命令

排查时不要只看“有没有 8 张 GPU”，而要同时看：

• GPU-GPU 是否走 NVLink/NVSwitch。
• GPU-NIC 是否跨 Socket。
• RDMA 是否启用 GPUDirect。
• NCCL 是否真的选择了 IB/RDMA，而不是 fallback 到 TCP。
• 是否存在某个慢 rank、慢 NIC 或拥塞路径拖累整体。

一句话面试版回答

NVLink/NVSwitch 是 8 卡 H100/A100 服务器内部 GPU-GPU 通信的主路径，带宽可达数百 GB/s 到 900 GB/s 级别；PCIe 是 CPU、GPU、NIC 之间的通用 I/O 总线，A100 常见是 Gen4，H100 常见是 Gen5，带宽明显低于 NVLink；RDMA 是跨服务器通信机制，通常通过 InfiniBand 或 RoCE，把本机 GPU memory 经 NIC 和网络直接写到远端 GPU memory，避免 CPU host memory staging。单机内训练通信应优先走 NVLink/NVSwitch，跨机时应使用 GPUDirect RDMA，并尽量保证 GPU-NIC 在同一 NUMA/SOCKET/PCIe switch 附近，否则跨 Socket 会增加延迟、降低有效带宽并引入抖动。

参考资料

1. NVIDIA H100 Tensor Core GPU
2. NVIDIA A100 Tensor Core GPU
3. PCI-SIG: PCIe 5.0 bit rates
4. NVIDIA Quantum-2 InfiniBand Platform
5. NVIDIA GPUDirect

一句话结论

CUDA 执行模型的核心链路是：CPU Host 发起 kernel launch，CUDA runtime 创建 grid，grid 拆成 block，block 被调度到 SM 上，block 内 thread 被组织成 warp，SM 以 warp 为单位发射指令并用多 warp 并发隐藏延迟。

系统链路

CPU Host 代码
  |
  |  kernel<<<gridDim, blockDim>>>(args)
  v
CUDA Kernel Launch
  |
  v
Grid：一次 kernel launch 对应一个 grid
  |
  +-- Block 0  ---- threads: 0, 1, 2, ...
  +-- Block 1  ---- threads: 0, 1, 2, ...
  +-- Block 2  ---- threads: 0, 1, 2, ...
  |
  v
GPU 硬件调度
  |
  +-- SM 0 执行若干 blocks
  +-- SM 1 执行若干 blocks
  +-- SM 2 执行若干 blocks
       |
       v
     每个 block 内的 threads 被切成多个 warp
     一个 warp 通常包含 32 个 threads

__global__ void add_kernel(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];
  }
}

int threads_per_block = 256;
int blocks = (n + threads_per_block - 1) / threads_per_block;
add_kernel<<<blocks, threads_per_block>>>(a, b, c, n);

一维 grid：处理向量
blockIdx.x * blockDim.x + threadIdx.x

二维 grid：处理矩阵
row = blockIdx.y * blockDim.y + threadIdx.y
col = blockIdx.x * blockDim.x + threadIdx.x

三维 grid：处理体数据或更复杂的张量切片
x/y/z 三个方向分别映射到数据维度

int global_id = blockIdx.x * blockDim.x + threadIdx.x;
if (global_id < n) {
  // 当前 thread 处理第 global_id 个元素
}

理想情况：同一个 warp 的 threads 访问连续地址
thread 0 -> a[0]
thread 1 -> a[1]
thread 2 -> a[2]
...
thread31 -> a[31]

较差情况：同一个 warp 的 threads 访问分散地址
thread 0 -> a[0]
thread 1 -> a[1024]
thread 2 -> a[17]
...
thread31 -> a[99999]

向量加法的执行链路：

1. CPU 调用 add_kernel<<<3907, 256>>>(a, b, c, n)
2. CUDA runtime 创建一个 grid，里面有 3907 个 block
3. GPU 把 block 分配给不同 SM
4. 每个 block 的 256 个 thread 被组织成 8 个 warp
5. 每个 thread 计算自己的 global_id
6. 每个 thread 从 HBM 读取 a[i] 和 b[i]
7. CUDA Core 执行加法
8. 结果写回 c[i]
9. 所有 block 完成后，这次 kernel 结束

常见误区

关联模块

• CUDA 内存模型与 Occupancy：继续看 block size、register、shared memory 对 occupancy 的影响。
• 利用率诊断：用 SM Active、Occupancy、Warp Stall 验证执行模型是否高效。
• Stream 与异步流水线：区分 kernel 内部调度和 kernel 之间的 stream 调度。

一句话结论

CUDA 内存和 occupancy 要一起看：register/shared memory/global memory 决定数据访问成本，block size、寄存器和 shared memory 用量决定一个 SM 能挂多少 warp。Occupancy 是隐藏延迟的条件，不是最终性能目标。

系统链路

硬件视角：基本单位是 SP（Streaming Processor，也叫 CUDA Core），每个 SP 有自己的 register 和 local memory（片下内存，应对寄存器不足），只能被自己访问；多个 SP 和一块 shared memory 构成一个 SM，shared memory 被 SM 内线程共享；多个 SM 和一块全局内存构成 GPU，global memory 被所有线程访问。

软件视角对应关系：

一句话：register/local 私有，shared memory 是 block 内协作的关键，global memory 全局共享但最慢；不同 block 的线程不能直接协作。

关键机制

block size 选择：

• 范围通常是 1 到 1024。
• 取 32（warp 大小）的倍数，避免最后一个 warp 只有部分线程有效。
• 考虑 occupancy：block size 要让 SM 上能挂足够多 warp。
• 考虑 register 数量：block 太大可能占用过多寄存器，降低同时驻留的 block 数。
• 常见起点是 128、256、512，但最终要 profile。

grid size 选择：

Element-wise 程序通常让 grid size = 数据量 / block size 向上取整，并保证 block 数足够多铺满所有 SM。

int block = 256;
int grid = (n + block - 1) / block;   // 向上取整
kernel<<<grid, block>>>(...);

资源模型

两者结合才能判断 GPU 是否真用满：block 数太少会导致 GPU utilization 低，即 SM 吃不满；block 内 warp 太少或寄存器占用过高会导致 occupancy 低，即单个 SM 隐藏延迟能力差。

内存墙

内存墙指处理器速度与内存访问速度的不匹配：处理器算力增长远快于内存访问速度，导致内存成为瓶颈。在 CUDA 里，内存墙通常指 global memory 的高延迟、相对低带宽会拖慢 GPU。

缓解手段是利用 shared memory 把频繁访问的数据缓存到片上，减少对 global memory 的访问次数，例如矩阵乘 tiling 和 FlashAttention 的分块。这也是很多 GPU kernel 优化的核心思路：不是算得更快，而是少搬数据。

PyTorch 自定义 CUDA 算子

三个步骤：

1. 编写 CUDA 算子和对应的 launch 调用函数。
2. 编写 torch cpp 函数建立 PyTorch 与 CUDA 的联系，用 pybind11 封装。
3. 用 PyTorch 的 cpp 扩展库编译调用。

Tensor Core 数据路径

Tensor Core 是专门做矩阵乘加（GEMM、卷积、Attention）的硬件单元。在 Volta / Turing / Ampere 架构上，计算所需输入输出数据都放在与 CUDA Core 共享的寄存器上；在 Hopper 架构上，为了获得更好的带宽，计算所需输入数据可以直接存放在共享内存上。

常见误区

关联模块

• CUDA 执行模型：block、warp、SM 是 occupancy 的前置概念。
• 性能指标：Roofline 和 arithmetic intensity 判断是否被内存墙限制。
• 利用率诊断：用 achieved occupancy、eligible warps、stall reason 做实证判断。

一句话结论

H2D/D2H 拷贝是 CPU 内存和 GPU 显存之间的数据搬运，常见性能问题不是“有拷贝”本身，而是拷贝频繁、小而碎、触发同步、无法和计算重叠，或者多 GPU 通信绕回 CPU。优化主线是少拷贝、批量拷贝、异步拷贝、拷贝计算重叠。

诊断入口

指标解释

Host memory   = CPU 内存
Device memory = GPU 显存

H2D: Host memory   -> Device memory
D2H: Device memory -> Host memory

1. DataLoader 从磁盘读取图片/样本
2. CPU 做 decode / augmentation / tokenize
3. batch 放在 CPU 内存
4. H2D: batch 从 CPU 拷贝到 GPU
5. GPU 前向计算
6. GPU 反向传播
7. GPU 更新参数
8. 可选 D2H: loss / metric / output 拷回 CPU 打日志或后处理

# 不推荐：每步都强制 D2H + 同步
loss_value = loss.item()

# 更好：降低频率，或多个指标聚合后再拷贝
if step % 100 == 0:
    loss_value = loss.detach().item()

loader = DataLoader(
    dataset,
    batch_size=128,
    shuffle=True,
    num_workers=8,
    pin_memory=True,
    persistent_workers=True,
    prefetch_factor=2,
)

for batch, label in loader:
    batch = batch.to("cuda", non_blocking=True)
    label = label.to("cuda", non_blocking=True)
    output = model(batch)

理想流水线：

时间轴 ---->
CPU/DataLoader:  prepare batch N+1   prepare batch N+2
H2D copy:              copy N+1           copy N+2
GPU compute: compute batch N     compute batch N+1

# 不推荐：很多小 H2D
for x in list_of_tensors:
    x = x.to("cuda")

# 更好：先合并，再一次 H2D
batch = torch.stack(list_of_tensors)
batch = batch.to("cuda", non_blocking=True)

x = x.cuda()
y = x.cpu().numpy()
z = torch.tensor(y).cuda()

差：GPU0 -> CPU -> GPU1
好：GPU0 -> GPU1
好：GPU -> NIC -> GPU

重点关注：
Memcpy HtoD
Memcpy DtoH
cudaStreamSynchronize
cudaDeviceSynchronize
kernel 之间的大空洞

关联模块

• Stream 与异步流水线：用 stream/event 和多缓冲把拷贝与计算重叠。
• 利用率诊断：timeline 空洞、memcpy 和同步是 H2D/D2H 瓶颈证据。
• GPU 互联与数据路径：理解 PCIe、NVLink、GPUDirect RDMA 的路径差异。

一句话结论

CUDA stream 是 GPU 侧的异步任务队列，同一个 stream 内顺序执行，不同 stream 在资源允许时可以并行。优化目标是用 pinned memory、cudaMemcpyAsync、event 和多缓冲，把 CPU 准备、H2D/D2H 拷贝和 kernel 计算流水线化。

核心概念

本节关注“如何用 stream/event 把拷贝、kernel 和 CPU 准备组织成异步流水线”。H2D/D2H 的触发场景和 DataLoader 诊断见 Host-Device 数据拷贝。

GPU 有自己的显存（HBM），CPU 也有自己的内存（DDR）。数据在两者之间搬运，就是 H2D（Host to Device）和 D2H（Device to Host）拷贝。

因为 H2D/D2H 走的是 PCIe 或 NVLink，带宽远低于 GPU 内部 HBM 带宽。一次不注意的同步拷贝可能让 GPU 空等数百微秒，这在推理低延迟场景下尤其致命。

系统链路

同步拷贝（cudaMemcpy）：CPU 发起后阻塞等待完成，GPU 和 CPU 都在等。适合初始化、模型加载等不在乎延迟的场景。

异步拷贝（cudaMemcpyAsync）：CPU 发起后立即返回，拷贝在 GPU 端异步执行。必须配合 CUDA stream 使用，才能实现计算和拷贝重叠。

// 同步：CPU 阻塞等待
cudaMemcpy(d_ptr, h_ptr, size, cudaMemcpyHostToDevice);

// 异步：CPU 立即返回，拷贝在 stream 上执行
cudaMemcpyAsync(d_ptr, h_ptr, size, cudaMemcpyHostToDevice, stream);

关键机制

Pinned Memory 为什么更快：

普通 malloc 分配的内存是可换页的，GPU DMA 无法直接访问，需要先拷贝到临时 pinned buffer 再传输，多了一次隐式拷贝。cudaMallocHost 分配的内存被锁定在物理页上，GPU DMA 可以直接传输，省掉中间环节。但 pinned memory 不可换页，会占用实际物理内存，不能无限分配。

CUDA stream 是 GPU 上的任务队列。同一个 stream 内的操作按提交顺序串行执行；不同 stream 之间的操作可以并行执行（如果硬件资源允许）。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// stream1：拷贝输入 + 计算
cudaMemcpyAsync(d_in1, h_in1, size, cudaMemcpyHostToDevice, stream1);
kernel_a<<<grid, block, 0, stream1>>>(d_in1, d_out1);

// stream2：同时拷贝另一组输入 + 计算
cudaMemcpyAsync(d_in2, h_in2, size, cudaMemcpyHostToDevice, stream2);
kernel_b<<<grid, block, 0, stream2>>>(d_in2, d_out2);

// 等待两个 stream 都完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

资源模型

推理或训练中，经典优化是让数据拷贝和计算重叠。三缓冲（triple buffering）是最常见的模式：

Buffer A: GPU 正在计算（kernel 执行中）
Buffer B: H2D 正在拷贝下一批数据（DMA 传输中）
Buffer C: CPU 正在准备再下一批数据

下一个时间步：
  A 计算完成 → D2H 取结果
  B 拷贝完成 → 开始计算
  C 准备完成 → 开始 H2D 拷贝

这样 GPU 不用等 CPU 准备数据，CPU 也不用等 GPU 算完，三者流水线化。

实现要点：

• 每个缓冲区用独立的 CUDA stream。
• 输入数据用 pinned memory 分配。
• 用 cudaMemcpyAsync 而不是 cudaMemcpy。
• 用 cudaStreamSynchronize 或 event 等待完成。
• 缓冲区数量通常 2 或 3，多了收益递减。

常见误区

关联模块

• Host-Device 数据拷贝：定位 H2D/D2H 触发点和 DataLoader 问题。
• CUDA 执行模型：理解 stream、kernel、block/warp 的层级差异。
• 利用率诊断：在 Nsight Systems timeline 中识别 memcpy、空洞和同步。

一句话结论

Linux CFS 和 CUDA thread/block 调度都叫“调度”，但它们解决的是完全不同层次的问题。

Linux CFS 是 CPU 上的操作系统调度器，调度对象是进程或线程，目标是公平性、响应性和 CPU 时间共享。它通过 vruntime、优先级权重、抢占和上下文切换来决定哪个任务运行。

CUDA thread block 调度主要由 GPU 硬件完成，调度对象是 kernel grid 中的 block，也叫 CTA。GPU 把 block 分配到 SM 上执行；block 一旦驻留在某个 SM 上，通常会运行到完成。SM 内部再通过 warp scheduler 在多个 ready warp 之间切换，以隐藏访存和执行延迟。

CUDA Stream/Event 又是另一层。它们不是 kernel 内部 thread 的调度器，而是 CUDA 程序员用来表达任务级异步执行和依赖关系的工具。

系统链路

为什么这个问题容易混淆？

因为“线程”这个词在 CPU 和 GPU 里含义不同。

所以面试回答时要先分层：

• CPU OS 调度层：进程/线程如何共享 CPU。
• CUDA 任务提交层：stream/event 如何组织 GPU work。
• GPU kernel 执行层：block/warp 如何映射到 SM。
• 集群资源调度层：K8s/Volcano/YARN 等如何分配 GPU 设备给任务。

不要把这四层混成一个“GPU 调度器”。

Linux CFS：公平地分 CPU 时间

CFS 全称是 Completely Fair Scheduler，目标不是让某个任务跑到最快，而是在多个 runnable task 之间尽量公平地分配 CPU 时间，同时兼顾交互响应。

CFS 调度对象

CFS 调度的是 Linux 内核里的 task_struct。对用户来说，它可以表现为：

• 一个进程；
• 一个线程；
• 一个容器里的某个线程；
• 一个 cgroup 下的一组任务。

在 Linux 里，线程和进程最终都可以作为调度实体参与 CPU 调度。CFS 不关心这个任务是不是 AI 训练、数据加载、推理服务、日志线程，它只看到“这个 runnable task 需要 CPU”。

vruntime：谁“欠 CPU 时间”最多

CFS 的核心指标是 vruntime，可以理解成“加权后的虚拟运行时间”。

• 任务真实运行越久，vruntime 越大。
• nice 值越低、优先级越高，权重越大，同样运行一段真实时间，vruntime 增长越慢。
• 调度器倾向于选择 vruntime 最小的任务运行，因为它看起来“拿到的公平份额最少”。

直觉上：

谁的 vruntime 小，说明谁相对更“饿”，应该优先给 CPU。
谁的 vruntime 大，说明谁已经跑得比较多，可以先等一等。

这和传统时间片轮转不一样。Round Robin 更像“每个人轮流拿固定时间片”；CFS 更像“持续维护每个人已经拿到的公平份额”。

nice、weight 和公平份额

Linux 的 nice 值会影响任务权重。

CFS 不是简单地“高优先级永远先跑”。它通过权重影响 vruntime 增长速度，让高权重任务在长期上获得更多 CPU 时间，但仍然允许其他任务运行。

runqueue 和红黑树

每个 CPU 通常有自己的 runqueue。CFS 会把 runnable task 按 vruntime 组织起来，经典实现使用红黑树。

这个结构的意义是：调度器能快速找到“最应该运行”的任务。

抢占和上下文切换

CFS 是抢占式调度。当前任务正在 CPU 上运行时，如果出现一个更应该运行的任务，调度器可以触发抢占。

常见触发点包括：

• 周期性 tick 或调度时钟更新；
• 当前任务主动阻塞，例如等待 IO、锁、网络；
• 新任务唤醒，例如交互请求到来；
• 当前任务运行超过合理粒度；
• 更高优先级或更小 vruntime 的任务需要运行。

CPU 上下文切换通常涉及：

• 保存当前任务寄存器状态；
• 切换内核栈；
• 切换或更新地址空间相关状态；
• 更新调度统计；
• 恢复下一个任务的执行上下文。

上下文切换不是免费的。线程数过多、锁竞争严重、频繁唤醒阻塞，都可能导致 CPU 时间花在调度和切换上，而不是有效计算上。

CUDA Stream/Event：任务级异步调度

对 CUDA 程序员来说，Stream/Event 是控制 GPU work 提交顺序和依赖的工具。

Stream 是 GPU 任务队列：

• 同一个 stream 内的操作按提交顺序执行；
• 不同 stream 的操作可以并发执行，但前提是硬件资源允许；
• 常见操作包括 H2D 拷贝、kernel launch、D2H 拷贝、event record/wait。

Event 是依赖和计时工具：

• 可以记录某个 stream 上的进度点；
• 可以让另一个 stream 等待这个 event；
• 可以用于测量 GPU 端耗时；
• 可以避免 CPU 端粗暴 cudaDeviceSynchronize()。

这里要强调：stream 不等于 SM 调度器。Stream 决定的是 kernel、memcpy 等任务之间的顺序和并发机会；block 和 warp 怎么在 SM 里执行，是更底层的硬件执行机制。

CUDA Block / Warp 调度：吞吐优先

一次 kernel launch 会产生一个 grid，grid 里包含很多 block。GPU 的工作是把这些 block 分配到 SM 上执行。

block 是调度到 SM 的基本单位

block 也常被称为 CTA（Cooperative Thread Array）。一个 block 里的 thread 可以：

• 共享 shared memory；
• 使用 __syncthreads() 做 block 内同步；
• 通过 thread/block index 处理不同数据。

GPU 硬件会把 block 调度到某个 SM。一个 block 一旦驻留到 SM 上，通常不会像 CPU task 那样被 CFS 时间片频繁抢占并迁移到另一个 SM，而是运行到完成。

这带来两个重要结论：

• 不同 block 之间默认没有执行顺序保证。
• block 数量要足够多，否则 SM 可能吃不满。

block residency：为什么不是想放多少就放多少

一个 SM 能同时驻留多少个 block，受多种资源限制：

所以 CUDA 调优不是“block 越大越好”。block 太小，单个 block 并行度不足；block 太大，可能占用太多 register/shared memory，导致 SM 上可同时驻留的 warp 变少。

warp scheduler：隐藏内存延迟

block 被放到 SM 后，block 内 thread 会被组织成 warp。NVIDIA GPU 上一个 warp 通常是 32 个 thread。SM 内部的 warp scheduler 会在多个 ready warp 之间选择并发射指令。

它的关键目标不是公平，而是吞吐：

• 某个 warp 等 HBM 访存时，切换到另一个 ready warp；
• 某个 warp 遇到长延迟指令时，让其他 warp 填补流水线；
• 通过足够多 active warp 隐藏内存和执行延迟。

这就是 GPU 和 CPU 的核心差异之一：CPU 用复杂控制逻辑优化单线程延迟；GPU 用大量 warp 并发隐藏延迟，追求整体吞吐。

CFS vs CUDA 调度：核心对比

面试中最容易犯的错是说“GPU 也像 CPU 一样靠 OS 调度每个 thread”。这不对。CUDA thread 是 GPU kernel 内的逻辑执行单元，不是 Linux CFS 直接调度的 OS thread。

Stream/Event vs Thread Block/Warp：不同层次

对 CUDA 程序员来说，最需要区分的是这两层：

例子：

stream1: H2D batch0 → kernel batch0 → D2H result0
stream2: H2D batch1 → kernel batch1 → D2H result1

这是 stream 层的并发组织。每个 kernel batchX 内部又会被拆成 grid/block/thread，并由 GPU 把 block 分配到 SM、把 thread 组织成 warp。

换句话说：

• Stream/Event 解决“多个 GPU 任务之间怎么排队、等待、重叠”。
• Block/Warp 调度解决“一个 kernel 内部怎么并行执行”。

和 AI Infra 的关系

这个对比不是纯 CUDA 八股。它能解释很多系统现象。

为什么 GPU 训练任务不适合像 CPU 一样频繁抢占？

CPU 线程抢占和上下文切换相对常见，CFS 正是为共享 CPU 时间设计的。但 GPU 训练任务通常有：

• 大量显存状态：模型参数、梯度、optimizer state、activation；
• NCCL communicator 和多 rank 同步；
• kernel 执行和通信 overlap；
• CUDA context、缓存分配器、通信 buffer；
• checkpoint 恢复成本。

如果频繁像 CPU 时间片一样抢占 GPU 训练任务，可能会导致上下文切换、显存迁移、通信重建和 checkpoint 回滚成本远大于收益。所以集群调度里 GPU 抢占往往要 checkpoint-aware、gang-aware，而不是简单时间片轮转。

为什么 Time Slicing 和 MPS 不等于 CUDA block 调度？

Kubernetes / NVIDIA device plugin 里的 GPU time-slicing、MPS 属于多进程/多 Pod 共享 GPU 的资源管理机制。

• Time Slicing：多个进程或容器按时间片共享 GPU。
• MPS：多个 CUDA 进程通过 MPS server 更高效共享 GPU 执行资源。
• CUDA block/warp 调度：一个 kernel 内部 block 和 warp 如何映射到 SM。

它们层次不同。不能说“开了 time-slicing 后，K8s 会调度 CUDA thread block”。K8s 只调度 Pod；NVIDIA device plugin/driver 负责 GPU 资源共享；kernel 内部 block/warp 仍由 GPU 硬件调度。

为什么 CPU 线程很多会拖慢 GPU 任务？

GPU kernel 虽然在 GPU 上执行，但 GPU 程序仍依赖 CPU：

• CPU 负责 dataloader、预处理、tokenization；
• CPU 发起 kernel launch；
• CPU 提交 H2D/D2H；
• CPU 管理 CUDA runtime、NCCL、RPC；
• CPU 处理服务端请求、排队和调度。

如果 CPU 侧线程过多、上下文切换严重、NUMA 绑定不合理或 dataloader 卡住，GPU 可能出现空洞：SM 没活干、GPU Util 周期性下降。此时问题不是 CUDA block 调度，而是 CPU 供给链路和 OS 调度压力。

高频面试问答

一句话结论

GPU 共享不是一种机制，而是一组取舍：MIG 追求硬件隔离，MPS 追求多进程并发，time-slicing 追求低成本复用，CUDA VMM 解决虚拟地址和物理显存映射弹性。面试里要先说清隔离边界，再谈利用率提升。

核心概念

系统链路

Kubernetes 里常见链路是：GPU Operator 或 NVIDIA Device Plugin 读取共享配置，把一张物理 GPU 暴露成多个逻辑扩展资源；kube-scheduler 只按资源数量调度 Pod；真正的共享语义由 driver、MPS daemon、CUDA runtime 或 MIG manager 执行。

在 Kubernetes 里，MPS 和 Time Slicing 通常不是通过改 kube-scheduler 实现，而是通过 NVIDIA Device Plugin / NVIDIA GPU Operator 把一张物理 GPU 暴露成多个可被 Pod 申请的逻辑 GPU slot。

例如把 nvidia.com/gpu 配成 replicas: 4 后，一张物理 GPU 会向 Kubernetes 上报为 4 个逻辑可申请资源；如果开启 renameByDefault: true，Pod 侧通常申请 nvidia.com/gpu.shared: 1，从语义上区分共享 GPU 和独占 GPU。

关键边界：scheduler 只看到逻辑资源数量，不理解每个 slot 的真实性能隔离。Time Slicing 和 MPS 都不是 MIG 那种硬件切分，生产使用时要特别关注显存 OOM、P99 延迟抖动、监控归因和 workload 之间的干扰。

显存碎片

GPU 显存碎片可以分为两类：外部碎片和内部碎片。外部碎片是总空闲显存足够，但没有足够大的连续块；内部碎片是为了对齐或预留最大长度而分配了比实际需求更大的块。

面试要点：OOM 不一定表示总显存不够，也可能是可用连续块不够、allocator pool 不可回收，或某个进程持有了大量碎片化缓存。

关键机制

显存碎片治理要分层看：

如果面试官追问“调度层能不能解决显存碎片”，可以回答：调度层只能减少碎片产生概率，比如保留大块资源、限制共享密度、做显存画像；真正的细粒度碎片复用要靠 runtime allocator、PagedAttention、CUDA VMM 这类机制。

对比表

一句话：MIG 是隔离优先，MPS 是并发效率优先，Time Slicing 是部署简单和提高密度优先。强隔离多租户不要把 MPS/Time Slicing 当 MIG 用。

常见误区

关联模块

• K8S GPU 共享：继续看 device plugin、replicas 和节点级配置。
• CUDA 内存模型与 Occupancy：理解显存、shared memory、register 对共享密度的影响。
• LLM 推理系统：KV cache、PagedAttention 和 CUDA VMM 是显存共享的重要场景。
• 调度与集群：共享 GPU 会引入干扰建模、碎片治理和拓扑感知调度问题。

一句话结论

K8S 默认把 GPU 当作不可分割的扩展资源，GPU 共享要靠 NVIDIA device-plugin / GPU Operator 把物理 GPU 表达成多个可调度资源。time-slicing、MPS 和 MIG 的本质差异在隔离级别、执行方式和调度语义，不要把逻辑 slot 当成稳定的 1/N 性能。

系统链路

# time-slicing ConfigMap：把每张卡复制成 4 份
apiVersion: v1
kind: ConfigMap
metadata:
  name: time-slicing-config
  namespace: gpu-operator
data:
  any: |-
    version: v1
    flags:
      migStrategy: none
    sharing:
      timeSlicing:
        renameByDefault: false
        failRequestsGreaterThanOne: true
        resources:
        - name: nvidia.com/gpu
          replicas: 4

kubectl patch clusterpolicy/cluster-policy \
  -n gpu-operator --type merge \
  -p '{"spec":{"devicePlugin":{"config":{"name":"time-slicing-config","default":"any"}}}}'

resources:
  limits:
    nvidia.com/gpu: 1   # 实际是 1/4 张卡的时间片

# MPS ConfigMap：把每张卡按 MPS 方式分成 4 份
apiVersion: v1
kind: ConfigMap
metadata:
  name: mps-config
  namespace: gpu-operator
data:
  any: |-
    version: v1
    sharing:
      mps:
        renameByDefault: false
        resources:
        - name: nvidia.com/gpu
          replicas: 4

kubectl patch clusterpolicy/cluster-policy \
  -n gpu-operator --type merge \
  -p '{"spec":{"devicePlugin":{"config":{"name":"mps-config","default":"any"}}}}'

# ConfigMap 里放多份配置
data:
  shared: |-
    version: v1
    sharing:
      timeSlicing:
        resources:
        - name: nvidia.com/gpu
          replicas: 8
  exclusive: |-
    version: v1
    flags:
      migStrategy: none

# 给节点打 label 选择对应配置
kubectl label node gpu-node-1 \
  nvidia.com/device-plugin.config=shared
kubectl label node gpu-node-2 \
  nvidia.com/device-plugin.config=exclusive

常见误区

关联模块

• 共享方式：理解 MIG/MPS/time-slicing 的底层取舍。
• Kubernetes 核心：device plugin、extended resource、node label 和调度语义。
• 调度与集群：共享密度、显存碎片、干扰建模和租户配额。

一句话结论

GPU 性能指标要按“算力、显存、利用率、互联、能耗”五条线理解，再用 Roofline 把 kernel 归类为 compute-bound 或 memory-bound。面试中最重要的判断是：不要把理论 TFLOPS、GPU-Util 或显存占用单独当成“性能好”的证据。

核心概念

ncu --set full ./your_program
# 或使用带 Roofline 的 section
ncu --section SpeedOfLight_RooflineChart ./your_program

关联模块

• 利用率诊断：把 GPU-Util 深挖到 SM Active、Occupancy、Warp Stall。
• 瓶颈分类：把指标组合映射到 compute-bound、memory-bound、communication-bound。
• GPU 互联与数据路径：解释 PCIe、NVLink、RDMA 对多卡性能的影响。
• 性能预测指标：把这些指标进一步转成特征和标签。

一句话结论

GPU 利用率诊断不能停在 nvidia-smi 的 GPU-Util。推荐链路是：GPU-Util 看时间上有没有 kernel，SM Active 看活有没有铺满 SM，Occupancy 看 SM 里 warp 是否足够，Warp Stall 看 warp 为什么发不出去，Compute/Memory Throughput 看瓶颈类型，最后用业务吞吐和延迟判断“忙得是否有效”。

诊断入口

指标解释

1. GPU 有没有活干？
   看 GPU-Util。

2. 显存只是占着，还是 GPU 真的在计算？
   区分 memory.used 和 utilization.memory。

3. 活有没有铺满 SM？
   看 SM Active。

4. SM 里面挂了多少 warp？
   看 achieved occupancy / active warps。

5. warp 是否真的能发射？
   看 eligible warps、issued warps、stall reasons。

6. 瓶颈在 compute 还是 memory？
   看 compute throughput 和 memory throughput。

7. AI workload 是否用上 Tensor Core？
   看 tensor pipe / tensor instruction。

8. 是否被功耗、温度、时钟限制？
   看 power、clock、temperature、p-state。

9. 端到端 timeline 有没有空洞？
   用 Nsight Systems 看 CPU、CUDA、memcpy、NCCL。

10. GPU 忙得是否有业务价值？
    训练看 tokens/sec、samples/sec、step time、MFU；推理看 QPS、TTFT、TPOT、P99。

# 快速看当前 GPU 状态
nvidia-smi

# 持续刷新，观察利用率、功耗、温度等变化
nvidia-smi dmon

# 查询指定字段
nvidia-smi --query-gpu=timestamp,index,name,utilization.gpu,utilization.memory,memory.used,memory.total,power.draw,temperature.gpu --format=csv -l 1

nvidia-smi --query-gpu=memory.used,memory.total,utilization.memory --format=csv -l 1

# 用 Nsight Compute 采集完整指标
ncu --set full ./your_program

# 只采集 SM Active 相关指标
ncu --metrics sm__cycles_active.avg.pct_of_peak_sustained_elapsed ./your_program

ncu --set full ./your_program

# 常看 Nsight Compute 中的这些 section
# Launch Statistics
# Occupancy
# Scheduler Statistics

# 也可以指定 active warps 指标
ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active ./your_program

active warp：
  驻留在 SM 上，还没结束。

eligible warp：
  当前周期已经准备好，可以发射指令。

selected warp：
  当前周期被 scheduler 选中发射。

stalled warp：
  因为内存、依赖、同步、barrier 等原因暂时不能发射。

ncu --set full ./your_program

# 重点看：
# Scheduler Statistics
# Warp State Statistics
# Source Counters

ncu --set full ./your_program

# 重点看：
# GPU Speed Of Light Throughput
# Compute Workload Analysis

# 示例指标
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./your_program

ncu --set full ./your_program

# 重点看：
# Memory Workload Analysis
# GPU Speed Of Light Throughput
# L1/TEX、L2、DRAM 相关指标

ncu --set full ./your_program

# 重点看：
# GPU Speed Of Light Throughput
# Compute Workload Analysis
# Instruction Statistics

# 示例 tensor 相关指标
ncu --metrics sm__inst_executed_pipe_tensor.sum,sm__pipe_tensor_cycles_active.avg.pct_of_peak_sustained_elapsed ./your_program

import torch

torch.backends.cuda.matmul.allow_tf32 = True
torch.backends.cudnn.allow_tf32 = True

nvidia-smi -q -d POWER,CLOCK,TEMPERATURE,PERFORMANCE

nvidia-smi --query-gpu=power.draw,power.limit,clocks.sm,clocks.mem,temperature.gpu,pstate --format=csv -l 1

nvidia-smi
nvidia-smi dmon -s pucvmet
nvidia-smi --query-gpu=timestamp,index,utilization.gpu,utilization.memory,memory.used,memory.total,power.draw,clocks.sm,temperature.gpu --format=csv -l 1

DCGM_FI_DEV_GPU_UTIL
DCGM_FI_DEV_FB_USED
DCGM_FI_DEV_POWER_USAGE
DCGM_FI_DEV_GPU_TEMP
DCGM_FI_DEV_SM_CLOCK
DCGM_FI_DEV_MEM_COPY_UTIL

nsys profile -t cuda,nvtx,osrt,cudnn,cublas -o report ./your_program

ncu --set full ./your_program
ncu --kernel-name regex:your_kernel_name --set full ./your_program

import torch
import torch.profiler as profiler

with profiler.profile(
    activities=[
        profiler.ProfilerActivity.CPU,
        profiler.ProfilerActivity.CUDA,
    ],
    record_shapes=True,
    profile_memory=True,
    with_stack=True,
) as prof:
    for step, batch in enumerate(loader):
        output = model(batch)
        loss = output.sum()
        loss.backward()
        if step > 10:
            break

print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=20))

nvidia-smi

nsys profile -t cuda,nvtx,osrt,cudnn,cublas -o report ./your_program

ncu --set full ./your_program

GPU-Util 低
SM Active 低
Nsight Systems 看到 GPU timeline 有空洞

GPU-Util 高
SM Active 低
Achieved Occupancy 低

GPU-Util 高
SM Active 高
Occupancy 高
Compute Throughput 低
Memory Throughput 也低
Warp Stall 高

SM Active 高
Memory Throughput 高
Compute Throughput 低
Long Scoreboard stall 高

SM Active 高
Compute Throughput 高
Memory Throughput 不高

GPU-Util 高
SM Active 高
Tensor Pipe / Tensor Core Util 低

排查路径

GPU-Util：
  时间上有没有 kernel。

SM Active：
  kernel 有没有铺满 SM。

Occupancy：
  SM 里 active warp 是否足够。

Eligible / Issued Warps：
  warp 是否真的能发射。

Warp Stall：
  不能发射是在等什么。

Compute Throughput：
  计算管线是否接近峰值。

Memory Throughput：
  是否被显存带宽或访问模式卡住。

Tensor Core Utilization：
  AI 矩阵算力是否真的用起来。

Power / Clock / Thermal：
  是否被功耗、温度、时钟限制。

业务指标：
  GPU 忙，是否真的换来了 tokens/sec、samples/sec、QPS 和低延迟。

关联模块

• 性能指标：理解 GPU-Util、SM Active、Occupancy、Roofline 的定义。
• 瓶颈分类：把诊断指标组合翻译成 compute/memory/communication 瓶颈。
• Host-Device 数据拷贝：timeline 空洞和 memcpy 往往来自 H2D/D2H。
• CUDA 内存模型与 Occupancy：解释 occupancy 受 block size、寄存器和 shared memory 影响。

一句话结论

GPU 慢不能只说“利用率低”或“显存满”，要先分类：计算瓶颈看 Tensor Core/CUDA Core 是否打满，显存瓶颈看 HBM 带宽和访存 stall，通信瓶颈看 NCCL/memcpy/拓扑等待。分类决定优化方向，方向错了会越调越慢。

诊断入口

指标解释

排查路径

计算瓶颈（Compute-bound）

定义：GPU 的计算单元（CUDA Core / Tensor Core）是瓶颈，计算请求已满，但数据供给充足。

典型场景：大矩阵乘法（GEMM）、大 batch 的卷积、大 batch 的 prefill 阶段 Attention。

优化方向：

• 使用 Tensor Core，确保 dtype、shape 和 layout 对齐。
• 降低精度，例如 FP16/BF16/FP8/INT8。
• 优化 tile 大小和指令 mix。
• 对 decode 类任务考虑 speculative decoding，用小模型生成候选，大模型并行验证。

显存带宽瓶颈（Memory-bandwidth-bound）

定义：GPU 的显存带宽是瓶颈，计算单元在等数据，大量时间花在 HBM 读写上。

典型场景：LLM decode 阶段、elementwise 算子、LayerNorm、Softmax、Embedding lookup、Gather/Scatter、小 batch 推理。

为什么 LLM decode 常是显存瓶颈：每生成一个 token，需要从 HBM 加载大量模型权重和 KV cache，但新增 token 的计算量相对少，算术强度低，容易落在 Roofline 的 memory-bound 区域。

优化方向：

• 提高算术强度：增大 batch size、融合算子。
• 减少 HBM 读写：FlashAttention、shared memory 数据复用、算子融合。
• 压缩数据量：KV cache 量化、权重量化。
• 改善访问模式：memory coalescing、提高 L2 cache hit rate。
• 用 PagedAttention 按需分配 KV cache 物理页，减少显存浪费。

显存容量瓶颈（Memory-capacity-bound）

容量瓶颈是指显存不够放，放不下模型权重、KV cache、activation、optimizer state 或足够的 batch 数据。它和带宽瓶颈不同：容量问题首先表现为 OOM 或并发受限，带宽问题表现为算子在等数据。

典型表现：

• OOM。
• 被迫降低 batch size、sequence length 或并发请求数。
• KV cache 容量不足导致 decode 阶段可服务请求数受限。

优化方向：

• 张量并行：模型分片到多卡。
• 量化：FP16 到 INT8/INT4，减少权重和 KV cache 占用。
• ZeRO / FSDP：切分参数、梯度、优化器状态。
• Offload：不活跃数据放到 CPU/NVMe。
• PagedAttention + KV cache 压缩：显存超配 + 按需分配。

通信瓶颈（Communication-bound）

定义：GPU 间的数据同步和传输成为瓶颈，GPU 时间花在等待通信完成上。

典型场景：大规模张量并行 all-reduce、流水线并行跨卡激活传递、跨节点训练梯度同步、KV cache 跨卡/跨节点迁移。

优化方向：

• 通信与计算重叠：NCCL 和 kernel 在不同 stream 上并行。
• 梯度压缩或量化：减少通信数据量。
• 拓扑优化：优先利用 NVLink/NVSwitch，再考虑 PCIe/RDMA。
• 减少通信次数：梯度累积、增大 bucket size。
• Ring Attention / 序列并行：减少跨卡通信量。

典型现象

排查顺序：先用 nvidia-smi 看有没有活，再用 Nsight Systems 看谁占时间，然后用 Nsight Compute 看单个 kernel 瓶颈在哪，最后用扩展性测试判断是否通信瓶颈。

关联模块

• 性能指标：Roofline、TFLOPS、HBM 带宽是分类基础。
• 利用率诊断：从 GPU-Util 继续深入到 SM Active、Occupancy、Warp Stall。
• GPU 互联与数据路径：通信瓶颈必须结合 PCIe/NVLink/RDMA/NUMA。
• LLM 推理系统：prefill/decode 的瓶颈类型不同。

一句话结论

GPU 性能预测的核心是把硬件、模型和部署策略转成可解释特征，再预测 step time、MFU、peak memory、communication time 等标签。面试里要强调：GPU 指标不是孤立监控项，而是模型性能、调度决策和容量规划的输入。

核心概念

系统链路

常见误区

关联模块

• 性能指标：提供 TFLOPS、HBM、Roofline 等基础指标定义。
• 利用率诊断：解释 GPU-Util、SM Active、Occupancy 的可预测性边界。
• 分布式训练：并行策略和通信时间是 step time 预测的关键部分。
• 调度与集群：预测结果可用于 bin packing、准入控制和性能回归检测。

一句话结论

GPU 面试不要只背硬件名词，要把问题落到三条主线：数据怎么流动、计算怎么铺满、资源怎么共享和诊断。能把这三条讲清楚，就能自然连接 LLM 推理、分布式训练、K8S GPU 调度和性能预测。

30 秒回答模板

GPU 适合 AI 的核心原因是矩阵计算密集、数据并行度高、显存带宽高。面试回答时我会先区分 CPU 和 GPU：CPU 擅长控制流和低延迟串行逻辑，GPU 擅长用大量 SM、warp 和 Tensor Core 做吞吐型计算。然后补充三类工程问题：第一，数据路径上要关注 HBM、PCIe、NVLink、RDMA 和 NUMA；第二，执行模型上要关注 kernel、grid/block/thread、warp、stream 和 occupancy；第三，性能诊断上要用 nvidia-smi、Nsight Systems、Nsight Compute、DCGM 去判断是计算、显存、通信、数据加载还是 kernel launch 瓶颈。

2 分钟展开模板

如果面试官继续追问，可以按下面顺序展开：

1. 硬件层：GPU 由多个 SM 组成，SM 内有 CUDA Core、Tensor Core、register file、shared memory 等资源；显存是 HBM，带宽远高于 CPU DDR，但容量更小。
2. 执行层：CPU 负责 host 侧控制和 kernel launch，GPU 执行 device 侧 kernel。Kernel 被组织成 grid、block、thread，warp 是硬件调度的基本执行单元。
3. 内存层：global memory 容量大但慢，shared memory/register 更快但容量小。优化目标是提高数据复用、访存连续性和 Tensor Core 利用率。
4. 数据路径：单机内优先走 NVLink/NVSwitch，CPU-GPU 和 GPU-NIC 通常走 PCIe，跨机训练依赖 RDMA/InfiniBand/RoCE 和 GPUDirect RDMA。
5. 共享隔离：MIG 是硬件切分，隔离强但粒度粗；MPS 是软件并发，灵活但有干扰；time-slicing 主要解决调度复用，不保证性能隔离。
6. 诊断层：GPU-Util 只能说明时间上是否有活，不能说明是否高效；要结合 SM Active、Occupancy、Memory Throughput、Tensor Core Util、Warp Stall 和业务吞吐延迟判断。

高频追问压缩表

排障 / 设计决策树

GPU-Util 长期低 | 先怀疑数据加载、CPU preprocessing、H2D、batch 太小、请求量不足
GPU-Util 高但业务慢 | 看 Nsight Systems timeline，判断 kernel、memcpy、NCCL、同步谁占时间
Kernel 占主时间 | 用 Nsight Compute 看 SM Active、Occupancy、Memory Throughput、Tensor Core、Warp Stall
Memory Throughput 高且 Compute 低 | memory-bound，优化访存、KV cache、算子融合、FlashAttention/PagedAttention
Compute Throughput 高且 Tensor Core 高 | compute-bound，考虑更低精度、更大 batch、算子选择和并行策略
NCCL/通信占比高 | 查拓扑、NVLink/RDMA、NUMA/NIC 亲和、bucket、overlap、并行策略
Memcpy 或同步很多 | 查 .item()/.cpu()、DataLoader、pinned memory、stream、CPU-GPU 来回转换

跨模块关联