AI 工程基础体系 · 第 65/100 篇。内容覆盖机器学习、深度学习与生成式 AI；模型、数据、评测、权限和成本会作为同一生产系统处理。

AI GPU 与 CUDA 基础：核函数、显存、带宽、算子和性能证据

GPU 性能问题经常被简化成“显卡不够快”或“显存不够大”。这种判断通常缺少中间层：一个 PyTorch 算子如何变成 CUDA 核函数，线程如何访问显存，访问量如何受带宽限制，计算如何受算力限制，以及一次优化是否真的改善了端到端延迟，都需要可验证的因果链。

本文覆盖机器学习、深度学习和生成式 AI 中常见的 GPU 执行基础，并把模型、输入数据、评测方法、权限边界和成本视为同一个生产系统的一部分。

一、先建立概念链：模型代码如何到达 GPU

1. AI GPU 是什么

AI GPU 不是一种严格统一的硬件标准，而是指适合机器学习工作负载的 GPU。它通常具有：

大容量、高带宽的设备内存，例如 HBM 或 GDDR；
大量可并行执行的计算单元；
面向矩阵乘法和低精度计算的专用单元，例如 Tensor Core；
支持 CUDA、ROCm 或其他 GPU 编程平台；
针对 FP32、FP16、BF16、TF32、INT8 等数据类型提供不同的计算能力。

GPU 的基本优势不是“任何代码都更快”，而是能够同时运行大量相似的计算。矩阵乘法、卷积、向量运算和 Transformer 中的注意力计算具有较强的数据并行性，因此适合 GPU。

如果一个任务主要包含串行分支、频繁的小对象分配、复杂指针追踪或大量 CPU 系统调用，GPU 可能无法取得优势。GPU 的并行计算能力必须与足够大的工作量、合适的内存访问和较低的调度开销共同成立。

2. CUDA 是什么

CUDA 是 NVIDIA GPU 的并行计算平台和编程模型。它包括：

编程语言扩展，例如 __global__、__device__；
编译工具链，例如 nvcc；
运行时 API，例如内存分配、内核启动、流和事件；
库，例如 cuBLAS、cuDNN、NCCL；
由 PyTorch 等框架调用的驱动和运行时接口。

CUDA 不是“GPU 本身”，也不是所有 GPU 都支持的通用标准。CUDA 程序依赖 NVIDIA 硬件和兼容的驱动、运行时及库版本。PyTorch 的 CUDA 版本、NVIDIA 驱动版本和 GPU 架构之间也存在兼容约束。

一个典型调用链是：

Python 模型代码
    ↓
PyTorch Tensor / Operator
    ↓
ATen、cuBLAS、cuDNN 或自定义扩展
    ↓
CUDA kernel launch
    ↓
GPU 上的线程块和线程
    ↓
寄存器、共享内存、显存中的数据读写

例如：

z = torch.matmul(x, w)

这不是直接执行一条“矩阵乘法指令”。PyTorch 先根据张量设备、数据类型、形状、布局和后端选择实现，随后可能调用 cuBLAS 或其他实现。底层实现会启动一个或多个 GPU 核函数；在支持的硬件和配置下，矩阵乘法的一部分可能使用 Tensor Core。

因此，看到 Python 代码中的一行操作，并不能直接知道它对应几个核函数、使用了哪些硬件单元或读取了多少显存。

二、核函数：GPU 真正执行的工作单元

1. 核函数的定义

核函数（kernel function） 是在 GPU 上执行的函数。CUDA 中通常使用 __global__ 声明：

