线程组织与内存层级

CUDA 把并行性分成三层：grid → block → thread。每个线程执行 同一个 kernel 函数，靠 threadIdx / blockIdx 区分自己处理的数据。

__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];
}
// 1 个 grid，(n+255)/256 个 block，每个 block 256 个线程
add<<<(n + 255) / 256, 256>>>(a, b, c, n);

执行模型：SM 与 warp

硬件上，一个 block 被整体调度到一个流多处理器（SM）上执行。 SM 以 warp（32 个线程） 为最小调度单位，同一 warp 内线程 锁步执行同一条指令：

graph TD
    Grid[Grid] --> B1[Block 0]
    Grid --> B2[Block 1]
    B1 --> W1[SM 按 32 线程切成 warp]
    W1 --> I[warp 内锁步执行同一指令]

分支发散是 warp 模型的第一课：if 两侧的路径会被串行执行， 两边的线程互相等待对方跑完。

内存层级

层级	位置	延迟（量级）	作用域
寄存器	每个线程	~1 周期	线程私有
共享内存	每个 SM	~30 周期	block 内
L2	芯片	~200 周期	全 GPU
全局内存	显存	~500 周期	全 GPU

优化方向几乎总是同一句话：把数据搬得离计算更近。 能放寄存器就别放共享内存，能放共享内存就别反复读全局内存。

小结

线程组织决定暴露给硬件的并行度，内存层级决定访存代价。 下一节看共享内存的正确用法与银行冲突。