3. CUDA编程模型与交付视角

在GPU加速项目中，CUDA编程模型是连接算法设计与硬件执行的桥梁。很多交付项目性能不达标、后期维护困难，根源往往不在算法本身，而在于对线程层次、内存层次和核函数执行模型的理解偏差。本节从交付视角出发，把这三个概念讲清楚，并说明它们如何直接影响性能与可维护性。

线程层次：从网格到线程的三级结构

CUDA把并行执行体组织成三级结构：网格（Grid）→ 线程块（Block）→ 线程（Thread）。一次核函数调用对应一个网格，网格由若干线程块组成，每个线程块又包含若干线程。线程通过内置变量定位自己：blockIdx标识所在块，threadIdx标识块内线程，blockDim给出块内线程数，gridDim给出网格内块数。

全局索引的常用写法是：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) { /* 处理第 idx 个元素 */ }

这里有一个交付中极易被忽视的约定：线程块内的线程可以同步（__syncthreads()），但不同线程块之间无法在核函数内直接同步。这意味着跨块的数据依赖必须通过多次核函数启动或原子操作来协调。如果算法设计时假设“所有线程一起完成某一步再进入下一步”，就必须把它拆成多个核函数，否则会出现难以复现的竞态错误。

线程块大小（blockDim）的选择也不是随意的。它应当是32的倍数，因为硬件以线程束（Warp）为调度单位，一个Warp固定32个线程。若块大小不是32的倍数，最后一个Warp会有空闲线程，浪费执行资源。常见取值是128、256、512。块太小则并行度不足、调度开销占比高；块太大则寄存器压力上升、占用率下降。交付时建议把块大小做成可配置参数，便于在不同GPU上调优。

内存层次：决定性能的关键

CUDA的内存层次从快到慢依次是：寄存器 → 共享内存/一级缓存 → 二级缓存 → 全局内存。核函数中每个线程的局部变量通常放在寄存器；共享内存由线程块内所有线程共享，延迟远低于全局内存；全局内存容量最大但延迟最高，是性能瓶颈的主要来源。

交付视角下，内存层次的核心结论是：能放共享内存的复用数据，不要反复读全局内存。以矩阵乘法为例，若每个线程都直接从全局内存读取A、B的元素，同一数据会被大量重复读取。采用分块（Tiling）策略，把子矩阵先加载到共享内存，块内线程复用后再计算，可显著减少全局内存访问次数。这是CUDA优化中最经典也最有效的手段之一。

另一个必须掌握的概念是合并访问（Coalesced Access）。当同一个Warp的32个线程访问全局内存中连续且对齐的地址时，硬件可以把这些访问合并成少量内存事务；若地址分散或跨步，则事务数成倍增加，带宽利用率骤降。因此数据布局应尽量让相邻线程访问相邻元素。若算法天然需要按列访问，可考虑先转置或改用共享内存中转。

核函数执行模型：异步与错误处理

核函数启动是异步的：主机端调用后立即返回，GPU在设备端排队执行。这意味着主机端代码不能假设核函数已经完成。需要结果时必须显式同步，例如cudaDeviceSynchronize()或通过cudaMemcpy的同步语义。交付中常见的“结果偶尔不对”往往就是漏了同步，主机在核函数还没算完时就去读数据。

核函数没有返回值，错误必须通过cudaGetLastError()或cudaPeekAtLastError()检查。启动失败（如块大小超限、共享内存超限）不会抛异常，只会静默设置错误码。建议在每次核函数启动后立即检查，否则错误会延迟到后续某个API调用才暴露，定位困难。

下面这张图把线程层次、内存层次与执行模型的关键关系放在一起，便于对照理解。

对性能与可维护性的实际影响

把上述三点落到交付项目上，可以形成几条明确的工程准则。

性能方面：块大小、共享内存使用、访问模式是三个最直接的调优杠杆。一个典型的向量加法核函数，若块大小设为256、每个线程处理一个元素，在主流GPU上通常能接近带宽上限；但若改成每个线程处理多个元素（网格跨步循环），可以减少块数量、提升指令级并行，往往还能再提升一截。交付时建议先用简单版本跑通正确性，再逐步引入分块、跨步、共享内存等优化，每步都用性能计数器（如Nsight Compute）确认瓶颈是否真的被消除，而不是凭感觉堆优化。

可维护性方面：线程层次和内存层次一旦在代码中硬编码，后续换GPU型号或调整问题规模时就要大改。建议把块大小、网格大小、共享内存尺寸抽成常量或配置项，并在核函数中统一用blockIdx * blockDim + threadIdx的索引模式，避免不同核函数各写一套。错误检查也应封装成宏或辅助函数，保证每次启动后都检查，而不是散落在各处、时有时无。

理解线程层次、内存层次和核函数执行模型，本质上是在理解“GPU如何组织并行、数据如何流动、主机与设备如何协作”。这三者决定了代码能不能跑对、跑得快、以及后续能不能改。在交付项目中，把它们作为设计评审和代码审查的固定检查项，比事后调优更省成本。