Nvidia PTX 指令理解
输入“/”快速插入内容
Nvidia PTX 指令理解
Nvidia PTX 指令理解
用户54用户542025年11月9日修改
2024年7月15日创建
13597
14438
PTX 全称 parallel thread execution。
1.
PTX 是一个通用的、虚拟的、低层次的指令集，不是特定 GPU 的硬件指令。
2.
核心模型是将 GPU 抽象为一个数据并行计算设备。理论上，任何支持此模型的硬件都可实现 PTX。
3.
PTX 的“一次编写，多处运行”，依赖 GPU driver 在运行时的即时编译（JIT）功能。
主要 cuda & PTX version
CUDA 版本 | 开始支持 | Compute Capability | PTX ISA version | NV driver | 点评 |
CUDA 8 | P4 / P40 | sm_61 | PTX 5.0 | r361 | PTX 5.0 与 6.0 之间没有小版本 |
CUDA 9 | V100 | sm_70 | PTX 6.0 | r384 | V100 是载入史册的版本 |
CUDA 10 | T4 | sm_75 | PTX 6.3 | r400 | 一代推理神卡 |
CUDA 11 | A100 | sm_80 | PTX 7.0 | r445 | 大模型的神卡 |
CUDA 11.8 | RTX 4090 L4 / L40 | sm_89 | PTX 7.8 | r520 | 支持不带加速 feature H100，sm_90 |
CUDA 12 | H100 H200 / GH200 | sm_90a | PTX 8.0 | r525 | adds wgmma and setmaxnreg |
CUDA 12.7 | B200 GB200 | sm_100 sm_100a | PTX ISA 8.6 | r565 | |
CUDA 12.8 | 5090 | sm_120 | PTX ISA 8.7 | r570 | |
语法体系
case sensitive. uses lowercase for keywords
each PTX module 都包含 .version 和 .target 字段。
指令中，可以写 optional guard predicate。@p, 取反是 @!p
interger constants 是 64-bits, 包括 signed .s64 和 unsigned .u64。
signed/unsigned 在处理 division 和 ordered comparison 时比较需要。
statements