NVPTX Machine Opcode Reference

This page is the master reference for NVPTX MachineInstr opcodes as they exist in cicc v13.0. These are the target-specific opcode numbers assigned during instruction selection and consumed by register allocation, instruction scheduling, the AsmPrinter, and every other machine-level pass. They are distinct from both LLVM IR opcodes (which live in the Instruction hierarchy) and from ISD/NVPTXISD SelectionDAG node opcodes (which exist only during lowering and are erased by ISel). A MachineInstr's opcode field is the 16-bit value at MachineInstr offset +68, and it indexes into the MCInstrDesc table to obtain operand counts, constraint classes, implicit defs/uses, and scheduling information.

Opcode Numbering Scheme

Opcodes 0--approximately 430 correspond to generic LLVM TargetOpcode values and standard LLVM machine pseudo-instructions (COPY, PHI, IMPLICIT_DEF, INLINEASM, etc.). These are identical to upstream LLVM 20.0.0. NVPTX target-specific opcodes begin around opcode 440 and extend into the thousands. The highest confirmed opcode numbers are in the 4900+ range (tcgen05 tensor core instructions for Blackwell).

The opcode numbering is generated by TableGen from the NVPTX .td instruction definitions and compiled into the MCInstrDesc table. Since cicc is a stripped binary, the symbolic names are lost. The identifications below come from behavioral analysis: matching the constraint table patterns, AsmPrinter string emission, and SelectionDAG lowering code against known PTX instruction semantics.

The Constraint Table: word_3F3E6C0

Every NVPTX machine opcode has an entry in the global constraint table at word_3F3E6C0. This is a flat array of 16-bit words, indexed by (opcode - 1). Each word packs two fields:

The access pattern, decompiled from sub_B612D0:

uint16_t entry = word_3F3E6C0[opcode - 1];
uint8_t constraint_class = entry & 0xFF;         // low byte
uint8_t register_class   = (entry >> 8) & 0xFF;  // high byte

switch (constraint_class) {
    case 0x00: ...  // simple 2-input ALU
    case 0x01: ...  // 3-input FMA
    ...
    case 0xB2: ...  // maximum observed class
}

The constraint class determines how many operands the instruction has, what register class each operand belongs to, and which operands are tied. Each case in the switch constructs a stack-allocated array of 16-byte constraint descriptors (see Pattern Database for the full descriptor layout) and calls sub_A78010 to emit them.

179 Constraint Classes

The constraint classes range from 0x00 through 0xB2 (179 values). Each class represents a distinct operand signature. Representative patterns:

The maximum observed operand count is 17 (constraint class 0xB0, associated with opcode 176), requiring 18 descriptor entries (17 inputs + 1 output) and 288 bytes of stack space in the constraint emitter's frame.

Register Class IDs in the High Byte

The high byte of each word_3F3E6C0 entry identifies the register class for the instruction's result. These IDs map to NVPTX's typed virtual register files:

Additional register class IDs observed in the constraint table (24, 27, 29, 32, 36, 39, 41, 67, 72, 76) likely correspond to sub-classes or aliased classes (e.g., Int32HalfRegs with ID related to 32 and prefix %hh), but their exact mappings have not been recovered. Instructions that produce no register result (stores, barriers, calls) have a zero or don't-care value in the high byte.

Identified Opcode Families

The following sections catalog every opcode range where the binary-to-PTX mapping has been confirmed. Opcodes are grouped by functional family. Where an opcode's identity is uncertain, it is marked with a question mark.

Copy and Move Family (440--503)

These are the NVPTX-specific copy instructions that the NVPTX register coalescer at sub_34AF4A0 processes. The standard LLVM RegisterCoalescer handles only the generic COPY pseudo (a generic TargetOpcode, not in this range); the NVPTX coalescer handles these target-specific copy families in a second pass.

The mapping function sub_3494EA0 contains a switch statement that classifies internal opcode IDs (1--0x12) into copy families:

