NVIDIA 的 Hopper 架构(以 H100 为代表,2022 年发布)在指令集和编程模型上相比此前的 Ampere 架构引入了不少新特性。严格来说,GPU 的“指令集”更多体现在 PTX(并行线程执行)指令 和 SASS(机器码) 层面,以及 CUDA 计算能力(Compute Capability) 的新特性上。
以下是 Hopper 架构主要支持或新增的重要指令集 / 指令特性:
1. 新的计算能力版本
- CUDA Compute Capability 9.0
- 引入了大量仅在该架构可用或优化的 PTX / SASS 指令
2. Tensor Core 相关新指令(关键更新)
Hopper 的第四代 Tensor Core 支持新数据类型和新指令:
新数据类型支持
- FP8(E4M3 / E5M2)
- FP16 / BF16 / TF32 继续支持
- INT8 / INT4 / FP4(部分通过后续更新)
新 Tensor Core 指令特性
- wgmma(Warp Group Matrix Multiply-Accumulate)
- 这是 Hopper 最重要新增 PTX 指令之一
- 支持 Warp Group(4 个 warp = 128 线程) 级矩阵乘
- 可异步执行,配合 TMA 使用
- 支持 异步 Tensor Core 执行
- 支持 FP8 Tensor Core 指令
3. TMA(Tensor Memory Accelerator)指令
Hopper 新增专用硬件 TMA,并暴露对应指令:
- cp.async.bulk.tensor
- cp.async.bulk.commit_group
- cp.async.bulk.wait_group
- 支持全局内存 ↔ 共享内存的 多维张量拷贝
- 大幅减少线程参与数据搬运
4. 线程块集群(Thread Block Cluster)相关指令
- 新增 cluster 级别编程模型
- 支持 cluster sync
- 相关指令 / 内置函数:
__cluster_sync()- 分布式共享内存访问(cluster shared memory)
5. 异步执行与调度指令
- 异步事务屏障(mbarrier)
- 更完善的 pipe / stage 控制
- 支持 kernel 内多 TMA + Tensor Core 流水线
6. 其他 PTX / SASS 层面更新
- 新的 memory scope(cluster / device)
- 更细粒度 memory consistency 控制
- 新增部分 整数 / 位运算优化指令
- 改进的 atomic 操作性能
7. 与 Ampere 不兼容的关键点
以下特性 仅 Hopper 及以上支持:
wgmma- FP8 Tensor Core
- TMA bulk 指令
- Thread Block Cluster
总结一句话
Hopper 架构在指令集上最核心的新增是:wgmma(Warp Group MMA)指令、TMA bulk 拷贝指令、FP8 Tensor Core 指令,以及 Thread Block Cluster 相关同步与共享内存指令。
如果你愿意,我也可以:
- 给你 具体 PTX 示例(wgmma / TMA)
- 对比 Ampere vs Hopper 指令差异表
- 说明 哪些指令在 H100 上性能提升最明显