__global__ void add_kernel(const float* a,
                           const float* b,
                           float* c,
                           int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

这里的函数体不是只执行一次。主机端启动核函数后，GPU 会创建许多线程，让不同线程处理不同的 i。

启动形式通常类似：

int threads = 256;
int blocks = (n + threads - 1) / threads;

add_kernel<<<blocks, threads>>>(a, b, c, n);

变量含义如下：

threadIdx.x：线程在线程块内的索引；
blockDim.x：一个线程块中的线程数；
blockIdx.x：线程块在网格中的索引；
blockIdx.x * blockDim.x + threadIdx.x：全局元素索引。

当 n = 1000、threads = 256 时：

blocks = ceil(1000 / 256) = 4

总共启动 1024 个线程，其中索引 1000 到 1023 的线程必须通过 if (i < n) 避免越界。这说明线程数量通常按块大小向上取整，而不是精确等于元素数量。

2. Grid、Block 和 Thread

CUDA 的执行组织有三层：

Grid
 ├── Block 0
 │    ├── Thread 0
 │    ├── Thread 1
 │    └── ...
 ├── Block 1
 └── ...

Grid：一次核函数启动的全部线程块；
Block：可以协作、同步并共享一块 shared memory 的线程集合；
Thread：执行核函数中一条控制路径的最小逻辑单位。

线程块会被调度到 GPU 的某个 SM（Streaming Multiprocessor）上。一个线程块不会在多个 SM 之间拆开执行；不同线程块通常可以并行执行，但实际并发数量受寄存器、shared memory、线程数和硬件资源限制。

线程块内可以使用：

__syncthreads();

进行屏障同步。但这个同步只对同一个 block 内的线程有效。不同 block 之间不能直接使用普通 CUDA 屏障同步；如果需要全局阶段同步，通常要结束一个核函数后再启动下一个核函数，或使用特定的协作组机制。

3. CUDA 调用是异步的

默认情况下，CPU 发起 CUDA 核函数或内存操作后，通常不会等待 GPU 完成：

y = model(x)
print("可能在 GPU 计算完成前执行")

因此 CPU 端的普通计时可能只测到“提交任务”的时间，而不是 GPU 执行时间。

常见同步边界包括：

torch.cuda.synchronize()

以及 CUDA Event。异步性还影响错误报告：某个核函数内部发生错误，可能直到后续 API 调用或同步点才被报告。

调试异步错误时，可以使用：

CUDA_LAUNCH_BLOCKING=1 python train.py

这会让 CUDA 调用更接近同步执行，使报错位置更容易定位，但会显著改变程序时序和性能，不能把这种模式下的耗时当作生产性能。

三、显存与内存层次

1. 显存不等于 CPU 内存

显存通常指 GPU 设备上的内存，例如 HBM 或 GDDR。GPU 运行核函数时，输入张量、权重和中间结果通常需要位于 GPU 可访问的设备内存中。

CPU 内存和 GPU 显存是不同的资源：

x_cpu = torch.randn(1024, 1024)
x_gpu = x_cpu.to("cuda")

x_cpu 存在主机内存，x_gpu 存在设备内存。这个 .to("cuda") 可能涉及 PCIe 或 NVLink 数据传输，通常远慢于 GPU 内部显存访问。

为了减少反复传输，训练和推理通常应保持如下数据流：

CPU 数据读取
  ↓
批量整理、预处理
  ↓
一次传入 GPU
  ↓
模型的多个算子在 GPU 上连续执行
  ↓
只在必要时把结果传回 CPU

但“全部放到 GPU”也不是无条件成立：数据集可能大于显存，或者多个服务需要共享设备，此时应采用分批加载、缓存、流水线或分布式策略。

2. 显存层次决定访问成本

GPU 内部不是只有一种内存：

层次	典型特征	主要用途
Register	每线程私有，容量小，访问快	临时标量、地址、累加器
Shared memory	每个 block 可见，片上，容量有限	block 内数据复用、通信
L1 / L2 cache	硬件缓存	缓存重复访问
Device memory	容量大，但延迟和带宽相对较低	模型权重、输入、中间张量
Host memory	CPU 内存	数据集、输入队列、结果

这里的“快慢”不是单个固定数字。访问延迟、并发请求数量、缓存命中率、访问是否合并、当前负载和 GPU 型号都会影响实际结果。

3. PyTorch 的 allocated、reserved 和可见显存

PyTorch CUDA 缓存分配器会保留部分已经申请过的显存，以便后续复用。因此以下概念不同：

allocated：当前由 PyTorch 张量实际占用的显存；
reserved：PyTorch 分配器向 CUDA 申请并保留的显存；
free：驱动视角下当前未被分配的显存；
显存容量：设备总物理容量。

可以查看：

import torch

if not torch.cuda.is_available():
    raise RuntimeError("CUDA 不可用")

print(torch.cuda.get_device_name())
print(torch.cuda.memory_summary())

nvidia-smi 展示的进程显存占用与 PyTorch 的 allocated 并不一定相等，因为缓存分配器、CUDA 上下文和其他库也可能占用显存。

显存不足有多种原因：

模型参数本身过大；
激活值随 batch size 或序列长度增长；
反向传播保存了训练所需的中间结果；
优化器状态，例如 Adam 的一阶、二阶矩；
临时工作区；
内存碎片；
其他进程占用显存。

因此，单纯删除一个 Python 变量未必立即让 nvidia-smi 的占用下降；缓存分配器可能仍然保留这部分空间用于复用。

四、带宽：数据移动能有多快

1. 带宽的定义

**带宽（bandwidth）**表示单位时间内可以传输的数据量，常见单位是 GB/s 或 TB/s。

对于一次数据移动，实际有效带宽可以近似计算为：

B_{\text{effective}} = \frac{D}{t}

其中：

$D$ 是实际读写的数据量；
$t$ 是完成这些读写所需的时间；
$B_{\text{effective}}$ 是有效带宽。

理论显存带宽通常由内存总线宽度和内存数据率推导：

B_{\text{theoretical}} = \frac{\text{bus width}}{8} \times \text{memory data rate}

其中总线宽度以 bit 计，除以 8 后得到 Byte。实际程序通常达不到理论值，因为还受到访问模式、缓存、指令调度、并发度、边界处理和其他竞争的影响。

PCIe、NVLink 和 GPU 显存带宽是不同链路的带宽，不能混为一谈。把 CPU 张量拷贝到 GPU 的瓶颈可能是 PCIe，而 GPU 内核读取权重的瓶颈可能是设备显存。

2. 合并访问

GPU 通常以一组相邻线程的方式处理内存请求。当相邻线程访问相邻地址时，硬件更容易合并内存事务：

// 适合连续数组
int i = blockIdx.x * blockDim.x + threadIdx.x;
c[i] = a[i] + b[i];

如果线程访问地址间隔很大，例如：

c[i] = a[i * stride] + b[i * stride];

当 stride 较大时，相邻线程访问的地址不连续，可能需要更多内存事务，导致有效带宽下降。

这也是张量布局重要的原因。一个二维张量的 shape 描述维度大小，stride 描述沿每个维度移动一个元素需要跨过多少底层存储位置。transpose 往往只改变视图和 stride，并不立即复制数据；后续算子若不适合这种非连续布局，可能触发隐式拷贝或采用更低效的访问方式。

可以检查：

import torch

x = torch.randn(1024, 2048, device="cuda")
y = x.t()

print(x.is_contiguous())  # 通常为 True
print(y.is_contiguous())  # 通常为 False
print(y.stride())

调用：

y_contiguous = y.contiguous()

会创建连续副本，改善某些后续访问，但也会增加一次显存读写和额外显存占用。因此不能把 .contiguous() 当作无成本修复。

五、算子：框架中的计算语义与实现单位

1. 算子的定义

**算子（operator）**是对张量执行特定数学变换的抽象，例如：

加法、乘法和归约；
矩阵乘法；
卷积；
Softmax；
LayerNorm；
Attention；
激活函数；
张量重排和类型转换。

算子描述“计算什么”，核函数描述“如何在 GPU 上执行其中一部分计算”。一个算子可能对应：

一个 CUDA 核函数；
多个核函数；
一个高度优化的库调用；
在编译或融合后生成的新核函数；
仅在 CPU 上执行的部分逻辑。

例如：

y = torch.relu(x + bias)

在朴素实现中可能先启动加法核函数，再启动 ReLU 核函数，并把中间结果写回显存。如果编译器或框架将它们融合，就可能由一个核函数完成：

读取 x 和 bias
  ↓
在线完成加法和 ReLU
  ↓
一次写出 y

融合的主要收益之一是减少中间张量的显存读写和核函数启动次数，而不是改变数学结果本身。

2. 算子正确性不只取决于公式

同一个数学算子，实际行为还取决于：

输入和输出 dtype；
张量设备；
shape；
stride 和布局；
是否允许广播；
是否需要梯度；
是否使用确定性算法；
是否存在数值稳定性处理；
GPU 架构和库版本。

例如 Softmax 不能简单实现为：

\operatorname{softmax}(x_i)=\frac{e^{x_i}}{\sum_j e^{x_j}}

当 $x_i$ 很大时，指数可能溢出。实际实现通常先减去最大值：

m=\max_j x_j

\operatorname{softmax}(x_i) = \frac{e^{x_i-m}} {\sum_j e^{x_j-m}}

减去同一个 $m$ 不改变数学结果，但能改善数值稳定性。低精度计算还可能使用更高精度累加，具体行为依赖实现和配置。

3. 算子融合的反例

融合并不总是更快。以下情况可能使融合收益消失：

融合后的核函数寄存器使用量过大，降低并发度；
融合逻辑包含分支，导致线程执行路径分化；
原先的中间结果可以被缓存复用，融合后反而重复计算；
融合核函数不再使用高度优化的专用库；
算子规模很小，启动开销占主导；
融合增加编译时间或缓存失效。

因此，“减少 kernel 数量”是一个机制假设，不是性能保证。必须用同一输入、同一精度和同一正确性标准测量。

六、计算量、带宽和算力：用 Roofline 建立第一判断

1. 算术强度

**算术强度（Arithmetic Intensity）**定义为：

I=\frac{F}{D}

其中：

$F$ ：执行的浮点操作数量，单位 FLOP；
$D$ ：从相关内存层次读取和写入的数据量，单位 Byte；
$I$ ：单位数据移动所完成的计算量，单位 FLOP/Byte。

算术强度低的操作通常更容易受带宽限制；算术强度高的操作通常更可能受计算吞吐限制。但这只是上界分析，实际还会受到访问效率和硬件利用率影响。

2. 向量加法的完整算例

考虑：

z_i=a_i+b_i

假设：

$a$ 、 $b$ 、 $z$ 都是 FP32；
每个元素需要读取 $a_i$ 和 $b_i$ ，写入 $z_i$ ；
每次加法计为 1 FLOP。

每个元素的数据移动量是：

D_{\text{element}}=4+4+4=12\text{ Byte}

每个元素的计算量是：

F_{\text{element}}=1\text{ FLOP}

所以：

I=\frac{1}{12}\approx0.0833\text{ FLOP/Byte}

如果某设备的有效显存带宽上限近似为 1 TB/s，则带宽给出的性能上界为：

P_{\text{bandwidth}} = 1\times10^{12}\times\frac{1}{12} \approx83.3\times10^9\text{ FLOP/s}

这个数看起来只有约 83 GFLOP/s，但这并不代表 GPU 的浮点单元很弱，而是向量加法的计算量太少，主要时间花在搬运 12 Byte 数据上。

如果把两个操作融合：

z_i = relu(a_i + b_i)

仍然可以一次读取两个输入、一次写出结果，而不是把加法中间结果写出再读回。融合减少了内存流量，可能提高有效带宽利用率。

3. 矩阵乘法的对照

对矩阵：

C_{m\times n}=A_{m\times k}B_{k\times n}

理想计算量约为：

F=2mnk

因为每个输出元素包含 $k$ 次乘法和 $k$ 次加法。

如果只按一次读取 A、一次读取 B、一次写出 C 的朴素数据量估算，FP32 数据移动量约为：

D=4(mk+kn+mn)

算术强度为：

I=\frac{2mnk}{4(mk+kn+mn)}

当 $m,n,k$ 较大时，分子按三维增长，分母按二维增长，算术强度会变高。实际高性能矩阵乘法还会把数据分块加载到 shared memory 和寄存器中，让同一元素被多个线程重复使用，从而减少对显存的访问。

这解释了一个常见现象：

大矩阵乘法通常能接近计算吞吐上限；
小矩阵乘法可能被核函数启动、布局转换或并行度不足限制；
形状不规则的矩阵乘法可能无法充分利用 Tensor Core；
即使 GPU 利用率很高，端到端延迟仍可能被 CPU、数据传输或其他算子限制。

4. Roofline 上界

Roofline 模型用以下不等式给出粗略性能上界：

P\leq\min(P_{\text{peak}}, B_{\text{effective}}\times I)

其中：

$P_{\text{peak}}$ ：设备在该数据类型和计算路径上的峰值计算吞吐；
$B_{\text{effective}}$ ：有效内存带宽；
$I$ ：算术强度。

当：

B_{\text{effective}}\times I<P_{\text{peak}}

任务更可能是带宽受限；反之则更可能是计算受限。

这不是性能证明。它忽略了缓存层次、指令效率、线程分歧、同步、占用率、库实现和数据布局。但它能帮助排除明显错误的优化方向：对一个带宽受限的逐元素算子只增加浮点计算单元，通常不会解决瓶颈。

七、并发、占用率与数据流

1. Occupancy 的含义

**Occupancy（占用率）**通常指一个 SM 上当前驻留线程数与该 SM 最大可驻留线程数的比例。它受以下资源共同限制：

每个 block 的线程数；
每个线程使用的寄存器数量；
每个 block 使用的 shared memory；
SM 的最大线程数和 block 数；
GPU 架构限制。

高 occupancy 有助于隐藏显存延迟，但不等于高性能。一个矩阵乘法核函数可能使用较多寄存器，却通过数据复用获得很高吞吐；强行降低寄存器使用量可能反而增加访存和溢出开销。

2. Stream 和依赖关系

CUDA stream 是 GPU 操作的有序队列。同一 stream 中，操作按照提交顺序建立依赖；不同 stream 中的操作在没有数据依赖和资源冲突时可能重叠。

一个数据流水线可能是：

Stream 0：H2D(batch 1) → kernel(batch 1) → D2H(result 1)
Stream 1：H2D(batch 2) → kernel(batch 2) → D2H(result 2)

但要实现有效重叠，通常还需要：

主机内存使用 pinned memory；
使用异步拷贝；
不存在隐式同步；
GPU 有足够资源同时执行传输和计算；
后续操作确实不依赖尚未完成的数据。

PyTorch DataLoader 中常见的相关配置是：

loader = torch.utils.data.DataLoader(
    dataset,
    batch_size=64,
    pin_memory=True,
    num_workers=4,
)

传输时可以写：

x = x.to("cuda", non_blocking=True)

non_blocking=True 并不保证任何情况下都异步；源内存是否 pinned、当前设备和操作路径都会影响实际行为。配置增加 worker 也不一定更快，过多 worker 可能造成 CPU 争用、文件系统压力和上下文切换。

八、一个可运行的 PyTorch 测量示例

下面示例比较 GPU 上的逐元素加法和矩阵乘法。它使用 CUDA Event 测量 GPU 时间，避免把 CPU 提交时间误认为 GPU 执行时间。

import torch

if not torch.cuda.is_available():
    raise RuntimeError("需要可用的 NVIDIA GPU、驱动和 CUDA 版 PyTorch")

device = torch.device("cuda")
print("device:", torch.cuda.get_device_name(device))

def benchmark(fn, warmup=20, repeat=100):
    # 预热：让库选择、缓存分配和首次初始化不进入正式结果
    for _ in range(warmup):
        fn()
    torch.cuda.synchronize()

    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)

    start.record()
    for _ in range(repeat):
        fn()
    end.record()

    # Event 记录的是 GPU 时间；必须同步等待 end 完成
    end.synchronize()
    total_ms = start.elapsed_time(end)

    return total_ms / repeat

