0. 心智模型：吞吐机器与延迟机器

GPU 不是“更快的 CPU”。它是一次芯片面积上的重新分配：砍掉让单线程变快的一切，换成更多的执行单元。理解了这个取舍，后面所有“为什么这么写”都有了解释。

0.1 两种不同的优化目标

CPU 的每一代改进，大多服务于“让一个线程更快”：乱序执行、分支预测、深流水、巨大的多级 cache、高主频。GPU 反过来：核心简单、频率低、每 SM 的 cache 很小，但把成百上千个执行单元堆在一起，用海量并行把延迟藏起来。

用一个粗糙的模型描述任务耗时：

其中 \(W\) 是总工作量，\(p\) 是有效并行度。CPU 把第一项压到极小（一次 cache 命中只要几纳秒）；GPU 主动接受很大的第一项（单线程慢、访存 500 ns、kernel 启动 5 µs），换取第二项的分母极大。

注意最后两行：GPU 的带宽约为 CPU 的 5–10 倍，算力约为 10–50 倍，但时钟更低、缓存更小。所以 GPU 的强项是规则、密集、可并行的计算，而不是“所有程序都变快”。

0.2 三个直接后果

1. 并行度必须足够。一个 SM 最多常驻 64 个 warp（2048 线程），靠在这些 warp 之间轮换，把一条指令 500 ns 的访存延迟藏起来。如果只有几百个线程，SM 会大量空转。
2. 分支变贵。一个 warp 的 32 个 lane 共享一条指令流。如果 lane 走了不同分支，硬件只能带着掩码把两条路径各执行一遍（见 1.2）。
3. 数据要自己搬。CPU 的 cache 由硬件自动管理；GPU 的 shared memory 是程序员手动管理的 scratchpad。更重要的是显存与主机内存物理分离，拷贝要过 PCIe，单向只有 32–64 GB/s。

0.3 为什么“延迟”这么难藏：Little 定律

要跑满带宽，必须有足够多的请求同时在路上。设带宽为 \(BW\)、单次访存延迟为 \(t_{\text{lat}}\)，则需要的在途字节数为

代入 H100：\(3.35\times10^{12}\times600\times10^{-9}\approx 2.0\) MB。也就是任意时刻要有 2 MB 的读请求在飞。平均到每个 SM 约 15 KB；一个 warp 一次读 128 B，就需要约 120 个 warp 同时在等——超过一个 SM 能常驻的 64 个 warp。所以真实 kernel 必须让每个线程一次读更多（比如 float4，一个 warp 512 B），或者让每个线程同时发多个独立请求（ILP），否则永远跑不满带宽。第 2.2 节会把这条账算完。

0.4 什么时候不该用 GPU

• 数据太小：\(N\) 只有几万时，传输加启动的开销远超计算本身（见 0.5 末尾的时间账）。
• 分支密集、指针追逐：树遍历、解析器、哈希表、图算法。warp 内路径一散，并行度就塌了。
• 低延迟单请求：交互式单条推理、交易系统。GPU 的尾延迟由 launch 和调度决定，不以微秒计。
• 需要复杂同步：全局锁、频繁的跨 block 通信。GPU 没有 mutex，跨 block 自旋还可能死锁（3.3）。
• 双精度密集：消费级显卡的 FP64 被大幅削减，需要 FP64 时先查规格。
• 显存放不下：显存不能超卖，放不下就是硬失败，不会像 CPU 那样换页。

反过来说，GPU 擅长：大矩阵乘与卷积、逐元素运算、归约、扫描、排序、随机数、规则网格上的 stencil。它们的共同点是同一段代码作用在大量数据上，且每个元素的工作量相近。

0.5 第一个程序：向量加

先跑通完整链路：分配显存、拷贝、启动 kernel、拷回、检查错误。这段代码刻意写全了错误检查，因为 GPU 的错误是异步的，不查就等于没写。

vecadd.cu

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

