Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions Instructions
5 PTX instructions in this category - click any row for encoding, pseudocode, and full documentation.
| Mnemonic | Syntax | PTX Version | Summary |
|---|---|---|---|
| wgmma.commit_group | wgmma.commit_group.sync.aligned; | PTX ISA 8.0 | 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; | PTX ISA 8.0 | 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; | PTX ISA 8.0 | 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; | PTX ISA 8.2 | 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; | PTX ISA 8.0 | 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. |