NVIDIA PTX Instructions GPU Virtual ISA NVIDIA
190 instruction families - showing 100 per page, page 1 of 2 - 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 |
|---|---|---|---|---|---|
| abs | abs.type d, a; | Arithmetic | PTX ISA 1.0 | sm_10 | Compute the absolute value of a signed or floating-point operand. |
| activemask | activemask.b32 d; | Parallel Synchronization and Communication Instructions | PTX ISA 6.2 | sm_30 | Query the bitmask of currently active (converged) lanes in the executing warp. |
| add | add.type d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Add two operands of the same type, with optional saturation for signed 32-bit integers. |
| add.cc | add.cc.type d, a, b; | Extended-Precision Integer Arithmetic Instructions | PTX ISA 1.2 | sm_20 | Performs integer addition and writes the carry-out value into the condition code register. |
| addc | addc{.cc}.type d, a, b; | Extended-Precision Integer Arithmetic Instructions | PTX ISA 1.2 | sm_20 | Performs integer addition with carry-in and optionally writes the carry-out value into the condition code register. |
| alloca | alloca.type ptr, size{, immAlign}; | Stack Manipulation Instructions | PTX ISA 7.3 | sm_52 | The alloca instruction dynamically allocates memory on the stack frame of the current function and updates the stack pointer accordingly. |
| and | and.type d, a, b; | Logic and Shift Instructions | PTX ISA 1.0 | sm_10 | Bitwise AND of two operands. |
| applypriority | applypriority{.global}.level::eviction_priority [a], size; | Data Movement and Conversion Instructions | PTX ISA 7.4 | sm_80 | The applypriority instruction applies the cache eviction priority specified by the.level::eviction_priority qualifier to the address range [a..a+size) |
| atom | atom.space.op.type d, [a], b; | Parallel Synchronization and Communication Instructions | PTX ISA 1.1 | sm_11 | Atomically read-modify-write a memory location and return the prior value. |
| bar.cta | bar.cta.sync a{, b}; | Parallel Synchronization and Communication Instructions | PTX ISA 1.0 | sm_20 | Performs barrier synchronization and communication within a CTA. |
| bar.warp.sync | bar.warp.sync membermask; | Parallel Synchronization and Communication Instructions | PTX ISA 6.0 | sm_30 | bar.warp.sync will cause executing thread to wait until all threads corresponding to membermask have executed a bar.warp.sync with the same membermask value before resuming execution. |
| barrier | barrier{.cta}.sync{.aligned} a{, b}; | Parallel Synchronization and Communication Instructions | PTX ISA 1.0 | sm_10 | Block threads in a CTA at a named barrier until the expected number of threads has arrived. |
| barrier.cluster | barrier.cluster.arrive{.sem}{.aligned}; | Parallel Synchronization and Communication Instructions | PTX ISA 7.8 | sm_90 | Performs barrier synchronization and communication within a cluster. |
| barrier.cta | barrier.cta.sync{.aligned} a{, b}; | Parallel Synchronization and Communication Instructions | PTX ISA 1.0 | sm_20 | Performs barrier synchronization and communication within a CTA. |
| bfe | bfe.type d, a, b, c; | Integer Arithmetic Instructions | PTX ISA 2.0 | sm_20 | Extract bit field from a and place the zero or sign-extended result in d. |
| bfi | bfi.type f, a, b, c, d; | Integer Arithmetic Instructions | PTX ISA 2.0 | sm_20 | Align and insert a bit field from a into b, and place the result in f. |
| bfind | bfind.type d, a; | Integer Arithmetic Instructions | PTX ISA 2.0 | sm_20 | Find the bit position of the most significant non-sign bit in a and place the result in d. |
| bmsk | bmsk.mode.b32 d, a, b; | Integer Arithmetic Instructions | PTX ISA 7.6 | sm_70 | Generates a 32-bit mask starting from the bit position specified in operand a, and of the width specified in operand b. |
| bra | @p bra{.uni} tgt; // tgt is a label | Control Flow Instructions | PTX ISA 1.0 | sm_10 | Continue execution at the target. |
| brev | brev.type d, a; | Integer Arithmetic Instructions | PTX ISA 2.0 | sm_20 | Perform bitwise reversal of input. |
| brkpt | brkpt; | Miscellaneous Instructions | PTX ISA 1.0 | sm_11 | Suspends execution. |
| brx.idx | @p brx.idx{.uni} index, tlist; | Control Flow Instructions | PTX ISA 6.0 | sm_30 | Index into a list of possible destination labels, and continue execution from the chosen label. |
| call | // direct call to named function, func is a symbol call{.uni} (ret-param), func, (param-list); |
Control Flow Instructions | PTX ISA 1.0 | sm_20 | The call instruction stores the address of the next instruction, so execution can resume at that point after executing a ret instruction. |
| clmad | clmad.mode.u64 d, a, b, c; | Integer Arithmetic Instructions | PTX ISA 9.3 | sm_80 | Performs a carryless multiplication of a and b, followed by a carryless addition of c, and writes the result into destination register d. |
| clusterlaunchcontrol.query_cancel | clusterlaunchcontrol.query_cancel.is_canceled.pred.b128 pred, try_cancel_response; | Parallel Synchronization and Communication Instructions | PTX ISA 8.6 | sm_100 | Instruction clusterlaunchcontrol.query_cancel can be used to decode opaque response written by instruction clusterlaunchcontrol.try_cancel. |
| clusterlaunchcontrol.try_cancel | clusterlaunchcontrol.try_cancel.async{.space}.completion_mechanism{.multicast::cluster::all}.b128 [addr], [mbar]; | Parallel Synchronization and Communication Instructions | PTX ISA 8.6 | sm_100 | The clusterlaunchcontrol.try_cancel instruction requests atomically cancelling the launch of a cluster that has not started running yet. |
| clz | clz.type d, a; | Integer Arithmetic Instructions | PTX ISA 2.0 | sm_20 | Count the number of leading zero bits in an integer operand. |
| cnot | cnot.type d, a; | Logic and Shift Instructions | PTX ISA 1.0 | sm_10 | Compute the logical negation using C/C++ semantics. |
| copysign | copysign.type d, a, b; | Floating-Point Instructions | PTX ISA 2.0 | sm_20 | Copy sign bit of a into value of b, and return the result as d. |
| cos | cos.approx{.ftz}.f32 d, a; | Floating-Point Instructions | PTX ISA 1.0 | sm_10 | Find the cosine of the angle a (in radians). |
| cp.async | cp.async.ca.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size} | Data Movement and Conversion Instructions | PTX ISA 7.0 | sm_80 | 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}; |
Data Movement and Conversion Instructions | PTX ISA 8.0 | sm_90 | 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; | Data Movement and Conversion Instructions | PTX ISA 8.0 | sm_90 | 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}; | Data Movement and Conversion Instructions | PTX ISA 8.0 | sm_90 | 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} |
Data Movement and Conversion Instructions | PTX ISA 8.0 | sm_90 | 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} |
Data Movement and Conversion Instructions | PTX ISA 8.0 | sm_90 | 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; | Data Movement and Conversion Instructions | PTX ISA 8.0 | sm_90 | 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 ; | Data Movement and Conversion Instructions | PTX ISA 7.0 | sm_80 | 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.mbarrier.arrive | cp.async.mbarrier.arrive{.noinc}{.shared{::cta}}.b64 [addr]; | Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | Causes an arrive-on operation to be triggered by the system on the mbarrier object upon the completion of all prior cp.async operations initiated by the executing thread. |
| cp.async.wait_all | cp.async.wait_all; | Data Movement and Conversion Instructions | PTX ISA 7.0 | sm_80 | 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; | Data Movement and Conversion Instructions | PTX ISA 7.0 | sm_80 | 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 | Data Movement and Conversion Instructions | PTX ISA 8.0 | sm_90 | 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} |
Data Movement and Conversion Instructions | PTX ISA 8.0 | sm_90 | 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; |
Data Movement and Conversion Instructions | PTX ISA 7.4 | sm_80 | 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; | Data Movement and Conversion Instructions | PTX ISA 1.0 | sm_10 | 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; | Data Movement and Conversion Instructions | PTX ISA 6.5 | sm_72 | 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 |
Data Movement and Conversion Instructions | PTX ISA 2.0 | sm_20 | 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; | Data Movement and Conversion Instructions | PTX ISA 7.4 | sm_80 | 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. |
| div | div.stype d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Divide the first operand by the second. |
| dp2a | dp2a.mode.atype.btype d, a, b, c; | Integer Arithmetic Instructions | PTX ISA 5.0 | sm_61 | Two-way 16-bit to 8-bit dot product which is accumulated in 32-bit result. |
| dp4a | dp4a.atype.btype d, a, b, c; | Integer Arithmetic Instructions | PTX ISA 5.0 | sm_61 | Four-way byte dot product which is accumulated in 32-bit result. |
| elect.sync | elect.sync d|p, membermask; | Parallel Synchronization and Communication Instructions | PTX ISA 8.0 | sm_90 | elect.sync elects one predicated active leader thread from among a set of threads specified by membermask. |
| ex2 | ex2.approx{.ftz}.f32 d, a; | Half Precision Floating-Point Instructions | PTX ISA 1.0 | sm_75 | Raise 2 to the power a. |
| exit | exit; | Control Flow Instructions | PTX ISA 1.0 | sm_10 | Ends execution of a thread. Barriers exclusively waiting on arrivals from exited threads are always released. |
| fabric.submit | fabric.submit{.submitop}; | Fabric Instructions | PTX ISA 9.3 | sm_100 | Submits prior fabric operations issued by the current thread. |
| fabric.try_get | fabric.try_get.async.dst.completion_mechanism.sem.scope.b128 [dst], [srcLeId, srcDataOff], size, [bar]; | Fabric Instructions | PTX ISA 9.3 | sm_100 | Asynchronously copies size bytes from fabric handle [srcLeId, srcDataOff] to destination memory [dst], where srcLeId is a 32-bit unsigned value denoting the logical endpoint identifier, and… |
| fabric.try_pullred | fabric.try_pullred.async.multimem.dst.completion_mechanism.sem.scope.redOpBit.typeBit.sync [dst], [srcLeId, srcDataOff], size, [bar], imm-membermask; | Fabric Instructions | PTX ISA 9.3 | sm_100 | Initiates asynchronous loads from multiple resources pointed to by multicast fabric handle [srcLeId, srcDataOff], of size bytes, and performs element- |
| fabric.try_put | fabric.try_put.async{.multimem}.src.completion_mechanism0.sem.scope.b128 [dstLeId, dstDataOff], [src], size, [bar]; | Fabric Instructions | PTX ISA 9.3 | sm_100 | Asynchronously copies size bytes from [src] to destination fabric handle [dstLeId, dstDataOff], where dstLeId is a 32-bit unsigned value denoting the logical endpoint identifier, and dstDataOff is a… |
| fabric.try_red | fabric.try_red.async{.multimem}.src.completion_mechanism0.sem.scope.redOpBit.typeBit [dstLeId, dstDataOff], [src], size, [bar]; | Fabric Instructions | PTX ISA 9.3 | sm_100 | Asynchronously copies size bytes from [src] to destination fabric handle [dstLeId, dstDataOff] with element-wise reduction, where dstLeId is a 32-bit unsigned value denoting the logical endpoint… |
| fabric.wait | fabric.wait.sync_restrict::reads; | Fabric Instructions | PTX ISA 9.3 | sm_100 | Fabric-read completion mechanism instruction fabric.wait waits on the local shared memory (.shared::cta ) reads of submitted fabric operations. |
| fence | // Thread fence: fence{.sem}.scope; |
Parallel Synchronization and Communication Instructions | PTX ISA 1.4 | sm_20 | The fence instruction establishes an ordering between memory accesses requested by this thread, as described by the memory consistency model. |
| fma | fma.rn.f64 d, a, b, c; | Arithmetic | PTX ISA 1.4 | sm_13 | Compute (a * b) + c with a single rounding step for improved precision over mad. |
| fns | fns.b32 d, mask, base, offset; | Integer Arithmetic Instructions | PTX ISA 6.0 | sm_30 | Given a 32-bit value mask and an integer value base (between 0 and 31), find the n-th (given by offset) set bit in mask from the base bit, and store the bit position in d. |
| getctarank | getctarank{.space}.type d, a; | Data Movement and Conversion Instructions | PTX ISA 7.8 | sm_90 | Write the destination register d with the rank of the CTA which contains the address specified in operand a. |
| griddepcontrol | griddepcontrol.action; | Parallel Synchronization and Communication Instructions | PTX ISA 7.8 | sm_90 | The griddepcontrol instruction allows the dependent grids and prerequisite grids as defined by the runtime, to control execution in the following way: |
| isspacep | isspacep.space p, a; // result is .pred | Data Movement and Conversion Instructions | PTX ISA 2.0 | sm_20 | Write predicate register p with 1 if generic address a falls within the specified state space window and with 0 otherwise. |
| istypep | istypep.type p, a; // result is .pred | Texture Instructions | PTX ISA 4.0 | sm_30 | Write predicate register p with 1 if register a points to an opaque variable of the specified type, and with 0 otherwise. Destination p has type.pred; |
| ld | ld.space.type d, [a]; | Data Movement and Conversion Instructions | PTX ISA 1.0 | sm_10 | 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}; | Data Movement and Conversion Instructions | PTX ISA 3.1 | sm_32 | 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. |
| ldmatrix | ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p]; | Warp Level Matrix Multiply-Accumulate Instructions | PTX ISA 6.5 | sm_75 | 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. |
| ldu | ldu{.ss}.type d, [a]; // load from address | Data Movement and Conversion Instructions | PTX ISA 2.0 | sm_13 | 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. |
| lg2 | lg2.approx.f32 d, a; | Arithmetic | PTX ISA 1.0 | sm_10 | Fast hardware approximation of log2(x). |
| lop3 | lop3.b32 d, a, b, c, immLut; | Logic and Shift Instructions | PTX ISA 4.3 | sm_50 | Compute bitwise logical operation on inputs a, b, c and store the result in destination d. |
| mad | mad.mode.stype d, a, b, c; | Arithmetic | PTX ISA 1.0 | sm_10 | Compute (a * b) + c as two rounding steps (unlike fma, which fuses them into one). |
| mad.cc | mad{.hi,.lo}.cc.type d, a, b, c; | Extended-Precision Integer Arithmetic Instructions | PTX ISA 3.0 | sm_10 | Multiplies two values, extracts either the high or low part of the result, and adds a third value. |
| mad24 | mad24.mode.type d, a, b, c; | Integer Arithmetic Instructions | PTX ISA 1.0 | sm_10 | Compute the product of two 24-bit integer values held in 32-bit source registers, and add a third, 32-bit value to either the high or low 32-bits of the 48-bit result. |
| madc | madc{.hi,.lo}{.cc}.type d, a, b, c; | Extended-Precision Integer Arithmetic Instructions | PTX ISA 3.0 | sm_10 | Multiplies two values, extracts either the high or low part of the result, and adds a third value along with carry-in. |
| mapa | mapa{.space}.type d, a, b; | Data Movement and Conversion Instructions | PTX ISA 7.8 | sm_90 | Get address in the CTA specified by operand b which corresponds to the address specified by operand a. |
| match.sync | match.any.sync.type d, a, membermask; | Parallel Synchronization and Communication Instructions | PTX ISA 6.0 | sm_70 | match.sync will cause executing thread to wait until all non-exited threads from membermask have executed match.sync with the same qualifiers and same membermask value before resuming execution. |
| max | max.type d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Select the larger of two operands. |
| mbarrier.arrive | mbarrier.arrive{.sem.scope}{.shared{::cta}}.b64 state, [addr]{, count}; | Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | A thread executing mbarrier.arrive performs an arrive-on operation on the mbarrier object at the location specified by the address operand addr. The 3 |
| mbarrier.arrive_drop | mbarrier.arrive_drop{.sem.scope}{.shared{::cta}}.b64 state, [addr] {, count}; | Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | A thread executing mbarrier.arrive_drop on the mbarrier object at the location specified by the address operand addr performs the following steps: Decrements the expected arrival count of the mbarrier object by the value specified by the 32-bit integer operand count. |
| mbarrier.check_layout | mbarrier.check_layout.layout{.ss}.b64 p, [addr]; | Parallel Synchronization and Communication Instructions | PTX ISA 9.3 | sm_90 | The layout of the opaque mbarrier object can be queried using mbarrier.check_layout. |
| mbarrier.complete_tx | mbarrier.complete_tx{.sem.scope}{.space}.b64 [addr], txCount; | Parallel Synchronization and Communication Instructions | PTX ISA 8.0 | sm_90 | A thread executing mbarrier.complete_tx performs a complete-tx operation on the mbarrier object at the location specified by the address operand addr. |
| mbarrier.expect_tx | mbarrier.expect_tx{.sem.scope}{.space}.b64 [addr], txCount; | Parallel Synchronization and Communication Instructions | PTX ISA 8.0 | sm_90 | A thread executing mbarrier.expect_tx performs an expect-tx operation on the mbarrier object at the location specified by the address operand addr. |
| mbarrier.init | mbarrier.init{.layout}{.shared{::cta}}.b64 [addr], count; | Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | mbarrier.init initializes the mbarrier object at the location specified by the address operand addr with the unsigned 32-bit integer count. |
| mbarrier.inval | mbarrier.inval{.shared{::cta}}.b64 [addr]; | Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | mbarrier.inval invalidates the mbarrier object at the location specified by the address operand addr. |
| mbarrier.pending_count | mbarrier.pending_count{.layout}.b64 count, state; | Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | The pending count can be queried from the opaque mbarrier state using mbarrier.pending_count. |
| mbarrier.test_wait | // without parity mbarrier.test_wait{.phase_type::primary}{.sem.scope}{.ss}.b64 waitComplete, [addr], state; |
Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | The test_wait and try_wait operations test for the completion of the current or the immediately preceding phase of an mbarrier object at the location specified by the operand addr. |
| mbarrier.try_wait | // without parity mbarrier.try_wait{.phase_type::primary}{.sem.scope}{.ss}.b64 waitComplete, [addr], state {, timeHint}; |
Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | The test_wait and try_wait operations test for the completion of the current or the immediately preceding phase of an mbarrier object at the location specified by the operand addr. |
| membar | membar.level; | Parallel Synchronization and Communication Instructions | PTX ISA 1.4 | sm_20 | Order this thread's prior memory accesses relative to later ones, visible to a given scope. |
| min | min.type d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Select the smaller of two operands. |
| mma | mma.sync.aligned.shape.row.col.dtype.atype.btype.ctype d, a, b, c; | Warp Level Matrix Multiply-Accumulate Instructions | PTX ISA 6.4 | sm_70 | Cooperative, warp-wide matrix-multiply-accumulate executed on tensor-core hardware. |
| mov | mov.type d, a; | Data Movement and Conversion Instructions | PTX ISA 1.0 | sm_10 | Copy a value into a register, or materialize an address/immediate. |
| movmatrix | movmatrix.sync.aligned.shape.trans.type d, a; | Warp Level Matrix Multiply-Accumulate Instructions | PTX ISA 7.8 | sm_75 | Move a row-major matrix across all threads in a warp, reading elements from source a, and writing the transposed elements to destination d. |
| mul | mul.mode.stype d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Multiply two operands, selecting the low, high, or widened part of an integer product. |
| mul24 | mul24.mode.type d, a, b; | Integer Arithmetic Instructions | PTX ISA 1.0 | sm_10 | Compute the product of two 24-bit integer values held in 32-bit source registers, and return either the high or low 32-bits of the 48-bit result. |
| multimem.cp.async.bulk | multimem.cp.async.bulk{.sem}.dst.src.completion_mechanism{.cp_mask} | Data Movement and Conversion Instructions | PTX ISA 9.1 | sm_90 | 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; | Data Movement and Conversion Instructions | PTX ISA 9.1 | sm_90 | 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]; |
Data Movement and Conversion Instructions | PTX ISA 8.1 | sm_90 | 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. |
Source
Normalized from NVIDIA's official Parallel Thread Execution ISA documentation. docs.nvidia.com/cuda/parallel-thread-execution ↗