CUDA 编程的基本原理是什么？怎么入门？

一、为什么需要先理解硬件

很多初学者直接上手写 kernel，能跑通，却说不清为什么快、为什么慢。CUDA 编程的本质是”用软件结构去匹配硬件结构”，如果对底层执行单元的布局没有概念，优化就无从谈起。本文以易懂性优先，在个别细节上做适度简化，先把最核心的几张”地图”讲清楚。

二、GPU 的宏观结构：多核阵列 + 分级存储

GPU 在抽象层面是一个大规模多核结构，可以类比 CPU 的多线程编程模型，但它有几个鲜明的新特点，而这些特点恰恰是后续学习的重点。

整体来看，一颗现代 GPU 由若干个 SM（Streaming Multiprocessor，流式多处理器） 组成，不同规格的芯片 SM 数量不同，从上百个到接近一百五十个不等。可以把一个 SM 粗略理解成 CPU 里”一组核”的集合体。

所有 SM 共享两条数据通路：

• L2 Cache（SRAM）：容量较小、带宽高、延迟低；
• DRAM（Global Memory，全局显存）：容量大、带宽相对受限、延迟高。

这一层级关系与 CPU 的缓存体系高度相似，区别在于：CPU 上缓存命中率通常由硬件和编译器自动优化，程序员较少直接操心；而在 GPU 上，提升 L2 命中率、减少对全局显存的访问，是 kernel 优化的核心手段之一。原因很直接——GPU 上并行线程数量巨大，一旦大量线程同时涌向全局显存，带宽就会成为瓶颈。

三、SM 内部：两类计算资源

拆开一个 SM，可以看到它被划分为若干子部分，每个子部分配备两类计算资源：

1. 矩阵运算单元（Tensor Core）

这是专为矩阵乘法设计的高吞吐计算单元，而矩阵乘法正是神经网络训练与推理中最主要的计算形态。它的设计思想可以追溯到脉动阵列（systolic array）这类经典的电路结构——通过让数据在计算单元阵列中有节奏地流动，减少重复的数据搬运，从而大幅提升单位能耗下的算力。学有余力的读者建议扩展阅读这一方向，它有助于理解”为什么矩阵乘能做得这么快”。

2. 通用运算单元（CUDA Core）

这类单元统称 CUDA Core，内部又按数据类型分为 INT32、FP32、FP64 等 ALU（算术逻辑单元）。每个 ALU 负责执行一次对应的标量算术运算。它们承担的是类似 CPU 所能做的通用算术与逻辑工作：地址计算、循环控制、标量运算、非线性函数等。

一句话总结：每个 SM 子部分同时拥有”主打矩阵乘法”和”主打通用标量运算”两类资源，前者服务于神经网络这类计算密集型负载，后者保证程序的通用性。理解这个分工，才能明白为什么同一份代码里，矩阵乘部分和逻辑控制部分的性能表现差异会那么大。

四、Warp Scheduler 与 warp：被很多入门课忽略的关键

SM 内部还有若干 Warp Scheduler（warp 调度器）。这是一个极其重要的概念，但不少入门课程并未涉及，导致学习者后期遇到性能问题时一头雾水。

warp 是什么？ 它类似 CPU 中的”一组线程”，但在 GPU 中，线程总数极其庞大（动辄数十万），为了让调度和硬件实现更高效，设计上要求线程具有较高的”同质性”——即一组线程在同一时刻执行同一条指令。

Warp Scheduler 的管理粒度就是 32 个线程为一组，这个组就叫一个 warp。这种”单指令、多数据”的结构统称为 SIMD（Single Instruction, Multiple Data），是一种经典的并行计算架构。

由此可以推导出几条对编程直接有用的结论：

• 分支发散是有代价的。如果一个 warp 内的 32 个线程因条件判断走向不同分支，硬件只能让不同分支串行执行，其他线程空等，效率成比例下降。因此应尽量让同一 warp 内的线程走相同路径。
• 线程数应按 32 的整数倍来规划。比如 block 大小取 128、256、512，而不是 100、250，否则最后一个不完整的 warp 会浪费执行资源。
• warp 是调度的基本单位，不是线程。谈论占用率、延迟隐藏时，实际上谈论的是”有多少个 warp 可以随时被切换进来执行”。

Warp Scheduler 的职责就是在多个 warp 之间快速切换：当某个 warp 因等待访存而停顿时，立刻切换到另一个就绪的 warp 继续发射指令，从而把访存延迟”藏”起来。这正是 GPU 能用相对简单的核换取高吞吐的根本机制。

五、软件如何映射到硬件：线程层次

CUDA 用三层抽象的线程组织来对应硬件：

• Thread（线程）：最小执行单元；
• Block（线程块）：一组线程，可进行块内同步与共享内存通信，被分配到一个 SM 上执行；
• Grid（网格）：一次 kernel 启动对应的全部线程块集合。

映射关系可以概括为：一个 Grid 的各个 Block 被分发到多个 SM 上；一个 Block 内的线程以 warp 为单位被 Warp Scheduler 调度执行。

一个最简的 kernel 示例，展示索引计算与边界保护这两个必备要素：

__global__ void saxpy(float* y, const float* x, float a, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {                 // 边界保护，避免越界访存
        y[idx] = a * x[idx] + y[idx];
    }
}

启动方式：

int threads = 256;                          // 取 32 的整数倍
int blocks  = (n + threads - 1) / threads;  // 向上取整覆盖全部元素
saxpy<<<blocks, threads>>>(d_y, d_x, 2.0f, n);

六、内存层级：优化的主战场

从快到慢、从小到大，大致是：

1. 寄存器：线程私有，最快，数量有限；
2. 共享内存（Shared Memory）：Block 内共享，可显式管理，是块内数据复用的关键；
3. L2 Cache：全 GPU 共享；
4. 全局显存：容量最大、延迟最高。

对应三条最实用的优化原则：

• 提高访存合并度：相邻线程访问相邻地址，使一次内存事务服务尽可能多的线程；
• 用共享内存做块内复用：典型如分块矩阵乘法，把子块搬到共享内存后反复使用，避免重复读全局显存；
• 提高 L2 命中率：通过合理安排数据布局与访问顺序，让热点数据留在 L2 中。

七、入门路径建议

1. 搭建环境：安装工具链，跑通官方示例，确认设备与驱动版本匹配；
2. 写第一个 kernel：从向量加法、SAXPY 这类逐元素运算开始，掌握索引与边界保护；
3. 理解线程层次：故意调整 block 大小，观察 32 的非整数倍带来的性能差异；
4. 学习同步与共享内存：实现分块矩阵乘法，对比优化前后的耗时；
5. 学会测量：用性能分析工具定位瓶颈，区分是计算受限、带宽受限还是延迟未被隐藏；
6. 进阶方向：矩阵运算单元的使用、异步拷贝与流水线、多流并发。

八、小结

CUDA 编程的基本原理可以压缩成一句话：用 Grid-Block-Thread 的软件层次去映射 SM-warp-线程的硬件层次，用 SIMD 的执行模型换取吞吐，用分级存储的合理利用换取带宽。

抓住三个关键词——SM 的分工、warp 的粒度、内存的层级——大部分性能问题的成因都能被定位到其中之一。入门阶段不必追求记住所有细节，先把这套”硬件地图”建立起来，后续的优化技巧自然会各归其位。