Appearance
CUDA 编程与内存模型
状态：入门理解 · 更新于 2026-09-14
CUDA 优化的核心不是“把代码搬到 GPU”，而是让足够多的线程并行执行，同时减少等待数据和同步的时间。本页先建立执行层级、内存层级和最基本的性能判断框架。
执行模型
一次 Kernel 启动会创建一个 Grid，Grid 由多个 Thread Block 组成，每个 Block 再包含多个 Thread。硬件执行时，Block 被调度到 SM，线程通常以 32 个一组的 Warp 发射指令。
| 软件层级 | 作用 | 关键约束 |
|---|---|---|
| Grid | 一次 Kernel 的全部工作 | 普通 Kernel 内没有低成本的全 Grid 屏障 |
| Block | 可协作的一组线程 | 同一 Block 才能直接共享 Shared Memory |
| Warp | SM 的基本执行批次 | 分支路径不同会产生 Warp Divergence |
| Thread | 处理一个或多个数据元素 | 寄存器使用过多会限制并发驻留量 |
Block 不是越大越好。选择线程数时要同时看 Warp 数量、寄存器、Shared Memory 和 Occupancy，最终以实际测量结果为准。
内存层级
| 存储位置 | 可见范围 | 典型特点 | 常见用途 |
|---|---|---|---|
| Register | 单线程 | 片上、延迟低、容量有限 | 索引、累加器、线程私有临时值 |
| Shared Memory | 单 Block | 片上、可编程缓存，需要显式同步 | Tile、Block 内数据复用 |
| Global Memory | 整个 Device | 容量大、延迟高 | 输入、输出和长期驻留数据 |
| Constant Memory | 整个 Device，只读 | 有缓存；Warp 读取同一地址时效率高 | 小型只读参数 |
| Local Memory | 逻辑上属于单线程 | 实际位于片外显存 | 寄存器溢出、大型线程局部数组 |
“Local”不代表片上。如果编译器发生 register spill，数据会落到 Local Memory，代价接近 Global Memory。
最小 Kernel
下面的向量加法展示最常见的索引方式：
cpp
__global__ void vector_add(const float* a, const float* b, float* c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
c[i] = a[i] + b[i];
}
}
int threads = 256;
int blocks = (n + threads - 1) / threads;
vector_add<<<blocks, threads>>>(a, b, c, n);这里相邻线程访问相邻的 float，有利于 Global Memory 合并访问。启动后仍需检查 CUDA 错误；计时时还要使用 CUDA Event 或显式同步，不能只测主机端的异步提交时间。
三个主要性能问题
1. Global Memory 访问是否合并
一个 Warp 的地址如果连续且合理对齐，硬件可以用较少的内存事务完成访问。大步长、散乱或未对齐访问通常会浪费带宽。优先调整数据布局和线程到数据的映射，再考虑更复杂的技巧。
2. 数据是否值得搬进 Shared Memory
Shared Memory 适合被同一 Block 多次复用的数据。只用一次的数据先搬进去通常得不偿失。使用时还要避免 Bank Conflict，并用 __syncthreads() 保证 Block 内生产者和消费者的顺序。
3. 并发度是否被资源限制
寄存器或 Shared Memory 占用过高，会减少一个 SM 上可同时驻留的 Block 和 Warp。Occupancy 太低可能无法隐藏延迟，但 Occupancy 最高也不等于性能最好；需要结合吞吐、访存和指令瓶颈判断。
基础优化顺序
- 先做 CPU 或框架基线，固定输入 shape、dtype 和误差标准。
- 检查相邻线程的 Global Memory 访问是否连续。
- 找到可复用数据，再设计 Shared Memory Tile。
- 观察寄存器、Shared Memory 和 Occupancy 的约束。
- 用 Nsight Systems 判断时间花在哪，再用 Nsight Compute 分析目标 Kernel。
- 每次只改一个变量，同时做正确性和性能回归。
常见误区
- Kernel 启动是异步的，主机计时可能只测到提交开销。
__syncthreads()只同步一个 Block，且不能放在部分线程无法到达的分支中。- Shared Memory 更快，但搬运、同步和 Bank Conflict 都有成本。
- 小输入可能主要受 Launch Overhead 影响，优化计算主体不一定有效。
- 不同 Compute Capability 的资源限制和最佳参数可能不同。
信息来源
延伸阅读
更新日志
- 2026-09-14：建立执行模型、内存层级与基础优化流程。