目录
- 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. 整数算术指令：add
- 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. 浮点指令：add
- 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. mma.m8n8k4 与 .f16 浮点类型的矩阵片段
- 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. 稀疏 mma.m16n8k16 与 .f16 和 .bf16 类型的矩阵片段
- 9.7.14.6.2.2. 稀疏 mma.m16n8k32 与 .f16 和 .bf16 类型的矩阵片段
- 9.7.14.6.2.3. 稀疏 mma.m16n8k16 与 .tf32 浮点类型的矩阵片段
- 9.7.14.6.2.4. 稀疏 mma.m16n8k8 与 .tf32 浮点类型的矩阵片段
- 9.7.14.6.2.5. 稀疏 mma.m16n8k32 与 .u8/.s8 整数类型的矩阵片段
- 9.7.14.6.2.6. 稀疏 mma.m16n8k64 与 .u8/.s8/.e4m3/.e5m2/.e3m2/.e2m3/.e2m1 类型的矩阵片段
- 9.7.14.6.2.7. 稀疏 mma.m16n8k64 与 .u4/.s4 整数类型的矩阵片段
- 9.7.14.6.2.8. 稀疏 mma.m16n8k128 与 .u4/.s4/.e2m1 类型的矩阵片段
- 9.7.14.6.3. 乘法和累加指令：mma.sp/mma.sp::ordered_metadata
-
9.7.15. 异步 Warpgroup 级别矩阵乘法累加指令
- 9.7.15.1. Warpgroup
- 9.7.15.2. 矩阵形状
- 9.7.15.3. 矩阵数据类型
- 9.7.15.4. 异步代理
- 9.7.15.5. 使用 wgmma.mma_async 指令的异步 Warpgroup 级别矩阵乘法累加操作
- 9.7.15.6. 使用 wgmma.mma_async.sp 指令的异步 Warpgroup 级别乘法和累加操作
- 9.7.15.7. 异步 wgmma 代理操作
-
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. 稀疏 tcgen05.mma.sp 与 .kind::tf32
- 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. 通知