目录
- 1. 简介
- 2. 编程模型
- 3. PTX机器模型
- 4. 语法
- 5. 状态空间、类型与变量
- 6. 指令操作数
- 7. 抽象化ABI
- 8. 内存一致性模型
-
9. 指令集
- 9.1. 指令描述的格式与语义
- 9.2. PTX指令集
- 9.3. 谓词执行
- 9.4. 指令与操作数的类型信息
- 9.5. 控制结构中线程的分歧
- 9.6. 语义
-
9.7. 指令说明
-
9.7.1. 整数算术指令
- 9.7.1.1. 整数算术指令：加法
- 9.7.1.2. 整数算术指令：sub
- 9.7.1.3. 整数算术指令：mul
- 9.7.1.4. 整数算术指令：mad
- 9.7.1.5. 整数算术指令：mul24
- 9.7.1.6. 整数算术指令：mad24
- 9.7.1.7. 整数算术指令：sad
- 9.7.1.8. 整数算术指令：div
- 9.7.1.9. 整数算术指令：rem
- 9.7.1.10. 整数算术指令：abs
- 9.7.1.11. 整数算术指令：neg
- 9.7.1.12. 整数算术指令：min
- 9.7.1.13. 整数算术指令：max
- 9.7.1.14. 整数算术指令：popc
- 9.7.1.15. 整数算术指令：clz
- 9.7.1.16. 整数算术指令：bfind
- 9.7.1.17. 整数算术指令：fns
- 9.7.1.18. 整数算术指令：brev
- 9.7.1.19. 整数算术指令：bfe
- 9.7.1.20. 整数算术指令：bfi
- 9.7.1.21. 整数算术指令：szext
- 9.7.1.22. 整数算术指令：bmsk
- 9.7.1.23. 整数算术指令：dp4a
- 9.7.1.24. 整数算术指令：dp2a
- 9.7.2. 扩展精度整数算术指令
-
9.7.3. 浮点指令
- 9.7.3.1. 浮点指令：testp
- 9.7.3.2. 浮点指令：copysign
- 9.7.3.3. 浮点指令：加法
- 9.7.3.4. 浮点指令：sub
- 9.7.3.5. 浮点指令：mul
- 9.7.3.6. 浮点指令：fma
- 9.7.3.7. 浮点指令：mad
- 9.7.3.8. 浮点指令：div
- 9.7.3.9. 浮点指令：abs
- 9.7.3.10. 浮点指令：neg
- 9.7.3.11. 浮点指令：min
- 9.7.3.12. 浮点指令：max
- 9.7.3.13. 浮点指令：rcp
- 9.7.3.14. 浮点指令：rcp.approx.ftz.f64
- 9.7.3.15. 浮点指令：sqrt
- 9.7.3.16. 浮点指令：rsqrt
- 9.7.3.17. 浮点指令：rsqrt.approx.ftz.f64
- 9.7.3.18. 浮点指令：sin
- 9.7.3.19. 浮点指令：cos
- 9.7.3.20. 浮点指令：lg2
- 9.7.3.21. 浮点指令：ex2
- 9.7.3.22. 浮点指令：tanh
- 9.7.4. 半精度浮点指令
- 9.7.5. 混合精度浮点指令
- 9.7.6. 比较与选择指令
- 9.7.7. 半精度比较指令
- 9.7.8. 逻辑与移位指令
-
9.7.9. 数据传输与转换指令
- 9.7.9.1. 缓存操作符
- 9.7.9.2. 缓存驱逐优先级提示
- 9.7.9.3. 数据移动与转换指令：mov
- 9.7.9.4. 数据移动与转换指令：mov
- 9.7.9.5. 数据移动与转换指令：shfl (已弃用)
- 9.7.9.6. 数据移动与转换指令：shfl.sync
- 9.7.9.7. 数据移动与转换指令：prmt
- 9.7.9.8. 数据移动与转换指令：ld
- 9.7.9.9. 数据移动与转换指令：ld.global.nc
- 9.7.9.10. 数据移动与转换指令：ldu
- 9.7.9.11. 数据移动与转换指令：st
- 9.7.9.12. 数据移动与转换指令：st.async
- 9.7.9.13. 数据移动与转换指令：st.bulk
- 9.7.9.14. 数据移动与转换指令：multimem.ld_reduce, multimem.st, multimem.red
- 9.7.9.15. 数据移动与转换指令：prefetch, prefetchu
- 9.7.9.16. 数据移动与转换指令：applypriority
- 9.7.9.17. 数据移动与转换指令：discard
- 9.7.9.18. 数据移动与转换指令：createpolicy
- 9.7.9.19. 数据移动与转换指令：isspacep
- 9.7.9.20. 数据移动与转换指令：cvta
- 9.7.9.21. 数据移动与转换指令：cvt
- 9.7.9.22. 数据移动与转换指令：cvt.pack
- 9.7.9.23. 数据移动与转换指令：mapa
- 9.7.9.24. 数据移动与转换指令：getctarank
- 9.7.9.25. 数据移动与转换指令：异步拷贝
- 9.7.9.26. 数据移动与转换指令：tensormap.replace
- 9.7.10. 纹理指令
- 9.7.11. 曲面指令
- 9.7.12. 控制流指令
-
9.7.13. 并行同步与通信指令
- 9.7.13.1. 并行同步与通信指令：bar, barrier
- 9.7.13.2. 并行同步与通信指令：bar.warp.sync
- 9.7.13.3. 并行同步与通信指令：barrier.cluster
- 9.7.13.4. 并行同步与通信指令：membar/fence
- 9.7.13.5. 并行同步与通信指令：atom
- 9.7.13.6. 并行同步与通信指令：red
- 9.7.13.7. 并行同步与通信指令：red.async
- 9.7.13.8. 并行同步与通信指令：vote (已弃用)
- 9.7.13.9. 并行同步与通信指令：vote.sync
- 9.7.13.10. 并行同步与通信指令：match.sync
- 9.7.13.11. 并行同步与通信指令：activemask
- 9.7.13.12. 并行同步与通信指令: redux.sync
- 9.7.13.13. 并行同步与通信指令：griddepcontrol
- 9.7.13.14. 并行同步与通信指令: elect.sync
-
9.7.13.15. 并行同步与通信指令：mbarrier
- 9.7.13.15.1. mbarrier对象的大小与对齐
- 9.7.13.15.2. mbarrier对象的内容
- 9.7.13.15.3. mbarrier对象的生命周期
- 9.7.13.15.4. mbarrier对象的阶段
- 9.7.13.15.5. 通过mbarrier对象跟踪异步操作
- 9.7.13.15.6. mbarrier对象的阶段完成
- 9.7.13.15.7. 对mbarrier对象执行Arrive-on操作
- 9.7.13.15.8. 支持共享内存的mbarrier
- 9.7.13.15.9. 并行同步与通信指令：mbarrier.init
- 9.7.13.15.10. 并行同步与通信指令：mbarrier.inval
- 9.7.13.15.11. 并行同步与通信指令：mbarrier.expect_tx
- 9.7.13.15.12. 并行同步与通信指令：mbarrier.complete_tx
- 9.7.13.15.13. 并行同步与通信指令：mbarrier.arrive
- 9.7.13.15.14. 并行同步与通信指令：mbarrier.arrive_drop
- 9.7.13.15.15. 并行同步与通信指令：cp.async.mbarrier.arrive
- 9.7.13.15.16. 并行同步与通信指令：mbarrier.test_wait/mbarrier.try_wait
- 9.7.13.15.17. 并行同步与通信指令：mbarrier.pending_count
- 9.7.13.16. 并行同步与通信指令：tensormap.cp_fenceproxy
- 9.7.13.17. 并行同步与通信指令：clusterlaunchcontrol.try_cancel
- 9.7.13.18. 并行同步与通信指令：clusterlaunchcontrol.query_cancel
-
9.7.14. Warp级矩阵乘加指令
- 9.7.14.1. 矩阵形状
- 9.7.14.2. 矩阵数据类型
- 9.7.14.3. 块缩放
- 9.7.14.4. 使用wmma指令的矩阵乘加运算
-
9.7.14.5. 使用mma指令的矩阵乘加运算
- 9.7.14.5.1. 使用.f16浮点类型的mma.m8n8k4矩阵片段
- 9.7.14.5.2. 针对mma.m8n8k4的矩阵片段（使用.f64浮点类型）
- 9.7.14.5.3. mma.m8n8k16的矩阵片段
- 9.7.14.5.4. 用于mma.m8n8k32的矩阵片段
- 9.7.14.5.5. 用于mma.m8n8k128的矩阵片段
- 9.7.14.5.6. 用于mma.m16n8k4的矩阵片段
- 9.7.14.5.7. mma.m16n8k8的矩阵片段
- 9.7.14.5.8. 浮点类型mma.m16n8k16的矩阵片段
- 9.7.14.5.9. mma.m16n8k16的矩阵片段
- 9.7.14.5.10. 用于mma.m16n8k32的矩阵片段
- 9.7.14.5.11. mma.m16n8k64 的矩阵片段
- 9.7.14.5.12. mma.m16n8k128的矩阵片段
- 9.7.14.5.13. 用于mma.m16n8k256的矩阵片段
- 9.7.14.5.14. 乘加指令：mma
- 9.7.14.5.15. Warp级矩阵加载指令：ldmatrix
- 9.7.14.5.16. Warp级矩阵存储指令：stmatrix
- 9.7.14.5.17. Warp级矩阵转置指令：movmatrix
-
9.7.14.6. 使用mma.sp指令执行稀疏矩阵A的矩阵乘积累加操作
- 9.7.14.6.1. 稀疏矩阵存储
-
9.7.14.6.2. 稀疏矩阵A的乘积累加运算矩阵片段
- 9.7.14.6.2.1. 针对.f16和.bf16类型的稀疏矩阵乘法m16n8k16的矩阵片段
- 9.7.14.6.2.2. 针对.f16和.bf16类型的稀疏矩阵乘法运算m16n8k32的矩阵片段
- 9.7.14.6.2.3. 针对稀疏矩阵乘法累加运算m16n8k16的矩阵片段（使用.tf32浮点类型）
- 9.7.14.6.2.4. 使用.tf32浮点类型的稀疏矩阵乘法m16n8k8的矩阵片段
- 9.7.14.6.2.5. 针对稀疏矩阵乘法累加运算m16n8k32的矩阵片段（支持.u8/.s8整数类型）
- 9.7.14.6.2.6. 针对稀疏矩阵乘法运算m16n8k64的矩阵片段，支持.u8/.s8/.e4m3/.e5m2/.e3m2/.e2m3/.e2m1数据类型
- 9.7.14.6.2.7. 针对稀疏矩阵乘法m16n8k64的矩阵片段，支持.u4/.s4整数类型
- 9.7.14.6.2.8. 针对.u4/.s4/.e2m1类型的稀疏矩阵乘法m16n8k128的矩阵片段
- 9.7.14.6.3. 乘加指令：mma.sp/mma.sp::ordered_metadata
- 9.7.15. 异步Warpgroup级矩阵乘加指令
-
9.7.16. TensorCore 第五代系列指令集
- 9.7.16.1. 张量内存
- 9.7.16.2. 矩阵与数据移动形状
- 9.7.16.3. 矩阵描述符
- 9.7.16.4. 问题粒度
- 9.7.16.5. 第五代TensorCore操作的内存一致性模型
- 9.7.16.6. 张量内存分配与管理指令
- 9.7.16.7. 张量内存与寄存器加载/存储指令
- 9.7.16.8. 张量内存数据移动指令
-
9.7.16.9. TensorCore 第五代矩阵乘加运算
- 9.7.16.9.1. 转置与取反运算
- 9.7.16.9.2. 矩阵布局组织
- 9.7.16.9.3. 类型大小、主序与混洗的有效组合
- 9.7.16.9.4. 张量和共享内存中元素的打包格式
- 9.7.16.9.5. 数据路径布局组织
- 9.7.16.9.6. 共享内存布局与数据重排
- 9.7.16.9.7. 块缩放
-
9.7.16.9.8. 稀疏矩阵
- 9.7.16.9.8.1. 使用.kind::tf32的稀疏tcgen05.mma.sp
- 9.7.16.9.8.2. 稀疏矩阵 tcgen05.mma.sp 支持 .kind::f16、.kind::f8f6f4、.kind::mxf8f6f4、.kind::i8 数据类型
- 9.7.16.9.8.3. 稀疏矩阵tcgen05.mma.sp，使用.kind::mxf4和.kind::mxf4nvf4
-
9.7.16.9.8.4. 稀疏选择器
- 9.7.16.9.8.4.1. 针对M=64的.kind::f16稀疏元数据矩阵布局
- 9.7.16.9.8.4.2. 针对M = 128 / M = 256的.kind::f16稀疏元数据矩阵布局
- 9.7.16.9.8.4.3. 针对M=64的.kind::tf32稀疏元数据矩阵布局
- 9.7.16.9.8.4.4. 针对M = 128 / M = 256的.kind::tf32稀疏元数据矩阵布局
- 9.7.16.9.8.4.5. 稀疏元数据矩阵布局（M = 64）适用于 .kind::f8f6f4、.kind::mxf8f6f4、.kind::i8、.kind::mxf4、.kind::mxf4nvf4
- 9.7.16.9.8.4.6. 针对M = 128 / M = 256的稀疏元数据矩阵布局（适用于.kind::f8f6f4、.kind::mxf8f6f4、.kind::i8、.kind::mxf4、.kind::mxf4nvf4类型）
- 9.7.16.9.8.5. 对齐限制
- 9.7.16.9.9. TensorCore 第五代MMA指令集
- 9.7.16.10. TensorCore 第五代专用同步操作
- 9.7.16.11. TensorCore 第五代异步同步操作
- 9.7.17. 堆栈操作指令
- 9.7.18. 视频教程
- 9.7.19. 其他指令
-
9.7.1. 整数算术指令
-
10. 特殊寄存器
- 10.1. 特殊寄存器：%tid
- 10.2. 特殊寄存器：%ntid
- 10.3. 特殊寄存器：%laneid
- 10.4. 特殊寄存器：%warpid
- 10.5. 特殊寄存器：%nwarpid
- 10.6. 特殊寄存器：%ctaid
- 10.7. 特殊寄存器：%nctaid
- 10.8. 特殊寄存器：%smid
- 10.9. 特殊寄存器：%nsmid
- 10.10. 特殊寄存器：%gridid
- 10.11. 特殊寄存器：%is_explicit_cluster
- 10.12. 特殊寄存器：%clusterid
- 10.13. 特殊寄存器：%nclusterid
- 10.14. 特殊寄存器：%cluster_ctaid
- 10.15. 特殊寄存器：%cluster_nctaid
- 10.16. 特殊寄存器：%cluster_ctarank
- 10.17. 特殊寄存器：%cluster_nctarank
- 10.18. 特殊寄存器：%lanemask_eq
- 10.19. 特殊寄存器：%lanemask_le
- 10.20. 特殊寄存器：%lanemask_lt
- 10.21. 特殊寄存器：%lanemask_ge
- 10.22. 特殊寄存器：%lanemask_gt
- 10.23. 特殊寄存器: %clock, %clock_hi
- 10.24. 特殊寄存器: %clock64
- 10.25. 特殊寄存器：%pm0..%pm7
- 10.26. 特殊寄存器：%pm0_64..%pm7_64
- 10.27. 特殊寄存器：%envreg<32>
- 10.28. 特殊寄存器：%globaltimer, %globaltimer_lo, %globaltimer_hi
- 10.29. 特殊寄存器：%reserved_smem_offset_begin, %reserved_smem_offset_end, %reserved_smem_offset_cap, %reserved_smem_offset_<2>
- 10.30. 特殊寄存器：%total_smem_size
- 10.31. 特殊寄存器：%aggr_smem_size
- 10.32. 特殊寄存器：%dynamic_smem_size
- 10.33. 特殊寄存器：%current_graph_exec
- 11. 指令
-
12. 版本说明
- 12.1. PTX ISA 版本 8.7 的变更
- 12.2. PTX ISA 版本 8.6 的变更
- 12.3. PTX ISA 版本 8.5 的变更
- 12.4. PTX ISA 版本 8.4 的变更
- 12.5. PTX ISA 版本 8.3 的变更
- 12.6. PTX ISA 版本 8.2 的变更
- 12.7. PTX ISA 版本 8.1 的变更
- 12.8. PTX ISA 版本 8.0 的变更
- 12.9. PTX ISA 版本 7.8 的变更
- 12.10. PTX ISA 版本 7.7 的变更
- 12.11. PTX ISA 版本 7.6 的变更
- 12.12. PTX ISA 版本 7.5 的变更
- 12.13. PTX ISA 7.4版本的变更
- 12.14. PTX ISA 版本 7.3 的变更
- 12.15. PTX ISA 版本 7.2 的变更
- 12.16. PTX ISA 版本 7.1 的变更
- 12.17. PTX ISA 版本 7.0 的变更
- 12.18. PTX ISA 版本 6.5 的变更
- 12.19. PTX ISA 版本 6.4 的变更
- 12.20. PTX ISA 版本 6.3 的变更
- 12.21. PTX ISA 版本 6.2 的变更
- 12.22. PTX ISA 版本 6.1 的变更
- 12.23. PTX ISA 版本 6.0 的变更
- 12.24. PTX ISA 版本 5.0 的变更
- 12.25. PTX ISA 版本 4.3 的变更
- 12.26. PTX ISA 版本 4.2 的变更
- 12.27. PTX ISA 版本 4.1 的变更
- 12.28. PTX ISA 版本 4.0 的变更
- 12.29. PTX ISA 版本 3.2 的变更
- 12.30. PTX ISA 版本 3.1 的变更
- 12.31. PTX ISA 版本 3.0 的变更
- 12.32. PTX ISA 版本 2.3 的变更
- 12.33. PTX ISA 版本 2.2 的变更
- 12.34. PTX ISA 版本 2.1 的变更
- 12.35. PTX ISA 2.0版本的变更
- 14. .pragma字符串描述
- 15. 通知