# 逐元素算子：通常有较低算术强度
n = 64 * 1024 * 1024
a = torch.randn(n, device=device, dtype=torch.float32)
b = torch.randn(n, device=device, dtype=torch.float32)

def vector_add():
    return a + b

vector_ms = benchmark(vector_add)
vector_out = vector_add()
torch.cuda.synchronize()

assert torch.allclose(vector_out, a + b)
vector_bytes = 3 * n * torch.tensor([], dtype=torch.float32).element_size()
vector_gbs = vector_bytes / (vector_ms / 1000) / 1e9

print(f"vector add: {vector_ms:.3f} ms, "
      f"estimated effective bandwidth: {vector_gbs:.1f} GB/s")

# 矩阵乘法：通常有更高算术强度
m = n_mat = k = 4096
x = torch.randn(m, k, device=device, dtype=torch.float16)
w = torch.randn(k, n_mat, device=device, dtype=torch.float16)

def matmul():
    return x @ w

matmul_ms = benchmark(matmul, warmup=10, repeat=50)
matmul_out = matmul()
torch.cuda.synchronize()

print(f"matmul: {matmul_ms:.3f} ms, output shape: {tuple(matmul_out.shape)}")

这个示例为什么成立

输入预先放在 GPU 上
测量的是算子执行，而不是 CPU 到 GPU 的拷贝。
先预热
首次调用可能包含 CUDA 上下文初始化、内存分配、库选择算法等额外成本。
使用 CUDA Event
Event 在 GPU 时间线上记录起止位置，比 time.perf_counter() 更适合测量单个 GPU 区间。
调用 end.synchronize()
如果不等待结束 Event 完成，CPU 可能在 GPU 尚未执行完时读取时间。
显式检查结果
性能优化不能以改变计算结果为代价。allclose 的容差也应根据 dtype 和任务数值要求设定。
带宽估算包含假设
vector_bytes = 3 * n * 4 假设一次读取两个输入并写出一个输出。缓存、写分配、融合和实际内存事务可能使这个估算与硬件计数器不同，因此它是有效带宽的近似，不是精确的总线流量。

