4. 内核调度机制：GPU内核调度器工作原理、并发内核执行、流（Stream）与事件（Event）的使用
4.1 GPU内核调度器的工作原理
GPU的硬件调度单元是GigaThread Engine（在NVIDIA架构中）或等效的硬件队列管理器（在AMD CDNA架构中）。该单元负责将用户提交的内核启动命令分发到各个GPU工作分发器（Work Distributor），后者再将线程块（Thread Block / CUs）调度到具体的SM（Streaming Multiprocessor）或CU（Compute Unit）上执行。
理解内核调度器的关键在于区分两个层级：
- 命令级调度（Command-level Scheduling）：由GPU前端硬件完成，决定下一个执行哪个内核、哪个复制操作。它遵循FIFO（先进先出）为主、但允许流内乱序的规则。
- 线程块级调度（Thread Block-level Scheduling）：由工作分发器完成，将内核的线程块分配到空闲的SM上。只要SM有足够的资源（寄存器、共享内存、线程槽位），线程块就会被立即调度，无需等待整个内核的所有块就绪。
关键认知：GPU不是“一个内核执行完才执行下一个”。只要资源允许，多个内核可以并发驻留在同一GPU上。调度器以线程块为粒度进行资源分配，而非以整个内核为粒度。
上图展示了典型的多内核并发场景：内核A、B、D、E以及一个内存复制操作C同时驻留在不同SM上。调度器根据每个SM的资源余量动态分配线程块，而非按启动顺序机械执行。
4.2 并发内核执行的条件与限制
要实现真正的内核并发执行，必须同时满足以下条件：
- 硬件支持：GPU必须支持并发内核执行（Compute Capability 3.5+的NVIDIA GPU均支持）。
- 资源充足：两个内核的寄存器占用、共享内存占用、线程块数量之和不能超过SM的物理上限。
- 无数据依赖：内核之间不能存在未同步的数据依赖，否则必须使用事件或流来保证顺序。
- 使用不同流（Stream）：默认流（Stream 0）上的内核是严格串行的，必须使用多个非默认流才能触发并发。
| 资源类型 | 典型上限（以A100为例） | 并发影响 |
|---|---|---|
| 线程块/SM | 32 | 超过则等待 |
| 线程/SM | 2048 | 决定最大驻留块数 |
| 寄存器/SM | 65536 | 寄存器用量大的内核并发度低 |
| 共享内存/SM | 164 KB | 共享内存大的内核独占SM |
4.3 流（Stream）与事件（Event）的使用方法
流是GPU操作（内核启动、内存复制）的有序序列。同一流内的操作按提交顺序执行；不同流之间的操作可以并发。CUDA中创建流的典型代码：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在流1中启动内核A
kernel_A<<<grid, block, 0, stream1>>>(args...);
// 在流2中启动内核B（可与A并发）
kernel_B<<<grid, block, 0, stream2>>>(args...);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);事件用于跨流同步或测量时间。典型用法是让流2等待流1中的某个事件完成：
cudaEvent_t event;
cudaEventCreate(&event);
// 流1中记录事件
kernel_A<<<grid, block, 0, stream1>>>(args...);
cudaEventRecord(event, stream1);
// 流2等待事件后启动内核B
cudaStreamWaitEvent(stream2, event, 0);
kernel_B<<<grid, block, 0, stream2>>>(args...);
cudaEventDestroy(event);最佳实践：使用优先级流（cudaStreamCreateWithPriority）可以让高优先级流中的内核抢占低优先级流的资源。对于延迟敏感的小内核，建议使用高优先级流；对于吞吐型大内核，使用普通优先级即可。
4.4 案例：多流并发矩阵乘法
假设需要计算 C = A × B，其中A、B为大型矩阵。将矩阵按行划分为4个分块，每个分块在不同流上独立计算：
const int NUM_STREAMS = 4;
cudaStream_t streams[NUM_STREAMS];
for (int i = 0; i < NUM_STREAMS; i++) {
cudaStreamCreate(&streams[i]);
}
// 每个流处理矩阵的一个行分块
for (int i = 0; i < NUM_STREAMS; i++) {
int offset = i * rows_per_stream;
matmul_kernel<<<grid, block, 0, streams[i]>>>(
d_A + offset, d_B, d_C + offset, rows_per_stream);
}
// 等待所有流完成
for (int i = 0; i < NUM_STREAMS; i++) {
cudaStreamSynchronize(streams[i]);
cudaStreamDestroy(streams[i]);
}实测在A100上，4流并发比单流执行快约2.8倍（受限于内存带宽和SM占用率）。若分块过小（如8个流），调度开销反而增加，性能下降。
4.5 常见误区与性能复盘
误区1：认为默认流与其他流可以并发。实际上，默认流（Stream 0）与所有其他流之间存在隐式同步。若内核A在默认流，内核B在自定义流，B必须等待A完成。解决方法是使用非阻塞默认流（cudaStreamNonBlocking标志）。
误区2：忽略资源占用对并发度的影响。一个内核若使用大量寄存器（如每个线程128个寄存器），则每个SM只能驻留少量线程块，并发内核可能因资源不足而串行执行。使用cudaOccupancyMaxPotentialBlockSize可辅助计算最佳块大小。
性能复盘清单：
- 使用Nsight Systems查看内核时间线，确认并发是否真正发生（多个内核时间轴重叠）。
- 检查GPU利用率是否达到90%以上；若低于70%，说明并发度不足。
- 对比单流与多流的执行时间，若提升小于20%，检查是否受内存带宽或PCIe传输瓶颈限制。
- 使用cudaEventElapsedTime精确测量各流耗时，定位负载不均问题。
- 对于短内核（执行时间<5μs），流切换开销可能超过收益，建议合并为单个内核。
通过合理使用流与事件，并理解底层调度器的资源分配逻辑，多内核GPU的吞吐量可以获得显著提升。关键在于平衡并发粒度与资源占用，并通过性能分析工具验证实际并发效果。