#define CK(x) do { cudaError_t e_ = (x); if (e_ != cudaSuccess) { \
    fprintf(stderr, "%s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(e_)); \
    std::exit(1); } } while (0)

__global__ void vecadd(const float* __restrict__ a,
                       const float* __restrict__ b,
                       float* __restrict__ c, size_t n) {
    size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
    size_t stride = (size_t)blockDim.x * gridDim.x;
    for (; i < n; i += stride)
        c[i] = a[i] + b[i];
}

int main() {
    const size_t n = 1u << 24;                 // 16M 个 float = 64 MiB
    const size_t bytes = n * sizeof(float);
    float* h = (float*)std::malloc(bytes);
    for (size_t i = 0; i < n; ++i) h[i] = (float)i;

    float *a, *b, *c;
    CK(cudaMalloc((void**)&a, bytes));
    CK(cudaMalloc((void**)&b, bytes));
    CK(cudaMalloc((void**)&c, bytes));
    CK(cudaMemcpy(a, h, bytes, cudaMemcpyHostToDevice));
    CK(cudaMemcpy(b, h, bytes, cudaMemcpyHostToDevice));

    int threads = 256;
    int blocks  = (int)((n + threads - 1) / threads);
    if (blocks > 65535) blocks = 65535;        // 更大的 N 交给 grid-stride
    vecadd<<<blocks, threads>>>(a, b, c, n);
    CK(cudaGetLastError());                   // 启动阶段的错误
    CK(cudaDeviceSynchronize());              // 执行阶段的错误

    CK(cudaMemcpy(h, c, bytes, cudaMemcpyDeviceToHost));
    printf("c[%zu] = %.1f\n", n - 1, h[n - 1]);   // 2 * (n-1)
    cudaFree(a); cudaFree(b); cudaFree(c); std::free(h);
    return 0;
}

编译与运行：

nvcc -O3 -arch=native vecadd.cu -o vecadd && ./vecadd

三处值得注意：grid-stride 循环让 grid 可以封顶，从而兼容任意 \(N\)；索引用 size_t，因为 \(2^{31}\) 个 float 就是 8 GiB，int 会溢出；启动错误和执行错误要分两次查，cudaGetLastError() 只能看到参数错误，越界和非法访存要等同步才报。

同样的计算用 CPU 写，只差一行：

vecadd.cpp

#pragma omp parallel for
for (size_t i = 0; i < n; ++i)
    c[i] = a[i] + b[i];

两段代码逐字对应，但执行模型完全不同：CPU 把循环的迭代分给 8–128 个线程，每个线程串行跑一大段；GPU 把同一个循环体发给几十万线程，每个线程只做很少几次迭代。写 kernel 的本质是把循环拆成“索引 + 单元素操作”，这一点在第 1 节会反复出现。

最后算一笔时间账，它解释了为什么 GPU 不是万能的。16M 个 float 的向量加：读 \(a,b\) 写 \(c\)，共 192 MiB 流量，在 3.35 TB/s 下约 60 µs；但一次性把数据从主机搬过去再搬回来，同样 192 MiB 走 PCIe 4.0 x16（32 GB/s）要 约 6 ms——传输是计算的 100 倍。只有数据常驻显存、被反复使用时，GPU 的优势才存在。

0.6 和 CPU 一样的地方

• 浮点规则：IEEE 754，单精度 24 位有效位、相对误差量级 \(10^{-7}\)；GPU 可以开启 FTZ 把非规格化数刷成 0（和 CPU 的默认行为可能不同，见 7.4）。
• 整数与字节序：补码、位运算、小端，与 x86/ARM 一致。
• C++ 语义与未定义行为：数据竞争在两边都是 UB，不会因为“在 GPU 上”就变安全。
• 局部性与 roofline：先看数据复用，再谈优化；这些分析方法原样适用（第 5、7 节）。
• 先测量再优化：直觉在 GPU 上错得更离谱，profiler 更重要。
• 优先用库：cuBLAS / cuDNN / CUTLASS 对应 BLAS / MKL / 手写 SIMD；99% 的场景不该自己写 GEMM。

0.7 容易写错的地方