矩阵乘法使用 FP16 不代表所有内部累加都必须是 FP16。库可能使用 FP32 累加，也可能根据硬件和配置选择 Tensor Core。具体路径需要通过 profiler、库文档和硬件信息确认。

九、性能证据：怎样证明一次优化有效

性能证据不是“GPU 利用率 99%”或“一次运行更快”，而是一组能够复现实验并支持因果判断的数据。

至少应记录：

GPU 型号、驱动和 PyTorch 版本；
dtype、张量 shape、layout；
batch size、序列长度和 padding 策略；
warmup 次数、重复次数和统计方法；
是否包含 H2D、D2H、数据加载和同步；
输出正确性标准；
单次延迟、吞吐量和尾延迟；
显存峰值；
运行期间的功耗、设备占用和实例成本。

1. 用 PyTorch Profiler 查看算子时间

import torch
from torch.profiler import profile, record_function, ProfilerActivity

if not torch.cuda.is_available():
    raise RuntimeError("CUDA 不可用")

device = "cuda"
x = torch.randn(2048, 2048, device=device)
w = torch.randn(2048, 2048, device=device)

# 先让上下文和输入准备完成
torch.cuda.synchronize()

with profile(
    activities=[
        ProfilerActivity.CPU,
        ProfilerActivity.CUDA,
    ],
    record_shapes=True,
    profile_memory=True,
) as prof:
    with record_function("matmul_relu_region"):
        y = torch.relu(x @ w)
        torch.cuda.synchronize()

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

