CS267 · 第 9 课

CUDA 与 GPU 导论

从 GPU 架构的演进出发，理解 CUDA 编程模型的核心抽象——线程、块、网格与内存层次，并掌握写出高效 GPU 代码的关键技术。

---

为什么要用 GPU？

GPU（Graphics Processing Unit）最初为图形渲染设计，但因其极高的吞吐能力，逐渐成为科学计算与机器学习的主力加速器。GPU 并非万能，它擅长的是大规模数据并行任务——对大量数据执行相同或相似的操作。

GPU 硬件的 5 个关键思想

现代 GPU 的设计建立在 5 个核心硬件思想之上，它们共同决定了 CUDA 编程模型：

#	思想	含义
1	简单核心	每个核心只做简单算术/访存，去掉乱序执行、分支预测等复杂逻辑
2	大量核心	上千个核心同时执行，吞吐远超 CPU
3	共享指令流（SIMT）	一组线程（warp）共享同一 PC，以锁步（lock-step）执行
4	分支掩码（masking）	warp 内线程走不同分支时，串行化执行两条路径，屏蔽不符合的线程
5	延迟隐藏（latency hiding）	当一个 warp 等待内存时，调度器切换到另一个就绪的 warp

GPU 硬件细节：SM 与 Warp

以 NVIDIA A100 为例：

• SM（Streaming Multiprocessor）：A100 有 108 个 SM，每个 SM 拥有寄存器文件、共享内存（shared memory）、L1 cache、64 个 FP32 CUDA core 等。
• Warp：SM 以 32 个线程为一组（称为一个 warp）调度执行。Warp 是 GPU 调度的最小单位。
• 线程块（thread block）：程序员指定的线程分组，一个 block 驻留在一个 SM 上，其中的线程可以协作（shared memory、同步）。

CUDA 编程模型

CUDA 是 NVIDIA 推出的并行计算平台与编程模型，本质上是 C/C++ 的一组扩展。核心抽象是 kernel——在 GPU 上执行的函数，由大量线程并行调用。

从 CPU 向量加法到 GPU kernel 的演化

先看一个朴素的 CPU 实现：

void vecAdd(float *a, float *b, float *c, int n) {
  for (int i = 0; i < n; i++)
    c[i] = a[i] + b[i];
}

GPU 版本把循环去掉，让每个线程处理一个元素：

__global__ void vecAdd(float *a, float *b, float *c, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n)
    c[i] = a[i] + b[i];
}

调用方式（host 代码）：

int n = 1<<20;
int threads = 256;
int blocks = (n + threads - 1) / threads;
vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);

线程层次：Thread、Block、Grid

• threadIdx：线程在 block 内的索引（x, y, z）。
• blockIdx：block 在 grid 内的索引。
• blockDim：每个 block 的线程数。
• gridDim：grid 中的 block 数。

全局线程 ID 的计算：i = threadIdx.x + blockIdx.x * blockDim.x。

GPU 内存层次

内存	位置	延迟	生命周期	可见性
寄存器（register）	SM 上	~1 cycle	线程	仅本线程
共享内存（shared memory）	SM 上	~5 cycles	block	同 block 线程
L1 / L2 cache	SM / 芯片	~10–300 cycles	—	自动管理
全局内存（global memory）	HBM 显存	~400–800 cycles	应用	所有线程 + host

在 CUDA 中：

__shared__ float s[256];   // 块内共享内存
__device__ float g[1024];  // 全局内存（设备端声明）
__constant__ float c[64];  // 常量内存（只读，有 cache）

Stencil 计算与 Halo Exchange

Stencil 是科学计算中常见的模式，每个输出元素是输入中一个局部邻域的函数。例如一维 3-point stencil：y[i] = c0*x[i-1] + c1*x[i] + c2*x[i+1]。

在 GPU 上实现时，一个 block 负责一段输出，但需要左右相邻的"halo"元素。通常有两种策略：

• 加载 halo 到 shared memory：每个线程加载自己的元素，额外线程加载 halo，用 __syncthreads() 同步后再计算。
• 每个线程加载多个输入：减少 halo 专用线程的开销。

__global__ void stencil_1d(float *in, float *out, int n) {
  extern __shared__ float s[];
  int gid = threadIdx.x + blockIdx.x * blockDim.x;
  int lid = threadIdx.x + 1;  // 留出 halo 位置

  // 加载内部元素
  s[lid] = (gid < n) ? in[gid] : 0.0f;
  // 加载左 halo
  if (threadIdx.x == 0)
    s[0] = (blockIdx.x > 0) ? in[gid - 1] : 0.0f;
  // 加载右 halo
  if (threadIdx.x == blockDim.x - 1)
    s[lid + 1] = (gid + 1 < n) ? in[gid + 1] : 0.0f;

  __syncthreads();  // 确保所有数据就绪

  if (gid < n)
    out[gid] = c0 * s[lid - 1] + c1 * s[lid] + c2 * s[lid + 1];
}

同步与通信

GPU 上的同步层次：

• Block 内：__syncthreads()、__threadfence_block()。
• Grid 内：传统上需要启动多个 kernel；Pascal+ 支持 cooperative groups 的 grid 级同步。
• Device 与 Host 之间：cudaMemcpy 隐式同步；cudaDeviceSynchronize() 显式同步。

访存合并（Memory Coalescing）

GPU 的内存访问以段（segment）为单位——通常 32 字节或 128 字节。如果 warp 中 32 个线程访问连续且对齐的地址，这些访问可以被合并为一次段事务。这就是coalesced access。

规则要点：

• warp 内线程访问连续地址最好。
• 首地址对齐到 128 字节最好。
• 跨步访问（strided access）要尽量避免；若无法避免，可通过 shared memory 重组。

CUDA 对 C 的扩展速览

扩展	用途
__global__	host 调用、device 执行的 kernel 函数
__device__	device 内部调用的函数
__host__	host 函数（默认）
__shared__	声明 block 内共享内存
__constant__	声明常量内存
__syncthreads()	block 内同步
<<< >>>	kernel 启动配置

此外，CUDA 还提供了：

• 统一内存（Unified Memory）：cudaMallocManaged 让 host 和 device 共享同一指针，由驱动自动迁移页面。
• Streams：允许并发执行多个 kernel 或重叠计算与数据传输。
• 库生态：cuBLAS、cuFFT、cuDNN、Thrust 等。

---

本课要点

---

知识检查

---