NVIDIA PTX Instructions GPU Virtual ISA NVIDIA
210 instruction families - showing 10 per page, page 3 of 3 - PTX (Parallel Thread Execution) is a portable, low-level virtual GPU instruction set. It is not native SASS machine code: NVIDIA's compiler stack lowers PTX text to a target architecture's real machine code at install- or run-time, so one PTX instruction can produce different native encodings on different targets. Each instruction may require both a minimum PTX ISA version and a minimum target architecture (compute capability).
| Mnemonic | Syntax | Category | PTX Version | Target | Summary |
|---|---|---|---|---|---|
| vsub | // 32-bit scalar operation, with optional secondary operation vsub.dtype.atype.btype{.sat} d, a{.asel}, b{.bsel}; |
Scalar Video Instructions | PTX ISA 2.0 | sm_20 | Perform scalar arithmetic operation with optional saturate, and optional secondary arithmetic operation or subword data merge. |
| vsub2 | // SIMD instruction with secondary SIMD merge operation vsub2.dtype.atype.btype{.sat} d{.mask}, a{.asel}, b{.bsel}, c; |
SIMD Video Instructions | PTX ISA 3.0 | sm_30 | Two-way SIMD parallel arithmetic operation with secondary operation. |
| vsub4 | // SIMD instruction with secondary SIMD merge operation vsub4.dtype.atype.btype{.sat} d{.mask}, a{.asel}, b{.bsel}, c; |
SIMD Video Instructions | PTX ISA 3.0 | sm_30 | Four-way SIMD parallel arithmetic operation with secondary operation. |
| wgmma.commit_group | wgmma.commit_group.sync.aligned; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.0 | sm_90a | wgmma.commit_group instruction creates a new wgmma-group per warpgroup and batches all prior wgmma.mma_async instructions initiated by the executing warp but not committed to any wgmma-group into the new wgmma-group. |
| wgmma.fence | wgmma.fence.sync.aligned; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.0 | sm_90a | wgmma.fence instruction establishes an ordering between prior accesses to any warpgroup registers and subsequent accesses to the same registers by a wgmma.mma_async instruction. |
| wgmma.mma_async | wgmma.mma_async.sync.aligned.shape.dtype.f16.f16 d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.0 | sm_90a | Instruction wgmma.mma_async issues a MxNxK matrix multiply and accumulate operation, D = A*B+D, where the A matrix is MxK, the B matrix is KxN, and the D matrix is MxN. |
| wgmma.mma_async.sp | wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16 d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.2 | sm_90a | Instruction wgmma.mma_async issues a MxNxK matrix multiply and accumulate operation, D = A*B+D, where the A matrix is MxK, the B matrix is KxN, and the D matrix is MxN. |
| wgmma.wait_group | wgmma.wait_group.sync.aligned N; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.0 | sm_90a | wgmma.wait_group instruction will cause the executing thread to wait until only N or fewer of the most recent wgmma-groups are pending and all the prior wgmma-groups committed by the executing threads are complete. |
| wmma | wmma.load.a.sync.aligned.layout.shape.type r, [p]; | Warp Level Matrix Multiply-Accumulate Instructions | PTX ISA 6.0 | sm_70 | Higher-level warp matrix-multiply-accumulate built from explicit load/mma/store steps. |
| xor | xor.type d, a, b; | Logic and Shift Instructions | PTX ISA 1.0 | sm_10 | Bitwise XOR of two operands. |
Source
Normalized from NVIDIA's official Parallel Thread Execution ISA documentation. docs.nvidia.com/cuda/parallel-thread-execution ↗