这个输出通常能帮助回答：

哪些算子占据 CUDA 时间；
一个 Python 区域是否对应多个底层操作；
是否发生了额外的类型转换或拷贝；
哪些算子分配了大量临时内存；
CPU 是否在等待 GPU。

self_cuda_time_total 等 profiler 字段和输出格式可能随 PyTorch 版本变化。生产脚本应固定 PyTorch 版本，并在升级后重新验证字段和报告解析逻辑。

2. 用证据区分几类瓶颈

现象	更可能的解释	需要验证的证据
GPU 利用率低、CPU 长时间忙	数据加载、Python 调度或 CPU 预处理受限	CPU profiler、数据加载耗时
GPU 利用率低、显存拷贝时间高	H2D/D2H 或跨 GPU 通信受限	profiler 中 memcpy、链路带宽
GPU 利用率高但算子吞吐低	可能是带宽、布局、分支或低效 kernel	内存吞吐、指令吞吐、访存合并
单个算子快但端到端不快	其他算子、同步或输入管线占主导	完整请求 trace
显存占用下降但延迟上升	低精度转换、重计算或更多 kernel	端到端 trace 和数值验证
`nvidia-smi` 显示高利用率但服务延迟差	利用率不代表请求级尾延迟	p50/p95/p99 和并发测试

