Warp 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 |
|---|---|---|---|
| ldmatrix | ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p]; | PTX ISA 6.5 | Collectively load one or more matrices across all threads in a warp from the location indicated by the address operand p, from.shared state space into destination register r. |
| mma | mma.sync.aligned.shape.row.col.dtype.atype.btype.ctype d, a, b, c; | PTX ISA 6.4 | Cooperative, warp-wide matrix-multiply-accumulate executed on tensor-core hardware. |
| movmatrix | movmatrix.sync.aligned.shape.trans.type d, a; | PTX ISA 7.8 | Move a row-major matrix across all threads in a warp, reading elements from source a, and writing the transposed elements to destination d. |
| stmatrix | stmatrix.sync.aligned.shape.num{.trans}{.ss}.type [p], r; | PTX ISA 7.8 | Collectively store one or more matrices across all threads in a warp to the location indicated by the address operand p, in.shared state space. |
| wmma | wmma.load.a.sync.aligned.layout.shape.type r, [p]; | PTX ISA 6.0 | Higher-level warp matrix-multiply-accumulate built from explicit load/mma/store steps. |