【CUDA编程入门①】CUDA编程体系全景：线程模型、内存层次与Warp执行机制

【CUDA编程入门①】CUDA编程体系全景：线程模型、内存层次与Warp执行机制

本文是「CUDA编程与HPC高性能计算」系列的第一篇。从这篇开始，我们正式进入GPU内核编程的世界。CUDA难吗？说难也难——它有一套全新的编程模型。说简单也简单——只要你搞懂了三件事：线程怎么组织、内存怎么分层、Warp怎么执行，CUDA就入门了。

---

一、写在前面：CUDA不是"在GPU上写C++"

很多初学者以为CUDA就是"把C++代码搬到GPU上跑"——这是最大的误解。

CUDA与普通C++最大的区别在于两点：

1. 线程组织模型：GPU不是一个线程在跑，而是几千上万个线程同时在跑，而且这些线程是按特定层级组织的
2. 内存层次结构：GPU不是"一块内存"，而是有寄存器、共享内存、全局内存等多层存储，速度和容量天差地别

忽略了这两点，写出来的kernel可能只有CPU级别的性能——那你用GPU干嘛呢？

先上一张全景图，建立整体印象：

---

二、CUDA线程组织模型：Grid/Block/Thread三级

2.1 三级结构

启动一个CUDA kernel时，会生成一个由线程组成的"网格"：

Grid（一个Kernel的所有线程） ├── Block (0,0) Block (1,0) ... │ ├── Thread(0,0) │ ├── Thread(1,0) │ └── ... └── Block (0,1) └── ...

从大到小三级别：

层级	含义	维度	协作能力
Grid	整个kernel的所有线程	最多3维（gridDim.x/y/z）	全局协作需要原子操作或多kernel
Block	一组线程	最多3维（blockDim.x/y/z）	同Block内可通过共享内存+__syncthreads()协作
Thread	单个执行单元	线程索引threadIdx	独立执行

最关键的一点：同一个Block内的线程可以协作（共享内存、同步），不同Block之间不能直接协作。

2.2 线程索引怎么算？

写CUDA kernel的第一步，就是搞清楚"我是哪个线程，我该处理哪个数据"。

一维索引（最简单）

intidx=blockIdx.x*blockDim.x+threadIdx.x;

就是：第几块 × 每块多少线程 + 块内第几号线程。

二维索引（图像处理等场景）

introw=blockIdx.y*blockDim.y+threadIdx.y;intcol=blockIdx.x*blockDim.x+threadIdx.x;intwidth=gridDim.x*blockDim.x;intidx=row*width+col;

行和列分别算，然后展平成一维索引。

💡小白提示：索引计算是kernel正确性的基础。写错了索引，结果肯定不对。建议刚开始学的时候，画个小例子（比如4个block，每个block 4个线程），手动算几个索引验证一下。

2.3 Block大小怎么选？

Block的大小（每个Block有多少线程）不是随便选的，有几个原则：

1. 必须是32的倍数——因为Warp大小是32（下一节讲）
2. 常见选择：128、256、512、1024
3. 受限于三个资源：
   • 寄存器数量（每个线程用多少寄存器）
   • 共享内存大小（每个Block用多少共享内存）
   • 最大线程数（通常一个Block最多1024线程）

256是最常用的"万金油"选择——平衡了资源占用和并行度。

---

三、CUDA内存模型：五层存储，逐级变慢

GPU的内存是分层的，从快到慢、从近到远：

┌─────────────────────────────────────┐ │ Host（CPU）内存 │ 可访问，通过cudaMemcpy与Device交互 ├─────────────────────────────────────┤ │ Global Memory（全局内存） │ 所有线程可访问，容量大，延迟高（几百cycle） ├─────────────────────────────────────┤ │ Constant Memory（常量内存） │ 只读，有缓存，适合广播式读取 ├─────────────────────────────────────┤ │ Shared Memory（共享内存） │ 同Block内线程共享，低延迟（~30 cycle） │ 分为32个bank │ 需要避免Bank Conflict ├─────────────────────────────────────┤ │ Registers（寄存器） │ 每个线程私有，最快，数量有限 └─────────────────────────────────────┘