设备利用率是粗粒度指标。它不能直接说明 Tensor Core 是否被使用，也不能说明内存访问是否高效，更不能代替请求级延迟和吞吐量。

3. 性能实验必须保持公平

以下比较通常是不公平的：

优化前：batch=1，FP32，包含数据拷贝
优化后：batch=32，FP16，不包含数据拷贝

正确比较至少要固定：

相同输入 shape 和数据分布；
相同精度与数值容差；
相同 warmup；
相同同步边界；
相同是否包含数据传输；
相同并发和服务队列条件。

对生成式 AI，还应固定输入长度、输出长度、KV cache 状态、采样参数和停止条件。Prefill 和 decode 是不同阶段：prefill 通常有较大矩阵计算，decode 每一步处理的 token 数较少，更容易受到 kernel launch、内存访问和 KV cache 带宽影响。

十、常见失败路径与诊断方式

1. “把数据搬到 GPU 就会更快”

错误原因是忽略了传输成本。如果只执行一次很小的算子：

x_gpu = x_cpu.to("cuda")
y_gpu = x_gpu + 1
y_cpu = y_gpu.cpu()

总耗时可能主要来自两次传输，而不是加法。GPU 加速应比较完整的计算阶段，或者让多个算子共享同一次传输。

2. “显存够用，所以不会 OOM”

