📚 课程目录第2章 CUDA基础回顾

2. CUDA基础回顾

在深入延迟分析之前，必须先建立对CUDA执行模型的准确认知。延迟（Latency）的根源在于硬件资源的调度方式与访存路径，而这些都直接由GPU架构、线程组织与内存层次决定。本节将系统梳理这四个层面的核心概念，为后续的延迟量化与优化打下基础。

一、GPU架构核心：SM与SIMT执行模型

GPU由多个流式多处理器（Streaming Multiprocessor, SM）构成，每个SM内部包含若干CUDA核心（FP32单元）、特殊函数单元（SFU）、加载/存储单元（LSU）以及共享内存和寄存器堆。以NVIDIA Ampere架构的GA102为例，一个SM包含128个CUDA核心，4个处理块（Processing Block），每个处理块拥有32个CUDA核心——恰好对应一个Warp的大小。

GPU执行的基本调度单位是Warp（32个线程）。硬件以Warp为单位进行指令发射：同一Warp内的32个线程在同一个指令周期执行同一条指令，这就是SIMT（单指令多线程）模型。若线程分支发散，则不同路径会被串行执行，造成指令吞吐下降——这是延迟分析中必须关注的第一类开销。

  关键理解：延迟的硬件基础是“线程→Warp→SM”的映射关系。一个线程从发出访存请求到数据返回，中间经历的时钟周期数即为访存延迟。GPU通过大量Warp的并行切换来隐藏这种延迟，而非像CPU那样依赖多级缓存和乱序执行。

二、线程层次与索引映射

CUDA将线程组织为三层结构：Grid（网格）→ Block（线程块）→ Thread（线程）。内核启动时通过执行配置指定网格与线程块的维度：

// 内核启动配置
dim3 blockDim(256);          // 每个Block含256个线程
dim3 gridDim((N + 255) / 256); // 覆盖N个数据元素所需的Block数
kernel<<<gridDim, blockDim>>>(data, N);

在设备端，每个线程通过内置变量获取自身全局索引：

int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N) { /* 执行计算 */ }

这种层次设计直接影响延迟：Block是SM上调度的基本单位，一个Block内的线程共享SM资源；而Warp是硬件执行的基本单位。若Block大小不是32的整数倍，会产生尾部效应——最后一个Warp中部分线程空闲，浪费执行槽位。

三、内存层次与访存延迟

CUDA内存模型是延迟分析的核心。下表总结了各内存类型的特性：

内存类型	作用域	生命周期	典型延迟	带宽
寄存器	单个线程	内核执行期间	≈0周期	极高
共享内存	Block内所有线程	内核执行期间	20-30周期	高（无广播冲突时）
L2缓存	整个GPU	内核执行期间	200-300周期	中高
全局内存	所有线程+主机	程序运行期间	400-800周期	受限于DRAM带宽
常量内存	所有线程（只读）	程序运行期间	命中时≈共享内存	广播优化

全局内存访问是延迟的主要来源。当Warp内32个线程访问连续地址时，硬件将其合并为一次或少数几次事务（Memory Transaction），这是实现高带宽的关键。若访问地址分散（如按列访问二维数组），则产生多次事务，有效带宽骤降——这就是访存合并（Memory Coalescing）问题。

四、CUDA编程模型与延迟隐藏

CUDA编程模型遵循异构计算范式：主机（CPU）负责串行逻辑与数据准备，设备（GPU）负责大规模并行计算。典型流程为：

分配设备内存（cudaMalloc）
将数据从主机拷贝到设备（cudaMemcpy，H2D）
启动内核（kernel<<<grid, block>>>）
将结果拷回主机（cudaMemcpy，D2H）
释放设备内存（cudaFree）

延迟隐藏的机制在于：当一个Warp因访存而阻塞时，SM调度器立即切换到另一个就绪（Ready）的Warp执行。要实现高效隐藏，需要满足两个条件：

足够的Warp数量：每个SM至少需要6-8个可切换Warp才能覆盖400-800周期的访存延迟。
高占用率（Occupancy）：占用率 = 活跃Warp数 / 最大Warp数。受限于寄存器与共享内存用量。

实践建议：在延迟分析中，首先检查cudaOccupancyMaxActiveBlocksPerMultiprocessor获取理论占用率，再结合ncu（NVIDIA Nsight Compute）的sm__warps_active.avg.pct_of_peak_sustained_active指标对比实际值。若实际占用率远低于理论值，说明存在资源限制（如寄存器溢出或共享内存超限）。

五、常见误区与复盘

误区1：认为线程越多性能越好。线程数量超过硬件可并行执行的Warp槽位后，增加线程只会增加调度开销，不会提升吞吐。应通过cudaDeviceProp查询SM数量和每SM最大线程数，合理设置Block大小（通常为128或256）。

误区2：忽略访存合并。即使计算量很小，非合并访存也会使内核变为访存密集型。例如二维数组按列访问时，应通过转置或改用共享内存进行矩阵分块（Tiling）来保证合并访问。

误区3：将共享内存当作缓存使用。共享内存需要程序员显式管理数据加载与同步（__syncthreads()）。若数据复用率低，直接使用L2缓存可能更高效，避免同步开销。

复盘案例：某内核处理1M浮点数求和，初始实现每个线程处理1个元素，Block=64，Grid=15625。实测延迟为12ms，占用率仅25%。分析发现：Block过小导致每个SM上Block数量受限，Warp切换不足。优化为Block=256，Grid=4096后，占用率提升至62%，延迟降至5ms。进一步使用向量化加载（float4）减少访存指令数，延迟降至3.2ms。最终通过多级归约（每个Block先归约到共享内存），延迟稳定在2.8ms。