一、并发调度基础：GPU架构概览、SM与CUDA Core、并发执行模型、Stream与Event概念
1.1 GPU架构概览——别把它当CPU用
很多人刚开始学GPU编程，总习惯用CPU的思维去理解GPU。我当年也犯过这个错。其实GPU和CPU的设计哲学完全不同。
CPU是「低延迟」设计，它要快速响应各种任务。GPU是「高吞吐」设计，它不在乎单个任务跑多快，只在乎单位时间内能处理多少任务。说白了，CPU是法拉利，GPU是公交车。法拉利能拉两个人跑得飞快，公交车能拉一车人慢慢走——但论总运力，公交车完胜。
GPU架构的核心思想就是：用大量的计算单元，去掩盖访存延迟。你想想看，当一个线程在等数据时，GPU不会闲着，它会立刻切换到另一个线程去算。这就是所谓的「零开销上下文切换」。
核心要点：GPU通过「大量线程 + 快速切换」来隐藏延迟，而不是像CPU那样靠大缓存和分支预测。
1.2 SM与CUDA Core——真正的执行单元
先说说SM（Streaming Multiprocessor，流式多处理器）。它是GPU内部的核心计算单元。你可以把SM想象成一个「小型的计算中心」。每个SM里包含多个CUDA Core、共享内存、寄存器文件、调度器等等。
那CUDA Core又是什么？它其实就是最基础的算术逻辑单元（ALU），负责执行浮点和整数运算。嗯，这里要注意：CUDA Core ≠ CPU Core。CPU Core是完整的处理器核心，能独立执行指令。CUDA Core只是个运算单元，它必须依赖SM里的调度器来分配任务。
我记得有一次做性能调优，发现某个kernel的occupancy（占用率）一直上不去。排查了半天，发现是每个线程用的寄存器太多了，导致SM能同时容纳的线程数大幅下降。这就是典型的「资源限制」问题。
| 组件 | 作用 | 个人经验 |
|---|---|---|
| SM | 管理线程块、调度warp | SM数量决定了并行度上限 |
| CUDA Core | 执行算术运算 | 别只看数量，要看利用率 |
| Warp Scheduler | 每周期选择一个warp发射指令 | warp divergence是性能杀手 |
| Shared Memory | 线程块内共享数据 | 用好了能省大量全局内存访问 |
避坑指南：我曾经在项目里盲目追求CUDA Core数量，结果发现大部分Core都在空转。后来才明白，利用率比数量更重要。建议用nvidia-smi或Nsight Compute先看看实际占用率。
1.3 并发执行模型——GPU到底怎么「同时」做事？
GPU的并发执行，其实分三个层次：
- 线程级并发：同一个warp里的32个线程，执行同一条指令（SIMT模型）。
- 线程块级并发：多个线程块可以分配到不同的SM上执行。
- kernel级并发：多个kernel可以同时运行在同一个GPU上（需要支持）。
为什么会设计成这种层次结构？说白了，就是为了「可扩展性」。你写代码时不用关心GPU有几个SM，CUDA运行时会自动把线程块分配到可用的SM上。我刚开始做CUDA开发时，总觉得这层抽象太「黑盒」，后来发现这恰恰是GPU能自动适配不同硬件的关键。
这里有个重要的概念叫occupancy（占用率）。它表示每个SM上活跃的warp数量与最大warp数量的比值。occupancy越高，SM越不容易因为等待数据而闲置。
注意：occupancy不是越高越好。我见过有人为了追求100% occupancy，把每个线程的工作量压得特别小，结果调度开销反而超过了计算收益。一般来说，50%-75%的occupancy配合合理的计算量，效果往往更好。
1.4 Stream与Event——并发调度的核心工具
Stream（流）是CUDA里实现并发执行的核心机制。你可以把Stream想象成一条「任务流水线」。默认情况下，所有操作都在默认流（default stream）里执行，它们是串行的。但如果你创建多个Stream，就可以让不同的操作并行执行。
举个例子：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 两个kernel可以同时执行
kernel1<<<grid, block, 0, stream1>>>(data1);
kernel2<<<grid, block, 0, stream2>>>(data2);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);
Event（事件）则用来做同步和计时。你可以往Stream里插入Event，然后查询Event是否完成，或者让某个Stream等待另一个Stream的Event。
我个人习惯用Event来测量kernel的执行时间，比用CPU计时器准得多：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream);
myKernel<<<grid, block, 0, stream>>>();
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);
printf("Kernel用时: %f ms\n", ms);
实用技巧：我曾经在做一个实时推理系统时，用Stream实现了「计算-传输」重叠。一个Stream负责从CPU拷贝数据到GPU，另一个Stream负责执行推理kernel，第三个Stream负责把结果拷回来。三个Stream流水线作业，吞吐量直接翻倍。
1.5 知识体系总览
下面这张图，是我梳理的本章核心逻辑。你可以把它当作一个「思维导图」来用：
这张图把本章的四个核心知识点串起来了。GPU架构是基础，SM和CUDA Core是执行单元，并发执行模型是运行机制，Stream和Event是调度工具。四者缺一不可。
总结一下：理解GPU并发调度，关键要抓住「SM是调度单位，warp是执行单位，Stream是并发单位」这条主线。我见过太多人一上来就调kernel参数，结果连occupancy和stream都没搞明白，走了不少弯路。
想要更多资料，联系可可