显存峰值还包括临时张量、梯度、优化器状态和库工作区。训练时，模型参数之外还可能保存：

参数
+ 梯度
+ 优化器状态
+ 激活值
+ 临时 workspace
+ CUDA/PyTorch 运行时开销

减小 batch size、使用梯度累积、激活检查点、混合精度或参数分片，解决的是不同部分的峰值来源，不能笼统称为“节省显存”。

3. “使用 FP16 一定更快”

FP16 或 BF16 可能减少显存流量并启用 Tensor Core，但不保证任何算子都更快：

shape 可能不适合 Tensor Core；
算子可能仍由 FP32 路径执行；
类型转换可能增加额外 kernel；
数值范围可能导致溢出或精度损失；
小规模工作负载可能受启动开销主导。

混合精度需要同时验证性能和数值结果。训练还需要关注 loss scaling、梯度溢出和收敛行为，具体机制取决于 PyTorch 版本和使用的 AMP API。

4. “异步报错位置就是根因位置”

例如某个越界核函数已经提交，但 Python 直到下一次 torch.cuda.synchronize() 才报错。诊断时可以：

CUDA_LAUNCH_BLOCKING=1 python reproduce.py

还应缩小输入、检查 index 边界、检查 dtype 和 device，并在必要时使用 NVIDIA 的内存检查工具。同步调试会改变时序，因此修复后必须关闭该环境变量重新测量。