The byte_444C4A0 operand-type classification table (16-byte entries, indexed by MVT enum) feeds the coalescer's type check:

struct OperandTypeEntry {    // 16 bytes at byte_444C4A0[16 * mvt - 16]
    uint8_t type_code;       // +0: 12=i32, 13=i64, 15=f32, etc.
    uint8_t size_class;      // +1: size in register-width units
    uint8_t register_bank;   // +2: bank identifier
    uint8_t constraint_flags; // +3: bit 0x10 = participates in coalescing
    uint8_t reserved[12];    // +4: padding
};

The constraint flag at offset +3 (mask 0x10) gates whether an operand participates in coalescing. Operands without this bit set (e.g., SpecialRegs) are excluded from the coalescer's worklist entirely.

Call ABI Family (505--573)

These opcodes implement the PTX .param-space calling convention. They are emitted by NVPTXTargetLowering::LowerCall (sub_3040BF0, 88KB) and form the backbone of every device-function call sequence.

The call sequence follows a strict emission order:

CallSeqBegin(315)
  for each argument:
    DeclareParam(505) or DeclareScalarParam(506)
    StoreV1/V2/V4(571/572/573) — store argument values
  DeclareRetParam(507) or DeclareRetScalarParam(508)  [if callee returns]
  CallProto(518)
  CallStart(514)                                       [actual call point]
  for each return value:
    LoadRetParam(515) or LoadRetParamLast(516)
  CallSeqEnd(517)
  DeclareRetParam_Ext(521)                             [if prototype present]
CallSeqEnd_Outer(316)

Vector Load/Store Family (568--573)

These opcodes handle vectorized .param-space data movement, emitted during argument passing and return value extraction:

The vector width selection logic in LowerCall (sub_3040BF0, lines 1429--1440):

accumulated_operand_count == 3  ->  StoreV1 (571), width=1
accumulated_operand_count == 4  ->  StoreV2 (572), width=2
accumulated_operand_count == 6  ->  StoreV4 (573), width=4
other                           ->  fatal error (unreachable)

The same pattern applies to LoadV1/V2/V4 on the return path. These opcodes are also used for by-value struct argument decomposition, where the struct is stored element-by-element into .param space using 8-byte chunks via StoreV1(571).

Atomic Family (294--317, 462)

Atomic opcodes are emitted by sub_20BED60 during DAG legalization and emitted as PTX by sub_21E5E70 (base) and sub_21E6420 (L2-hinted variant for SM 80+):

Within the PTX emission layer, the atomic operation is encoded in a packed operand word:

Note that operation codes 0x02 and 0x04 are absent -- there is no signed atomic add or a second OR variant, matching the PTX ISA specification.

On Ampere (SM 80+), each atomic operation has an L2 cache-hinted variant emitted by sub_21E6420. The PTX format becomes atom[.scope].op.L2::cache_hint.type, instructing the GPU to retain or evict data in L2 after the atomic completes.

Barrier and Fence Family (287--290)

The emission function sub_21E94F0 dispatches on the low 4 bits of the operand word. The fence.sc.cluster instruction requires SM 90 (Hopper) and provides sequentially-consistent fence semantics at cluster scope.

Cluster barrier instructions (SM 90+, emitted by sub_21E8EA0):

NVPTXISD Custom DAG Opcodes (22--499)

These are SelectionDAG-level opcodes used during lowering. After instruction selection, they are replaced by concrete MachineInstr opcodes. They are documented here because the DAG opcode numbers appear in the binary's lowering functions and serve as the conceptual identity of each instruction family:

The rounded math opcodes (245--274) follow a systematic pattern. The intrinsic lowering switch at sub_33B0210 maps NVVM intrinsic IDs to NVPTXISD opcodes:

MMA / Tensor Core Opcodes

Tensor core MachineInstr opcodes occupy a large range and are organized by generation. The central MMA instruction builder at sub_21E74C0 reads a packed 64-bit descriptor to determine the specific instruction variant.

