Data Movement and Conversion Instructions Instructions
39 PTX instructions in this category - click any row for encoding, pseudocode, and full documentation.
| Mnemonic | Syntax | PTX Version | Summary |
|---|---|---|---|
| applypriority | applypriority{.global}.level::eviction_priority [a], size; | PTX ISA 7.4 | The applypriority instruction applies the cache eviction priority specified by the.level::eviction_priority qualifier to the address range [a..a+size) |
| cp.async | cp.async.ca.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size} | PTX ISA 7.0 | cp.async is a non-blocking instruction which initiates an asynchronous copy operation of data from the location specified by source address operand src to the location specified by destination address operand dst. |
| cp.async.bulk | // global -> shared::cta cp.async.bulk{.sem}.dst.src.completion_mechanism{.level::cache_hint}{.ignore_oob} [dstMem], [srcMem], size{, ignoreBytesLeft, ignoreBytesRight}, [mbar] {, cache_policy}; |
PTX ISA 8.0 | cp.async.bulk is a non-blocking instruction which initiates an asynchronous bulk-copy operation from the location specified by source address operand srcMem to the location specified by destination address operand dstMem. |
| cp.async.bulk.commit_group | cp.async.bulk.commit_group; | PTX ISA 8.0 | cp.async.bulk.commit_group instruction creates a new per-thread bulk async-group and batches all prior cp{.reduce}.async.bulk{.prefetch}{.tensor} instructions satisfying the following conditions into… |
| cp.async.bulk.prefetch | cp.async.bulk.prefetch.L2.src{.level::cache_hint} [srcMem], size {, cache_policy}; | PTX ISA 8.0 | cp.async.bulk.prefetch is a non-blocking instruction which may initiate an asynchronous prefetch of data from the location specified by source address operand srcMem, in.src statespace, to the L2 cache. |
| cp.async.bulk.prefetch.tensor | // global -> L2: cp.async.bulk.prefetch.tensor.dim.L2.src{.load_mode}{.level::cache_hint} [tensorMap, tensorCoords] {, im2colInfo } {, cache_policy} |
PTX ISA 8.0 | cp.async.bulk.prefetch.tensor is a non-blocking instruction which may initiate an asynchronous prefetch of tensor data from the location in.src statespace to the L2 cache. |
| cp.async.bulk.tensor | // global -> shared::cta cp.async.bulk.tensor.dim.dst.src{.load_mode}.completion_mechanism{.cta_group}{.level::cache_hint} [dstMem], [tensorMap, tensorCoords], [mbar]{, im2colInfo} {, cache_policy} |
PTX ISA 8.0 | cp.async.bulk.tensor is a non-blocking instruction which initiates an asynchronous copy operation of tensor data from the location in.src state space to the location in the.dst state space. |
| cp.async.bulk.wait_group | cp.async.bulk.wait_group{.read} N; | PTX ISA 8.0 | cp.async.bulk.wait_group instruction will cause the executing thread to wait until only N or fewer of the most recent bulk async-groups are pending and all the prior bulk async-groups committed by the executing threads are complete. |
| cp.async.commit_group | cp.async.commit_group ; | PTX ISA 7.0 | cp.async.commit_group instruction creates a new cp.async-group per thread and batches all prior cp.async instructions initiated by the executing thread but not committed to any cp.async-group into the new cp.async-group. |
| cp.async.wait_all | cp.async.wait_all; | PTX ISA 7.0 | cp.async.wait_all instruction will cause the executing thread to wait until all the prior cp.async operations are complete. It is equivalent to cp.async.commit_group immediately followed by cp.async.wait_group 0. |
| cp.async.wait_group | cp.async.wait_group N; | PTX ISA 7.0 | cp.async.wait_group instruction will cause executing thread to wait till only N or fewer of the most recent cp.async-group s are pending and all the prior cp.async-group s committed by the executing threads are complete. |
| cp.reduce.async.bulk | cp.reduce.async.bulk{.sem.scope}.dst.src.completion_mechanism.redOp.type | PTX ISA 8.0 | cp.reduce.async.bulk is a non-blocking instruction which initiates an asynchronous reduction operation on an array of memory locations specified by the destination address operand dstMem with the source array whose location is specified by the source address operand srcMem. |
| cp.reduce.async.bulk.tensor | // shared::cta -> global: cp.reduce.async.bulk.tensor.dim.dst.src.redOp{.load_mode}.completion_mechanism{.level::cache_hint} [tensorMap, tensorCoords], [srcMem] {,cache_policy} |
PTX ISA 8.0 | cp.reduce.async.bulk.tensor is a non-blocking instruction which initiates an asynchronous reduction operation of tensor data in the.dst state space with tensor data in the.src state space. |
| createpolicy | // Range-based policy createpolicy.range{.global}.level::primary_priority{.level::secondary_priority}.b64 cache_policy, [a], primary-size, total-size; |
PTX ISA 7.4 | The createpolicy instruction creates a cache eviction policy for the specified cache level in an opaque 64-bit register specified by the destination operand cache_policy. |
| cvt | cvt.dtype.atype d, a; | PTX ISA 1.0 | Convert a value between integer and/or floating-point types with an explicit rounding mode. |
| cvt.pack | cvt.pack.sat.convertType.abType d, a, b; | PTX ISA 6.5 | Convert two 32-bit integers a and b into specified type and pack the results into d. |
| cvta | // convert const, global, local, or shared address to generic address cvta.space.size p, a; // source address in register a |
PTX ISA 2.0 | Convert a const, Kernel Function Parameters (.param ), global, local, or shared address to a generic address, or vice-versa. |
| discard | discard{.global}.level [a], size; | PTX ISA 7.4 | Semantically, this behaves like a weak write of an unstable indeterminate value: reads of memory locations with unstable indeterminate values may return different bit patterns each time until the memory is overwritten. |
| getctarank | getctarank{.space}.type d, a; | PTX ISA 7.8 | Write the destination register d with the rank of the CTA which contains the address specified in operand a. |
| isspacep | isspacep.space p, a; // result is .pred | PTX ISA 2.0 | Write predicate register p with 1 if generic address a falls within the specified state space window and with 0 otherwise. |
| ld | ld.space.type d, [a]; | PTX ISA 1.0 | Load a value from the specified state space into a register. |
| ld.global.nc | ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache_policy}; | PTX ISA 3.1 | Load register variable d from the location specified by the source address operand a in the global state space, and optionally cache in non-coherent read-only cache. |
| ldu | ldu{.ss}.type d, [a]; // load from address | PTX ISA 2.0 | Load read-only data into register variable d from the location specified by the source address operand a in the global state space, where the address is guaranteed to be the same across all threads in the warp. |
| mapa | mapa{.space}.type d, a, b; | PTX ISA 7.8 | Get address in the CTA specified by operand b which corresponds to the address specified by operand a. |
| mov | mov.type d, a; | PTX ISA 1.0 | Copy a value into a register, or materialize an address/immediate. |
| multimem.cp.async.bulk | multimem.cp.async.bulk{.sem}.dst.src.completion_mechanism{.cp_mask} | PTX ISA 9.1 | Instruction multimem.cp.async.bulk initiates an asynchronous bulk-copy operation from source address range [srcMem, srcMem + size) to memory locations residing on each GPU’s memory referred to by the destination multimem address range [dstMem, dstMem + size). |
| multimem.cp.reduce.async.bulk | multimem.cp.reduce.async.bulk{.sem.scope}.dst.src.completion_mechanism.redOp.type [dstMem], [srcMem], size; | PTX ISA 9.1 | Instruction multimem.cp.reduce.async.bulk initiates an element-wise asynchronous reduction operation with elements from source memory address range [srcMem, srcMem + size) to memory locations residing on each GPU’s memory referred to by the multimem destination address range [dstMem, dstMem + size). |
| multimem.ld_reduce | // Integer type: multimem.ld_reduce{.ldsem}{.scope}{.ss}.op.type d, [a]; |
PTX ISA 8.1 | The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to. Multimem addresses can be accessed only by multimem.* operations. Accessing a multimem address with ld, st or any other memory operations results in undefined behavior. Refer to CUDA programming guide for creation and management of the multimem addresses. |
| multimem.red | // Integer type: multimem.red{.redsem}{.scope}{.ss}.op.type [a], b; |
PTX ISA 8.1 | The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to. Multimem addresses can be accessed only by multimem.* operations. Accessing a multimem address with ld, st or any other memory operations results in undefined behavior. Refer to CUDA programming guide for creation and management of the multimem addresses. |
| multimem.st | // Integer type: multimem.st{.stsem}{.scope}{.ss}.type [a], b; |
PTX ISA 8.1 | The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to. Multimem addresses can be accessed only by multimem.* operations. Accessing a multimem address with ld, st or any other memory operations results in undefined behavior. Refer to CUDA programming guide for creation and management of the multimem addresses. |
| multimem.st.async | multimem.st.async.sem.scope{.ss}.type [a], b; | PTX ISA 9.3 | multimem.st.async is a non-blocking instruction which initiates an asynchronous store operation that stores the value specified by source operand b to the memory locations residing on each GPU’s memory referred to by the destination multimem address operand a. |
| prefetch | prefetch{.space}.level [a]; // prefetch to data cache | PTX ISA 2.0 | The prefetch instruction brings the cache line containing the specified address in global or local memory state space into the specified cache level. |
| prefetchu | prefetchu.L1 [a]; // prefetch to uniform cache | PTX ISA 2.0 | The prefetchu instruction brings the cache line containing the specified generic address into the specified uniform cache level. A prefetch to a shared memory location performs no operation. |
| prmt | prmt.b32{.mode} d, a, b, c; | PTX ISA 2.0 | Pick four arbitrary bytes from two 32-bit registers, and reassemble them into a 32-bit destination register. |
| shfl | shfl.mode.b32 d[|p], a, b, c; | PTX ISA 6.0 | Exchange a value directly between lanes of the same warp. |
| st | st.space.type [a], b; | PTX ISA 1.0 | Store a register value into the specified state space. |
| st.async | st.async{.weak}{.ss}.completion_mechanism{.vec}.type [a], b, [mbar]; | PTX ISA 8.1 | st.async is a non-blocking instruction which initiates an asynchronous store operation that stores the value specified by source operand b to the destination memory location specified by operand a. |
| st.bulk | st.bulk{.weak}{.shared::cta} [a], size, initval; // initval must be zero | PTX ISA 8.6 | st.bulk instruction initializes a region of shared memory starting from the location specified by destination address operand a. |
| tensormap.replace | tensormap.replace.mode.field1{.ss}.b1024.type [addr], new_val; | PTX ISA 8.3 | The tensormap.replace instruction replaces the field, specified by.field qualifier, of the tensor-map object at the location specified by the address operand addr with a new value. |