GPU架构基础：现代GPU硬件架构与多内核并行执行原理
2.1 从“核数”迷思到真实硬件视图
在讨论多内核GPU工作负载分配之前，必须先建立正确的硬件抽象。许多开发者将GPU视为“拥有数千个核心的处理器”，这种说法虽然营销上正确，却容易导致性能调优方向的根本性误判。现代NVIDIA GPU（如A100、H100、RTX 40系列）的真实组织方式是：GPU由多个流式多处理器（SM）构成，每个SM内部包含若干CUDA核心、共享内存、寄存器堆、调度器与特殊功能单元。CUDA核心只是SM内部的整数/浮点执行单元，它本身不具备独立的取指或调度能力。
以A100为例，其GA100芯片包含108个SM，每个SM拥有64个FP32 CUDA核心、32个FP64核心、4个第三代Tensor Core，以及256KB的寄存器文件（可拆分为64K个32位寄存器）。因此，从线程调度的视角看，真正决定并行执行能力的是SM数量与每个SM的硬件线程槽位（thread slots），而非CUDA核心总数。一个线程块（block）被调度到某个SM后，其所有线程共享该SM的寄存器与共享内存资源，直到整个块执行完毕才会释放。
关键认知：CUDA核心是“执行单元”，SM是“调度与资源容器”。多内核并行时，内核之间竞争的是SM上的执行资源（线程槽、寄存器、共享内存），而不是抽象的“核心数”。
2.2 内存层次：带宽与延迟的权衡
GPU内存层次的设计目标非常明确：用高带宽隐藏延迟，用局部性减少片外访问。从离SM最近到最远，依次是：寄存器文件、共享内存/一级缓存（L1）、二级缓存（L2）、显存（HBM）。每个层级在容量、带宽、延迟和可见性上都有本质差异。
| 内存层级 | 容量（A100典型值） | 带宽/延迟 | 作用域 | 多内核影响 |
|---|---|---|---|---|
| 寄存器文件 | 256KB / SM | ~19 TB/s（合计），<1周期 | 单个线程私有 | 寄存器溢出会强制访问局部内存 |
| 共享内存 / L1 | 192KB / SM（可配置） | ~33 TB/s（合计），~30周期 | 线程块内共享 | 块间隔离，但容量竞争 |
| L2 缓存 | 40MB（全芯片） | ~5 TB/s，~200周期 | 芯片内所有SM共享 | 内核间数据复用与污染 |
| HBM 显存 | 40~80GB | 1.6~3.35 TB/s，~500周期 | 主机与GPU共享 | PCIe/片外带宽瓶颈 |
在多内核并行场景中，L2缓存是内核间相互影响最关键的层级。当两个内核交替执行时，如果它们访问的数据集都大于L2容量，就会互相驱逐对方的缓存行，导致缓存抖动（cache thrashing）。反之，如果内核间存在数据复用（例如内核A的产出恰好是内核B的输入），则可以通过L2缓存实现“近内存计算”，避免数据落回显存再读出的往返开销。
2.3 多内核并行执行原理：从硬件调度到资源竞争
现代GPU支持多内核并发执行（Concurrent Kernel Execution），其底层机制是：GPU硬件工作调度器（GigaThread Engine）将不同内核的线程块分发到不同的SM上。只要SM还有空闲的资源槽位（线程槽、寄存器、共享内存），就可以接纳来自另一个内核的线程块。这意味着多个内核可以在同一时刻、不同SM上真正并行执行，而非时间片轮转。
内核并行的前提条件包括：
- 资源充足性：每个内核的线程块所需资源（寄存器数×线程数、共享内存字节数）必须小于SM剩余可用资源。
- 独立执行流：内核之间不能有未同步的数据依赖（除非通过流事件或全局内存栅栏显式同步）。
- 流（Stream）机制：不同流中的内核才有并发可能；同一流中的内核默认串行执行。
从SM内部视角看，每个SM包含4个硬件调度器（processing block），每个调度器管理一组CUDA核心。线程以warp（32线程）为单位被调度。当一个warp因访存或同步而阻塞时，调度器会立即切换到另一个就绪的warp，这种零开销上下文切换是GPU隐藏延迟的核心手段。多内核并行时，不同内核的warp会交错占用同一SM的调度器，因此warp级调度策略直接影响内核间的公平性与吞吐量。
2.4 案例：双内核并发执行与资源分配
假设GPU有20个SM，每个SM最多2048个线程槽。现在提交两个内核：
- 内核A：gridDim=(40,1)，blockDim=(256,1)，每线程使用32个寄存器，无共享内存。
- 内核B：gridDim=(20,1)，blockDim=(512,1)，每线程使用64个寄存器，每块使用8KB共享内存。
计算资源占用：
内核A每块资源：
线程数 = 256
寄存器 = 256 × 32 = 8192 个（占SM寄存器堆 8192/65536 = 12.5%）
共享内存 = 0 KB
每SM可容纳块数 = min(2048/256, 65536/8192, 无共享内存限制) = min(8, 8) = 8 块
内核A总块数40，需要至少 40/8 = 5 个SM
内核B每块资源：
线程数 = 512
寄存器 = 512 × 64 = 32768 个（占50%）
共享内存 = 8KB（占192KB的4.2%）
每SM可容纳块数 = min(2048/512, 65536/32768, 192KB/8KB) = min(4, 2, 24) = 2 块
内核B总块数20，需要至少 20/2 = 10 个SM如果两个内核放入不同流并同时启动，硬件调度器会尝试将内核A的40个块和内核B的20个块分配到20个SM上。理想情况下，内核A占用5个SM（每SM 8块），内核B占用10个SM（每SM 2块），剩余5个SM空闲。但实际调度可能更分散——例如内核A的块散布在15个SM上，每个SM只分到2~3块，导致剩余资源碎片化，无法容纳内核B的块。此时内核B只能等待部分SM释放资源。
调优启示：使用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 查询每个内核的占用率，再结合 cudaStreamCreate 创建多个流，通过 cudaEventQuery 检测内核是否真正并发执行。若发现并发度不足，可尝试减小blockDim或调整寄存器使用量（如 __launch_bounds__）来降低每块资源占用。
2.5 常见误区与复盘
误区一：认为“内核并行 = 所有SM同时执行所有内核”。实际上，硬件调度以线程块为粒度，一个SM在某一时刻可能只执行一个内核的块，也可能交错执行多个内核的块。并行程度完全取决于资源余量。
误区二：忽略L2缓存竞争。两个内核即使在不同SM上并行，它们仍然共享L2缓存和显存带宽。如果内核A是访存密集型（如矩阵转置），内核B是计算密集型（如FFT），两者并行时内核A可能占满L2带宽，导致内核B的访存延迟显著上升。此时需要借助计算能力（compute capability）与内存优先级，或使用 cudaCtxResetPersistingL2Cache 等API管理L2持久化窗口。
误区三：将“流”等同于“并发保证”。流只是提交顺序的容器，硬件是否真正并发取决于资源。若内核A的块占用了所有SM的线程槽，即使内核B在不同流中，也必须等待。正确做法是先用占用率API计算理论并发可能性，再通过事件回调或cudaStreamQuery验证实际行为。
复盘案例：某图像处理管线中，开发者将“高斯模糊”和“边缘检测”放入两个流期望并行。实测发现性能反而下降15%。分析后发现：两个内核都使用256线程块、每线程64寄存器，导致每个SM最多容纳4块（2048/256=8，但寄存器限制65536/(256×64)=4）。两个内核合计需要40个SM，而GPU只有20个SM，因此它们实际是串行执行的。同时，由于两个内核都访问同一张超大图像（超过L2容量），L2命中率从85%降至40%，加剧了显存带宽压力。解决方案是将两个内核合并为一个内核（融合），或降低寄存器使用量使每SM容纳更多块，从而让两个内核真正并发。