NVIDIA PTX Instructions GPU Virtual ISA NVIDIA
190 instruction families - showing 90 per page, page 2 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 |
|---|---|---|---|---|---|
| multimem.red | // Integer type: multimem.red{.redsem}{.scope}{.ss}.op.type [a], b; |
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. |
| multimem.red.async | multimem.red.async.sem.scope{.ss}.op.type [a], b; | Parallel Synchronization and Communication Instructions | PTX ISA 9.3 | sm_90 | multimem.red.async is a non-blocking instruction which initiates an asynchronous reduction operation specified by.op, with operand b and the value at memory locations residing on each GPU’s memory referred to by the destination multimem address operand a. |
| multimem.st | // Integer type: multimem.st{.stsem}{.scope}{.ss}.type [a], b; |
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. |
| multimem.st.async | multimem.st.async.sem.scope{.ss}.type [a], b; | Data Movement and Conversion Instructions | PTX ISA 9.3 | sm_90 | 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. |
| nanosleep | nanosleep.u32 t; | Miscellaneous Instructions | PTX ISA 6.3 | sm_70 | Suspends the thread for a sleep duration approximately close to the delay t, specified in nanoseconds. |
| neg | neg.type d, a; | Arithmetic | PTX ISA 1.0 | sm_10 | Negate a signed or floating-point operand. |
| not | not.type d, a; | Logic and Shift Instructions | PTX ISA 1.0 | sm_10 | Bitwise complement of an operand. |
| or | or.type d, a, b; | Logic and Shift Instructions | PTX ISA 1.0 | sm_10 | Bitwise OR of two operands. |
| pmevent | pmevent a; // trigger a single performance monitor event | Miscellaneous Instructions | PTX ISA 1.4 | sm_20 | Triggers one or more of a fixed number of performance monitor events, with event index or mask specified by immediate operand a. |
| popc | popc.type d, a; | Integer Arithmetic Instructions | PTX ISA 2.0 | sm_20 | Count the number of set bits in an integer operand. |
| prefetch | prefetch{.space}.level [a]; // prefetch to data cache | Data Movement and Conversion Instructions | PTX ISA 2.0 | sm_20 | 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 | Data Movement and Conversion Instructions | PTX ISA 2.0 | sm_20 | 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; | Data Movement and Conversion Instructions | PTX ISA 2.0 | sm_20 | Pick four arbitrary bytes from two 32-bit registers, and reassemble them into a 32-bit destination register. |
| rcp | rcp.approx{.ftz}.f32 d, a; // fast, approximate reciprocal | Floating-Point Instructions | PTX ISA 1.0 | sm_20 | Compute 1/a, store result in d. |
| rcp.approx.ftz.f64 | rcp.approx.ftz.f64 d, a; | Floating-Point Instructions | PTX ISA 2.1 | sm_20 | Compute a fast, gross approximation to the reciprocal as follows: extract the most-significant 32 bits of.f64 operand a in 1.11.20 IEEE floating-point format (i.e., ignore the least-significant 32… |
| red | red.space.op.type [a], b; | Parallel Synchronization and Communication Instructions | PTX ISA 1.2 | sm_11 | Atomically read-modify-write a memory location without returning the prior value. |
| red.async | // Increment and Decrement reductions red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar]; |
Parallel Synchronization and Communication Instructions | PTX ISA 8.1 | sm_90 | red.async is a non-blocking instruction which initiates an asynchronous reduction operation specified by.op, with the operand b and the value at destination shared memory location specified by operand a. |
| redux.sync | redux.sync.op.type dst, src, membermask; | Parallel Synchronization and Communication Instructions | PTX ISA 7.0 | sm_80 | redux.sync will cause the executing thread to wait until all non-exited threads corresponding to membermask have executed redux.sync with the same qualifiers and same membermask value before resuming execution. |
| rem | rem.stype d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Compute the integer remainder of division. |
| ret | ret{.uni}; | Control Flow Instructions | PTX ISA 1.0 | sm_10 | Return execution to caller’s environment. |
| rsqrt | rsqrt.approx.f32 d, a; | Arithmetic | PTX ISA 1.0 | sm_10 | Fast hardware approximation of 1/sqrt(x). |
| rsqrt.approx.ftz.f64 | rsqrt.approx.ftz.f64 d, a; | Floating-Point Instructions | PTX ISA 4.0 | sm_20 | Compute a double-precision (.f64 ) approximation of the square root reciprocal of a value. The least significant 32 bits of the double-precision (.f64 |
| sad | sad.type d, a, b, c; | Integer Arithmetic Instructions | PTX ISA 1.0 | sm_10 | Adds the absolute value of a-b to c and writes the resulting value into d. |
| selp | selp.type d, a, b, p; | Comparison and Selection Instructions | PTX ISA 1.0 | sm_10 | Select between two operands based on a predicate, without branching. |
| set | set.CmpOp.dtype.stype d, a, b; | Comparison and Selection Instructions | PTX ISA 1.0 | sm_10 | Compare two operands and write a numeric (not predicate) 0/1 or all-ones/all-zeros result. |
| setmaxnreg | setmaxnreg.action.sync.aligned.u32 imm-reg-count; | Miscellaneous Instructions | PTX ISA 8.0 | sm_90a | setmaxnreg provides a hint to the system to update the maximum number of per-thread registers owned by the executing warp to the value specified by the imm-reg-count operand. |
| setp | setp.CmpOp.type p, a, b; | Comparison and Selection Instructions | PTX ISA 1.0 | sm_10 | Compare two operands and write the boolean result to a predicate register. |
| shf | shf.l.mode.b32 d, a, b, c; // left shift | Logic and Shift Instructions | PTX ISA 3.1 | sm_32 | Shift the 64-bit value formed by concatenating operands a and b left or right by the amount specified by the unsigned 32-bit value in c. |
| shfl | shfl.mode.b32 d[|p], a, b, c; | Data Movement and Conversion Instructions | PTX ISA 6.0 | sm_30 | Exchange a value directly between lanes of the same warp. |
| shl | shl.type d, a, b; | Logic and Shift Instructions | PTX ISA 1.0 | sm_10 | Shift bits left, filling with zero. |
| shr | shr.type d, a, b; | Logic and Shift Instructions | PTX ISA 1.0 | sm_10 | Shift bits right, arithmetic or logical depending on the operand's signedness. |
| sin | sin.approx.f32 d, a; | Arithmetic | PTX ISA 1.0 | sm_10 | Fast hardware approximation of sin(x). |
| slct | slct.dtype.s32 d, a, b, c; | Comparison and Selection Instructions | PTX ISA 1.0 | sm_13 | Conditional selection. |
| sqrt | sqrt.approx{.ftz}.f32 d, a; // fast, approximate square root | Floating-Point Instructions | PTX ISA 1.0 | sm_20 | Compute sqrt( a ) and store the result in d. |
| st | st.space.type [a], b; | Data Movement and Conversion Instructions | PTX ISA 1.0 | sm_10 | Store a register value into the specified state space. |
| st.async | st.async{.weak}{.ss}.completion_mechanism{.vec}.type [a], b, [mbar]; | Data Movement and Conversion Instructions | PTX ISA 8.1 | sm_90 | 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 | Data Movement and Conversion Instructions | PTX ISA 8.6 | sm_100 | st.bulk instruction initializes a region of shared memory starting from the location specified by destination address operand a. |
| stackrestore | stackrestore.type a; | Stack Manipulation Instructions | PTX ISA 7.3 | sm_52 | Sets the current stack pointer to source register a. |
| stacksave | stacksave.type d; | Stack Manipulation Instructions | PTX ISA 7.3 | sm_52 | Copies the current value of stack pointer into the destination register d. |
| stmatrix | stmatrix.sync.aligned.shape.num{.trans}{.ss}.type [p], r; | Warp Level Matrix Multiply-Accumulate Instructions | PTX ISA 7.8 | sm_90 | 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. |
| sub | sub.type d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Subtract the second operand from the first, with optional saturation for signed 32-bit integers. |
| sub.cc | sub.cc.type d, a, b; | Extended-Precision Integer Arithmetic Instructions | PTX ISA 1.2 | sm_20 | Performs integer subtraction and writes the borrow-out value into the condition code register. |
| subc | subc{.cc}.type d, a, b; | Extended-Precision Integer Arithmetic Instructions | PTX ISA 1.2 | sm_20 | Performs integer subtraction with borrow-in and optionally writes the borrow-out value into the condition code register. |
| suld | suld.b.geom{.cop}.vec.dtype.clamp d, [a, b]; // unformatted | Surface Instructions | PTX ISA 1.5 | sm_20 | suld.b.{1d,2d,3d} Load from surface memory using a surface coordinate vector. |
| suq | suq.query.b32 d, [a]; | Surface Instructions | PTX ISA 1.5 | sm_20 | Query an attribute of a surface. |
| sured | sured.b.op.geom.ctype.clamp [a,b],c; // byte addressing | Surface Instructions | PTX ISA 2.0 | sm_20 | Reduction to surface memory using a surface coordinate vector. |
| sust | sust.b.{1d,2d,3d}{.cop}.vec.ctype.clamp [a, b], c; // unformatted | Surface Instructions | PTX ISA 1.5 | sm_20 | sust.{1d,2d,3d} Store to surface memory using a surface coordinate vector. |
| szext | szext.mode.type d, a, b; | Integer Arithmetic Instructions | PTX ISA 7.6 | sm_70 | Sign-extends or zero-extends an N-bit value from operand a where N is specified in operand b. |
| tanh | tanh.approx.f32 d, a; | Half Precision Floating-Point Instructions | PTX ISA 7.0 | sm_75 | Take hyperbolic tangent value of a. |
| tcgen05.alloc | tcgen05.alloc.cta_group.sync.aligned{.shared::cta}.b32 [dst], nCols; | Tensor Memory Allocation and Management Instructions | PTX ISA 8.6 | sm_100a | tcgen05.alloc is a blocking instruction which dynamically allocates the specified number of columns in the Tensor Memory and writes the address of the allocated Tensor Memory into shared memory at the location specified by address operand dst. |
| tcgen05.commit | tcgen05.commit.cta_group.completion_mechanism{.shared::cluster}{.multicast}.b64 | TensorCore 5th Generation Family Instructions | PTX ISA 8.6 | sm_100a | The instruction tcgen05.commit is an asynchronous instruction which makes the mbarrier object, specified by the address operand mbar, track the completion of all the prior asynchronous tcgen05 operations, as listed in mbarrier based completion mechanism, initiated by the executing thread. |
| tcgen05.cp | tcgen05.cp.cta_group.shape{.multicast}{.dst_fmt.src_fmt} [taddr], s-desc; | Tensor Memory Data Movement Instructions | PTX ISA 8.6 | sm_100a | Instruction tcgen05.cp initiates an asynchronous copy operation from shared memory to the location specified by the address operand taddr in the Tensor Memory. |
| tcgen05.dealloc | tcgen05.dealloc.cta_group.sync.aligned.b32 taddr, nCols; | Tensor Memory Allocation and Management Instructions | PTX ISA 8.6 | sm_100a | tcgen05.dealloc is a blocking instruction which de-allocates the Tensor Memory specified by the Tensor Memory address taddr. The operand nCols specifies the number of columns to be de-allocated. |
| tcgen05.fence | tcgen05.fence::before_thread_sync ; | TensorCore 5th Generation Family Instructions | PTX ISA 8.6 | sm_100a | The instruction tcgen05.fence::before_thread_sync orders all the prior asynchronous tcgen05 operations with respect to the subsequent tcgen05 and the execution ordering operations. |
| tcgen05.ld | // Base load instruction: tcgen05.ld.sync.aligned.shape1.num{.pack}.b32 r, [taddr]; |
Tensor Memory and Register Load/Store Instructions | PTX ISA 8.6 | sm_100a | 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.mma | // 1. Floating-point type without block scaling: tcgen05.mma.cta_group.kind [d-tmem], a-desc, b-desc, idesc, { disable-output-lane }, enable-input-d {, scale-input-d}; |
TensorCore 5th Generation of MMA Instructions | PTX ISA 8.6 | sm_100a | Instruction tcgen05.mma is an asynchronous instruction which initiates an 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. |
| tcgen05.mma.sp | // 1. Floating-point type without block scaling: tcgen05.mma.sp.cta_group.kind [d-tmem], a-desc, b-desc, [sp-meta-tmem] , idesc, { disable-output-lane }, enable-input-d{, scale-input-d}; |
TensorCore 5th Generation of MMA Instructions | PTX ISA 8.6 | sm_100a | Instruction tcgen05.mma.sp is an asynchronous instruction which initiates an MxNxK matrix multiply and accumulate operation of the form D = A*B+D where the A matrix is Mx(K/2), the B matrix is KxN, and the D matrix is MxN. |
| tcgen05.mma.ws | // 1. Floating-point type without block scaling: tcgen05.mma.ws.cta_group::1.kind{.collector_usage} [d-tmem], a-desc, b-desc, idesc, enable-input-d {, zero-column-mask-desc }; |
TensorCore 5th Generation of MMA Instructions | PTX ISA 8.6 | sm_100a | Instruction tcgen05.mma.ws is an asynchronous instruction which initiates an 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. |
| tcgen05.mma.ws.sp | // 1. Floating-point type without block scaling: tcgen05.mma.ws.sp.cta_group::1.kind{.collector_usage} [d-tmem], a-desc, b-desc, [sp-meta-tmem] , idesc, enable-input-d {, zero-column-mask-desc}; |
TensorCore 5th Generation of MMA Instructions | PTX ISA 8.6 | sm_100a | Instruction tcgen05.mma.ws.sp is an asynchronous instruction which initiates an MxNxK matrix multiply and accumulate operation, D = A*B+D where the A |
| tcgen05.relinquish_alloc_permit | tcgen05.relinquish_alloc_permit.cta_group.sync.aligned; | Tensor Memory Allocation and Management Instructions | PTX ISA 8.6 | sm_100a | tcgen05.relinquish_alloc_permit specifies that the CTA of the executing thread is relinquishing the right to allocate Tensor Memory. |
| tcgen05.shift | tcgen05.shift.cta_group.down [taddr]; | Tensor Memory Data Movement Instructions | PTX ISA 8.6 | sm_100a | Instruction tcgen05.shift is an asynchronous instruction which initiates the shifting of 32-byte elements downwards across all the rows, except the last, by one row. |
| tcgen05.st | tcgen05.st.sync.aligned.shape1.num{.unpack}.b32 [taddr], r; | Tensor Memory and Register Load/Store Instructions | PTX ISA 8.6 | sm_100a | 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; | Tensor Memory and Register Load/Store Instructions | PTX ISA 8.6 | sm_100a | Instruction tcgen05.wait::st causes the executing thread to block until all prior tcgen05.st operations issued by the executing thread have completed. |
| tensormap.cp_fenceproxy | tensormap.cp_fenceproxy.cp_qualifiers.fence_qualifiers.sync.aligned [dst], [src], size; | Parallel Synchronization and Communication Instructions | PTX ISA 8.3 | sm_90 | The tensormap.cp_fenceproxy instructions perform the following operations in order: Copies data of size specified by the size argument, in bytes, from the location specified by the address operand… |
| tensormap.replace | tensormap.replace.mode.field1{.ss}.b1024.type [addr], new_val; | Data Movement and Conversion Instructions | PTX ISA 8.3 | sm_90a | 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. |
| testp | testp.op.type p, a; // result is .pred | Floating-Point Instructions | PTX ISA 2.0 | sm_20 | testp tests common properties of floating-point numbers and returns a predicate value of 1 if True and 0 if False. |
| tex | tex.geom.v4.dtype.ctype d, [a, c] {, e} {, f}; | Texture Instructions | PTX ISA 1.0 | sm_20 | tex.{1d,2d,3d} Texture lookup using a texture coordinate vector. |
| tld4 | tld4.comp.2d.v4.dtype.f32 d[|p], [a, c] {, e} {, f}; | Texture Instructions | PTX ISA 2.2 | sm_20 | Texture fetch of the 4-texel bilerp footprint using a texture coordinate vector. |
| trap | trap; | Miscellaneous Instructions | PTX ISA 1.0 | sm_10 | Abort execution and generate an interrupt to the host CPU. |
| txq | txq.tquery.b32 d, [a]; // texture attributes | Texture Instructions | PTX ISA 1.5 | sm_20 | Query an attribute of a texture or sampler. |
| vadd | // 32-bit scalar operation, with optional secondary operation vadd.dtype.atype.btype{.sat} d, a{.asel}, b{.bsel}; |
Scalar Video Instructions | PTX ISA 2.0 | sm_20 | Perform scalar arithmetic operation with optional saturate, and optional secondary arithmetic operation or subword data merge. |
| vadd2 | // SIMD instruction with secondary SIMD merge operation vadd2.dtype.atype.btype{.sat} d{.mask}, a{.asel}, b{.bsel}, c; |
SIMD Video Instructions | PTX ISA 3.0 | sm_30 | Two-way SIMD parallel arithmetic operation with secondary operation. |
| vadd4 | // SIMD instruction with secondary SIMD merge operation vadd4.dtype.atype.btype{.sat} d{.mask}, a{.asel}, b{.bsel}, c; |
SIMD Video Instructions | PTX ISA 3.0 | sm_30 | Four-way SIMD parallel arithmetic operation with secondary operation. |
| vmad | // 32-bit scalar operation vmad.dtype.atype.btype{.sat}{.scale} d, {-}a{.asel}, {-}b{.bsel}, {-}c; |
Scalar Video Instructions | PTX ISA 2.0 | sm_20 | Calculate (a*b) + c, with optional operand negates, plus one mode, and scaling. The source operands support optional negation with some restrictions. |
| vote | vote.mode.pred d, {!}a; | Parallel Synchronization and Communication Instructions | PTX ISA 6.0 | sm_12 | Combine a per-lane predicate across the warp using any/all/ballot reduction. |
| vset | // 32-bit scalar operation, with optional secondary operation vset.atype.btype.cmp d, a{.asel}, b{.bsel}; |
Scalar Video Instructions | PTX ISA 2.0 | sm_20 | Compare input values using specified comparison, with optional secondary arithmetic operation or subword data merge. |
| vset2 | // SIMD instruction with secondary SIMD merge operation vset2.atype.btype.cmp d{.mask}, a{.asel}, b{.bsel}, c; |
SIMD Video Instructions | PTX ISA 3.0 | sm_30 | Two-way SIMD parallel comparison with secondary operation. |
| vset4 | // SIMD instruction with secondary SIMD merge operation vset4.atype.btype.cmp d{.mask}, a{.asel}, b{.bsel}, c; |
SIMD Video Instructions | PTX ISA 3.0 | sm_30 | Four-way SIMD parallel comparison with secondary operation. |
| vshl | // 32-bit scalar operation, with optional secondary operation vshl.dtype.atype.u32{.sat}.mode d, a{.asel}, b{.bsel}; |
Scalar Video Instructions | PTX ISA 2.0 | sm_20 | vshl Shift a left by unsigned amount in b with optional saturate, and optional secondary arithmetic operation or subword data merge. |
| vshr | // 32-bit scalar operation, with optional secondary operation vshr.dtype.atype.u32{.sat}.mode d, a{.asel}, b{.bsel}; |
Scalar Video Instructions | PTX ISA 2.0 | sm_20 | vshl Shift a left by unsigned amount in b with optional saturate, and optional secondary arithmetic operation or subword data merge. |
| vsub | // 32-bit scalar operation, with optional secondary operation vsub.dtype.atype.btype{.sat} d, a{.asel}, b{.bsel}; |
Scalar Video Instructions | PTX ISA 2.0 | sm_20 | Perform scalar arithmetic operation with optional saturate, and optional secondary arithmetic operation or subword data merge. |
| vsub2 | // SIMD instruction with secondary SIMD merge operation vsub2.dtype.atype.btype{.sat} d{.mask}, a{.asel}, b{.bsel}, c; |
SIMD Video Instructions | PTX ISA 3.0 | sm_30 | Two-way SIMD parallel arithmetic operation with secondary operation. |
| vsub4 | // SIMD instruction with secondary SIMD merge operation vsub4.dtype.atype.btype{.sat} d{.mask}, a{.asel}, b{.bsel}, c; |
SIMD Video Instructions | PTX ISA 3.0 | sm_30 | Four-way SIMD parallel arithmetic operation with secondary operation. |
| wgmma.commit_group | wgmma.commit_group.sync.aligned; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.0 | sm_90a | 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; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.0 | sm_90a | 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; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.0 | sm_90a | 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; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.2 | sm_90a | 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; | Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions | PTX ISA 8.0 | sm_90a | 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. |
| wmma | wmma.load.a.sync.aligned.layout.shape.type r, [p]; | Warp Level Matrix Multiply-Accumulate Instructions | PTX ISA 6.0 | sm_70 | Higher-level warp matrix-multiply-accumulate built from explicit load/mma/store steps. |
| xor | xor.type d, a, b; | Logic and Shift Instructions | PTX ISA 1.0 | sm_10 | Bitwise XOR of two operands. |
Source
Normalized from NVIDIA's official Parallel Thread Execution ISA documentation. docs.nvidia.com/cuda/parallel-thread-execution ↗