NVIDIA PTX Instructions GPU Virtual ISA NVIDIA
38 instruction families — 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; | Warp-Level | 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. |
| and | and.type d, a, b; | Logic & Shift | PTX ISA 1.0 | sm_10 | Bitwise AND of two operands. |
| atom | atom.space.op.type d, [a], b; | Atomic Memory | PTX ISA 1.0 | sm_11 | Atomically read-modify-write a memory location and return the prior value. |
| barrier | bar.sync a{, b}; | Synchronization | PTX ISA 1.0 | sm_10 | Block threads in a CTA at a named barrier until the expected number of threads has arrived. |
| clz | clz.type d, a; | Bit Manipulation | PTX ISA 2.0 | sm_20 | Count the number of leading zero bits in an integer operand. |
| cvt | cvt.dtype.atype d, a; | Data Movement | PTX ISA 1.0 | sm_10 | Convert a value between integer and/or floating-point types with an explicit rounding mode. |
| div | div.stype d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Divide the first operand by the second. |
| 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. |
| ld | ld.space.type d, [a]; | Memory | PTX ISA 1.0 | sm_10 | Load a value from the specified state space into a register. |
| lg2 | lg2.approx.f32 d, a; | Arithmetic | PTX ISA 1.0 | sm_10 | Fast hardware approximation of log2(x). |
| 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). |
| max | max.type d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Select the larger of two operands. |
| membar | membar.level; | Synchronization | PTX ISA 2.0 | 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; | Matrix Operations | PTX ISA 6.4 | sm_70 | Cooperative, warp-wide matrix-multiply-accumulate executed on tensor-core hardware. |
| mov | mov.type d, a; | Data Movement | PTX ISA 1.0 | sm_10 | Copy a value into a register, or materialize an address/immediate. |
| 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. |
| 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 & Shift | PTX ISA 1.0 | sm_10 | Bitwise complement of an operand. |
| or | or.type d, a, b; | Logic & Shift | PTX ISA 1.0 | sm_10 | Bitwise OR of two operands. |
| popc | popc.type d, a; | Bit Manipulation | PTX ISA 2.0 | sm_20 | Count the number of set bits in an integer operand. |
| red | red.space.op.type [a], b; | Atomic Memory | PTX ISA 1.0 | sm_11 | Atomically read-modify-write a memory location without returning the prior value. |
| rem | rem.stype d, a, b; | Arithmetic | PTX ISA 1.0 | sm_10 | Compute the integer remainder of division. |
| rsqrt | rsqrt.approx.f32 d, a; | Arithmetic | PTX ISA 1.0 | sm_10 | Fast hardware approximation of 1/sqrt(x). |
| selp | selp.type d, a, b, p; | Comparison & Selection | PTX ISA 1.2 | sm_10 | Select between two operands based on a predicate, without branching. |
| set | set.CmpOp.dtype.stype d, a, b; | Comparison & Selection | PTX ISA 1.0 | sm_10 | Compare two operands and write a numeric (not predicate) 0/1 or all-ones/all-zeros result. |
| setp | setp.CmpOp.type p, a, b; | Comparison & Selection | PTX ISA 1.0 | sm_10 | Compare two operands and write the boolean result to a predicate register. |
| shfl | shfl.mode.b32 d[|p], a, b, c; | Warp-Level | PTX ISA 3.0 | sm_30 | Exchange a value directly between lanes of the same warp. |
| shl | shl.type d, a, b; | Logic & Shift | PTX ISA 1.0 | sm_10 | Shift bits left, filling with zero. |
| shr | shr.type d, a, b; | Logic & Shift | 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). |
| st | st.space.type [a], b; | Memory | PTX ISA 1.0 | sm_10 | Store a register value into the specified 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. |
| vote | vote.mode.pred d, {!}a; | Warp-Level | PTX ISA 1.2 | sm_12 | Combine a per-lane predicate across the warp using any/all/ballot reduction. |
| wmma | wmma.load.a.sync.aligned.layout.shape.type r, [p]; | Matrix Operations | 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 & Shift | 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 ↗