我们逐层来看：

3.1 全局内存（Global Memory）

• 所有线程都能访问，容量就是显存大小（比如80GB）
• 访问延迟高（几百个cycle）
• **合并访问（Coalesced Access）**可以高效利用带宽——相邻线程访问相邻地址时，硬件会把多个请求合并成一个

💡 全局内存是"最大但最慢"的一层。优化的核心就是：尽量少访问全局内存，尽量让数据停留在离计算更近的地方。

3.2 共享内存（Shared Memory）

• 同一个Block内的所有线程都能访问
• 容量很小（每个SM几十KB到一百多KB）
• 速度很快（~30 cycles，比全局内存快一个数量级还多）
• 按32-bit划分为32个bank——如果同一个Warp内多个线程同时访问同一个bank的不同地址，会触发Bank Conflict（存储体冲突），访问被串行化

典型用途：

• 矩阵乘法分块（Tiling）
• 数据复用（多次读取的数据先加载到共享内存）
• 线程间通信（同Block内交换数据）

3.3 寄存器（Registers）

• 每个线程私有，其他线程看不到
• 速度最快（1个cycle）
• 数量有限——如果用太多，会发生寄存器溢出，多余的数据放到Local Memory（其实就是全局内存的一部分），性能急剧下降

🔍怎么看寄存器用了多少：编译时加上--ptxas-options=-v参数，编译器会打印每个kernel用了多少寄存器、多少共享内存。

3.4 常量内存（Constant Memory）

• 只读的，有常量缓存
• 适合广播式读取（同一个Warp内所有线程读同一个地址）
• 如果线程访问不同地址，会串行化，反而更慢

---

四、Warp执行模型：32个线程齐步走

4.1 Warp是什么？

Warp（ warp是"一绺"的意思，这里翻译成"线程束"）是GPU调度的最小单位。

• 一个Warp固定包含32个线程
• 同一个Warp内的线程执行完全相同的指令
• 线程以**lock-step（锁步）**方式执行——大家步调一致，一起走

可以这样理解：GPU的"调度器"调度的不是单个线程，而是一个个Warp。Warp才是GPU调度的基本单位。

4.2 分支发散（Branch Divergence）

既然同一个Warp内的线程执行相同的指令，那如果它们走了不同的分支呢？比如：

