GPU 架构与指令集

warp 是 NVIDIA GPU 中组织线程执行的单位。CUDA Core 是算术执行资源。调度器按可执行状态发射 warp 的指令，算术、访存和矩阵指令走各自的数据通路。

线程保留自己的寄存器和执行状态，但不会在整个 kernel 生命周期里独占一个 CUDA Core。理解吞吐时，要看指令类型、发射能力和依赖。

普通 thread block 在一个 SM 上执行。只要线程数、寄存器、shared memory 等资源允许，一个 SM 可以同时驻留多个 block。

cluster 把多个 block 组织起来协作。跨 block 访问 distributed shared memory 需要满足 cluster 的生命周期和同步要求。

寄存器保存线程的操作数与状态。shared memory 由程序显式管理，供 block 内线程协作；L1 由硬件缓存访问的数据。部分架构让 L1 与 shared memory 共享物理容量，但两者的访问和管理方式不同。

Blackwell SM100 的 TMEM 是矩阵指令使用的专用存储，累加结果可以保存在其中，再搬到寄存器处理。CUDA local memory 表示线程私有的地址空间，发生寄存器溢出时可能产生设备内存访问。

PTX 是虚拟指令集。编译器把它转换成目标 GPU 的机器代码，过程中会安排指令、分配寄存器并选择具体实现。

研究实际指令和寄存器使用时，要检查对应构建产物的 SASS。比较两个 kernel 时，同时记录 GPU、编译器版本与编译选项。

先查所用指令和目标架构。SM100 的 tcgen05 / TMEM 路径不能因为另一张卡也叫 Blackwell，就直接套到 SM120。两者的矩阵指令、资源限制和最佳 tile 配置都有差异。

编译成功只解决兼容性的一部分。数据类型、布局、对齐和调度仍要按目标芯片验证。带 a 或 f 后缀的编译目标还具有各自的兼容范围。

Occupancy 描述常驻 active warp 数相对硬件上限的比例。更多可调度的 warp 有助于隐藏等待，但增大并行度也会消耗寄存器和 shared memory。

减少每个线程的寄存器，可能让更多 warp 驻留，也可能增加溢出或破坏数据复用。先看执行时间，再结合 eligible warps、发射空槽和内存流量判断。

发出指令和操作完成是两个事件。消费者必须等待对应的数据就绪；生产者复用缓冲区前，还要确认消费者已经结束读取。

TMA、WGMMA 与普通线程读写可能涉及不同的完成机制和内存代理。按该指令规定使用 barrier、wait 和 fence，不能用一个普通线程同步代替全部异步依赖。

先对齐输入精度、累加类型以及 dense / sparse 口径。峰值矩阵吞吐只描述其中一条计算路径。

实际任务还要搬运数据、执行其他算子并等待依赖。小 batch 的权重读取、长上下文的 KV 访问，以及跨卡通信，都可能限制最终吞吐。

依赖链测量前一条结果何时能供下一条使用。多条独立指令可以重叠执行，用来测稳定发射和完成工作的速率。

复核微基准时，要看依赖关系、warp 数、频率和数据所在的缓存层。同一条指令的高吞吐，可以与较长的单条延迟同时存在。

想用一个真实 kernel 串起 eligible warps、依赖链、访存并行度和缓存命中率，区分带宽不足与延迟无法隐藏。