GPGPU 架构与大规模并行编程

GPU 最初为图形渲染设计——为屏幕上每个像素独立计算颜色。这个任务的天然并行性决定了 GPU 的架构方向：海量简单核心，而非少量复杂核心。

2006 年 NVIDIA 推出 CUDA，标志着 GPU 从专用图形处理器走向通用并行计算（GPGPU）。今天的 AI 训练、科学模拟、密码学等领域，GPU 已成为标配。

形象地说：CPU 是跑车——少量、昂贵、每辆都很快；GPU 是货运列车——大量、便宜、总运量巨大。你要运一个人，用跑车；你要运一万吨货，用火车。

GPU 的并行执行模型叫做 SIMT（Single Instruction, Multiple Threads），可以看作 SIMD 的「线程化」版本。在 SIMD 中，一条指令同时操作多个数据元素；在 SIMT 中，一条指令同时被多个线程执行，每个线程操作自己的数据。

SIMT 的关键概念：

• Thread（线程）：最小的执行单元，每个线程有自己的寄存器状态和程序计数器
• Warp（束）：32 个线程组成一个 Warp，同一 Warp 内的线程锁步执行同一条指令
• Block（块）：多个 Warp 组成一个 Block，Block 内的线程可以通过共享内存和同步原语协作
• Grid（网格）：多个 Block 组成一个 Grid，对应一次内核启动的所有线程

CUDA 编程的核心是内核函数（Kernel）——在 GPU 上并行执行的函数。调用内核时，程序员指定 Grid 和 Block 的维度：

// 定义内核：每个线程执行一次 add()
__global__ void add(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];
}

// 启动内核：256 个线程/Block，足够多的 Block 覆盖 n 个元素
int blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize;
add<<>>(a, b, c, n);

三级索引的计算：$\text{global\_idx} = \text{blockIdx} \times \text{blockDim} + \text{threadIdx}$。每个线程根据全局索引处理不同的数据元素。

CUDA 的内置变量：

GPU 的内存层级比 CPU 更复杂，但也提供了更多优化空间：

共享内存（Shared Memory）是 CUDA 优化的关键武器。它位于芯片上，延迟接近寄存器，可以被 Block 内所有线程读写。典型用法是先把数据从 Global Memory 加载到 Shared Memory，然后在 Shared Memory 上反复计算，最后写回 Global Memory——类似 CPU 的缓存分块（Tiling）。

SIMT 的一个关键约束：同一 Warp 内的所有线程必须执行同一条指令。如果 Warp 内部分线程走 if 分支，部分走 else 分支，GPU 不能跳过任何线程——它只能串行执行两个分支，禁用不走该分支的线程。

__global__ void bad_branch(float *data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (data[i] > 0)       // 分支发散！
        data[i] = sqrt(data[i]);
    else
        data[i] = data[i] * data[i];
}

如果 Warp 内 16 个线程 data[i]>0，16 个线程 data[i]≤0，那么这个 Warp 需要执行两条路径，有效性能减半。

1. CUDA 的线程模型与 CPU 多线程（pthread/OpenMP）有什么本质区别？为什么 GPU 线程切换是「零开销」而 CPU 线程切换代价高昂？
2. 矩阵乘法是 GPU 的经典 benchmark。分析 naive CUDA 实现 vs 使用 Shared Memory Tiling 的性能差距，瓶颈在哪里？
3. NVIDIA 的 Tensor Core 在 Volta/Turing/Ampere 架构中引入，它与传统 CUDA Core 有什么区别？为什么对 AI 训练至关重要？
4. GPU 的内存带宽（如 H100 的 3 TB/s）远高于 CPU，但延迟也高得多。在什么类型的问题中，高带宽可以弥补高延迟？什么问题不适合 GPU？