Pre-Blackwell (SM 70--90) families:

Each family exists in two copies: the AsmPrinter-side at 0x21Dxxxx--0x21Exxxx and the NVPTX backend-side at 0x36Exxxx.

Blackwell tcgen05 (SM 100+):

Opcodes 4905--4940 cover 10 shape variants of tcgen05.mma. The packed descriptor encodes:

Modifiers include block_scale, weight_stationary, and scaleInputAccumulator. The architecture gate is subtarget+340 >= 0x3E8 (SM 100 decimal).

MMA Shape and Type Encoding

The MMA instruction builder uses enumerated shape and type codes embedded in the packed descriptor:

Shape codes (bits [39:32]):

Data type codes (in aty/bty fields):

Special Register Access

Special register read instructions map to PTX special registers. The AsmPrinter function sub_21E86B0 dispatches on a single-byte operand:

Cluster special registers (SM 90+, sub_21E9060) add 15 registers: %is_explicit_cluster, %cluster_ctarank, %cluster_nctarank, %cluster_ctaid.{x,y,z}, %cluster_nctaid.{x,y,z}, %clusterid.{x,y,z}, %nclusterid.{x,y,z}.

Address Space Conversion

The cvta instruction family is emitted by sub_21E7FE0:

Direction is determined by a separate operand: value 0 emits "a" (to-generic), value 1 emits "b" (to-specific).

Constraint Emission Pipeline

The full path from opcode to emitted constraint:

sub_B612D0(emitter_state, opcode):
    // Step 1: Table lookup
    entry = word_3F3E6C0[opcode - 1]
    reg_class = entry >> 8
    constraint_class = entry & 0xFF

    // Step 2: Build descriptor array on stack
    switch (constraint_class):
        case 0x00:
            // Simple 2-input ALU: {op0=RC, op1=RC, result=RC}
            desc[0] = {kind=0, value=sub_A778C0(state, reg_class, flags)}
            desc[1] = {kind=1, value=sub_A778C0(state, reg_class, flags)}
            desc[2] = {kind=-1, value=sub_B5BA00(state, reg_class)}
            sub_A78010(state, desc, 3)
        case 0x01:
            // Ternary FMA: {op0, op1, op2, result}
            desc[0..2] = three input constraints
            desc[3] = {kind=-1, value=sub_B5BA00(state, reg_class)}
            sub_A78010(state, desc, 4)
        ...
        case 0xB0:
            // 17-input complex: 17 input constraints + 1 output
            for i in 0..16:
                desc[i] = {kind=i, value=...}
            desc[17] = {kind=-1, value=sub_B5BA00(state, reg_class)}
            sub_A78010(state, desc, 18)

Key helper functions:

The constraint descriptors are purely stack-allocated within sub_B612D0's approximately 0x160-byte frame. No heap allocation occurs during constraint emission.

Complete Identified Opcode Summary

The following table consolidates every opcode where the binary-to-PTX mapping has been confirmed or strongly inferred. This represents a partial inventory -- the total opcode space extends to at least 4940, and many opcodes in the gaps (particularly in the load/store, texture, surface, and extended intrinsic ranges) remain unidentified.

Gaps and Unknown Ranges

The following opcode ranges are known to contain NVPTX instructions but have not been fully mapped:

Recovering these ranges requires systematic analysis of the sub_33B0210 intrinsic lowering switch (343KB, the single largest function in the binary) and correlation with the AsmPrinter's printInstruction dispatch table.

Function Map

Global Data References

Cross-References

• Pattern Database -- detailed constraint descriptor layout and emission sub-functions
• Register Coalescing -- the NVPTX-specific coalescer that processes copy family opcodes 440--503
• Code Generation -- pipeline overview including ISel, RA, and machine-level passes
• InstrEmitter -- how SDNodes become MachineInstrs with these opcodes
• Register Allocation -- greedy RA that consumes constraint table data
• AsmPrinter -- the PTX emission layer that converts these opcodes to text