if(threadIdx.x%2==0){// 路径A：偶数线程走这里}else{// 路径B：奇数线程走这里}

答案是：硬件会先执行路径A，走路径B的线程被mask掉（不干活）；然后再执行路径B，走路径A的线程被mask掉。

结果虽然正确，但性能下降了——因为两个路径是串行执行的。

这就叫分支发散（Branch Divergence）。

怎么避免？

• 让同一个Warp内的线程尽量走同一路径
• 按Warp对齐分配任务
• 用Warp内的优化技巧（比如Warp Shuffle，后面讲）

💡小白误区：不是所有分支都会导致发散。如果整个Warp都走同一个分支，就不会发散。比如if (idx < N)这种边界判断，通常只有边界上的那个Warp会发散，大部分Warp不受影响。

---

五、三大关键优化概念

5.1 合并访问（Coalesced Access）

全局内存访问的理想模式：

理想：Thread i 访问地址 base + i * 4 访存请求可被合并为少量cache line传输 带宽利用率高 ✅ 糟糕：Thread i 访问地址 base + i * 1024 每个线程触发独立cache line 带宽利用率低 ❌

简单说就是：相邻线程访问相邻地址，硬件可以把多个请求合并，带宽利用率最高。

这是CUDA性能优化的第一课——写kernel的时候，一定要注意访存模式。

5.2 Bank Conflict（存储体冲突）

共享内存按4字节划分为32个bank（存储体），bank 0存地址0,32,64…，bank 1存地址1,33,65…，以此类推。

同一个Warp（32个线程）内，如果多个线程同时访问同一个bank的不同地址，就会发生bank conflict——访问被串行化，性能下降。

怎么避免？

• 加padding（填充）：比如共享内存数组多开一列，让跨一行的访问错开bank
• 调整访问模式：让同一个Warp内的线程访问不同的bank

💡 举个例子：共享内存矩阵转置时，按行写入后按列读取就会有bank conflict——因为同一列的元素在同一个bank里。加一列padding就能解决。

5.3 Occupancy（占用率）

Occupancy是什么？

Occupancy = 每个SM上实际驻留的Warp数 / 最大可驻留Warp数

可以理解为SM的"饱满程度"——越高说明SM越"忙"。

影响Occupancy的因素：

• 每个线程的寄存器数量（用得多，驻留的Warp就少）
• 每个Block的共享内存大小（用得多，驻留的Block就少）
• 每个Block的线程数
• 最大并发Block数限制

Occupancy越高越好吗？不一定。

• 高Occupancy有助于隐藏内存延迟——当一些Warp在等内存的时候，其他Warp可以继续计算
• 但如果kernel已经是compute-bound（计算瓶颈），再提高Occupancy收益不大
• 而且，为了提高Occupancy而减少寄存器使用，可能反而让计算变慢

结论：Occupancy是重要指标，但不是越高越好，要结合具体kernel分析。

---

六、Kernel Launch：怎么启动一个kernel

6.1 启动语法

kernel<<<grid_dim,block_dim,shared_mem_bytes,stream>>>(args...);

四个参数：

1. grid_dim：Grid的维度（有多少个Block）
2. block_dim：Block的维度（每个Block有多少线程）
3. shared_mem_bytes：动态共享内存大小（可选，静态共享内存用__shared__声明）
4. stream：CUDA Stream（可选，默认0）

6.2 配置示例

一维配置（最常用）：

// N个元素，每个block 256线程add_kernel<<<(N+255)/256,256>>>(d_a,d_b,d_c,N);

为什么是(N + 255) / 256？因为要向上取整——比如N=257，就需要2个block。

二维配置（图像处理等）：

dim3block(16,16);// 每个block 16×16 = 256线程dim3grid((W+15)/16,(H+15)/16);img_kernel<<<grid,block>>>(d_img,W,H);

---

七、常见误区

误区一：线程数越多越快

线程太多 → 寄存器压力大 → Occupancy下降 → 性能反而下降。

不是线程越多越好，要根据资源限制找到"甜点"。

误区二：共享内存一定比全局内存快

共享内存快的前提是：bank conflict少、数据复用率高。

如果每个线程只读一次数据，而且访问是合并的——直接读全局内存可能更快，省了"搬到共享内存"这一步。

误区三：分支发散无所谓

Warp内分支发散会导致串行执行，性能可能直接腰斩。写kernel时要注意避免。

---

八、总结一下

今天这篇文章，我们讲了CUDA编程体系的基础：

1. 线程组织模型：Grid/Block/Thread三级结构，线程索引计算是kernel正确性的基础
2. 内存模型：全局内存、常量内存、共享内存、寄存器——从大到小，从慢到快
3. Warp执行模型：Warp是最小调度单位，32个线程lock-step执行，分支发散会损害性能
4. 三大优化概念：合并访问（全局内存）、避免Bank Conflict（共享内存）、合理Occupancy
5. Kernel Launch：grid、block、共享内存、stream四个配置参数

下一篇预告：我们来聊CUDA性能优化的基础技术——内存访问优化、Warp Shuffle、Occupancy调优、同步与原子操作。掌握了这些，你才能写出"跑得动"的kernel。

---

📚本系列导航：

• 第1篇：CUDA编程体系全景：线程模型、内存层次与Warp执行机制（本文）
• 第2篇：CUDA性能优化基础：内存访问、Warp Shuffle与Occupancy
• 第3篇：经典算子实战：Reduce归约——从朴素到高性能
• 第4篇：经典算子实战：GEMM矩阵乘法——从Tiling到Tensor Core
• 更多内容持续更新中……