Tensor Memory and Register Load/Store Instructions Instructions
3 PTX instructions in this category - click any row for encoding, pseudocode, and full documentation.
| Mnemonic | Syntax | PTX Version | Summary |
|---|---|---|---|
| tcgen05.ld | // Base load instruction: tcgen05.ld.sync.aligned.shape1.num{.pack}.b32 r, [taddr]; |
PTX ISA 8.6 | Instruction tcgen05.ld asynchronously loads data from the Tensor Memory at the location specified by the 32-bit address operand taddr into the destination register r, collectively across all threads of the warps. |
| tcgen05.st | tcgen05.st.sync.aligned.shape1.num{.unpack}.b32 [taddr], r; | PTX ISA 8.6 | Instruction tcgen05.st asynchronously stores data from the source register r into the Tensor Memory at the location specified by the 32-bit address operand taddr, collectively across all threads of the warps. |
| tcgen05.wait | tcgen05.wait_operation.sync.aligned; | PTX ISA 8.6 | Instruction tcgen05.wait::st causes the executing thread to block until all prior tcgen05.st operations issued by the executing thread have completed. |