5. “多开 stream 就一定并发”

不同 stream 只有在没有依赖、没有资源冲突且操作粒度适合时才可能重叠。两个都占满 SM 或显存带宽的 kernel 放到不同 stream，通常不会带来线性加速，反而可能互相争用。

十一、从算子性能到生产成本

生产系统的目标通常不是单个 kernel 的最短时间，而是满足正确性和权限约束下的服务级目标：

\text{成本/请求} \approx \frac{\text{GPU 小时价格}\times\text{GPU 运行时间}} {\text{完成请求数}}

这只是粗略模型。实际还要加入 CPU、存储、网络、数据预处理、重试和空闲时间。

因此优化目标可能有不同优先级：

在线推理关注 p95/p99 延迟和并发下的稳定性；
离线训练关注每步时间、样本吞吐量和总训练成本；
生成式推理还关注首 token 延迟、每 token 延迟和 KV cache 占用；
评测任务关注结果可复现性，不能为了速度随意改变采样、精度或数据顺序。

数据访问权限也属于性能实验的一部分。若基准数据包含受限样本，应保证执行身份有明确读取权限，日志中不要泄露原始数据或敏感输入。否则即使测得更快，也不能作为可部署的生产证据。

十二、把优化判断落实为一条证据链

一个可靠的 GPU 优化过程可以按以下因果顺序进行：

flowchart TD
    A[固定模型、数据、shape、dtype 和并发] --> B[验证输出正确性]
    B --> C[用 CUDA Event 测 GPU 时间]
    C --> D[用 Profiler 定位算子、拷贝和同步]
    D --> E{判断瓶颈}
    E -->|带宽或布局| F[检查访问模式、stride、融合和数据移动]
    E -->|计算吞吐| G[检查 dtype、矩阵形状、库路径和 Tensor Core]
    E -->|启动或调度| H[合并小算子、减少同步、调整批量]
    E -->|CPU 或输入管线| I[检查 DataLoader、预处理和 H2D]
    F --> J[重新测量并检查数值]
    G --> J
    H --> J
    I --> J
    J --> K[比较端到端延迟、吞吐、显存和成本]

关键路径是：

先固定实验条件，避免比较失真；
先验证结果，再谈速度；
用 GPU Event 测量设备时间；
用 profiler 找到具体算子、拷贝或同步；
用算术强度和数据布局解释瓶颈；
只改变一个主要因素；
重新测量端到端指标，而不是只看局部 kernel；
将显存峰值、尾延迟、稳定性和成本一起纳入结论。

GPU、CUDA、核函数、显存、带宽和算子之间的关系可以归纳为：

算子定义数学工作
  ↓
实现选择库或 CUDA 核函数
  ↓
核函数组织线程块和线程
  ↓
线程通过寄存器、shared memory 和显存完成数据流动
  ↓
算术强度决定更接近带宽上限还是计算上限
  ↓
Profiler、事件、硬件计数器和端到端指标提供性能证据

只有当这条链条闭合时，“优化有效”才不再是凭感觉的判断。

系列导航与关联阅读

系列入口：AI 工程完整学习路线：从机器学习与 Transformer 到 RAG、Agent 和生产治理
上一篇：分布式训练：数据并行、模型并行、梯度同步与故障恢复
下一篇：混合精度训练：FP16、BF16、Loss Scaling、溢出与精度验证

官方资料

本文依据研究论文、标准组织与主流框架官方文档重新梳理；正文、示例与工程清单由 WR BLOG 编写。