Tensor Core Intrinsics

All addresses in this page apply to ptxas v13.0.88 (CUDA 13.0). Other versions will differ.

ptxas supports five generations of tensor core operations spanning SM 70 through SM 100. The binary contains three major codegen handlers -- sub_5C7A50 (173KB, WMMA), sub_5C10A0 (120KB, MMA), and sub_5BBC30 (90KB, tcgen05.mma) -- plus four WGMMA handlers, eleven tcgen05 instruction handlers, and ~400 numeric MMA hash table entries. Together these constitute ~500KB of code generation logic dedicated to tensor core instructions, making this the single largest functional subsystem in ptxas.

Tensor Core Generations

ptxas tracks five distinct tensor core generations, each introducing new SASS opcodes, data types, and matrix shapes. The internal scheduling counters (visible in DUMPIR statistics output) reveal the SASS-level taxonomy.

Scheduling Unit Names

The binary's statistics printer functions (clones at 0x700-byte intervals from sub_ABBA50) emit per-unit throughput counters using the internal SASS operation classification:

Format strings from the binary:

# [est hmma1688=%d] [est hmma1688f16=%d] [est hmmaSp1688=%d] [est hmmaSp1688f16=%d]
# [est hmma16816=%d] [est hmma16816f16=%d]
# [est imma16816=%d] [est imma16832=%d] [est immaSp8832=%d] [est immaSp16832=%d]
# [est dmma=%d] [est fma64=%d]
# [hmma1688 thru=%f] [hmma1688f16 thru=%f] [hmmaSp1688 thru=%f] [hmmaSp1688f16 thru=%f]
# [hmma16816 thru=%f] [hmma16816f16 thru=%f]
# [imma16816 thru=%f] [imma16832 thru=%f] [immaSp8832 thru=%f] [immaSp16832 thru=%f]
# [dmma thru=%f] [fma64 thru=%f]

These counters are emitted by the post-scheduling statistics pass. The scheduler treats each counter as a separate functional unit throughput class, with dedicated latency table entries:

PTX-Level Instruction Lowering

WMMA Instructions

The WMMA (Warp Matrix Multiply-Accumulate) API is the oldest tensor core interface. Five PTX instructions are registered in the opcode dispatch table at sub_5D4190:

All five handlers allocate a 50,000-byte code generation buffer. The load/store handlers are ~8--10KB each and cover the combinatorial product of shape, layout, data type, and address space. The wmma.mma handler at 173KB is the largest single codegen handler in ptxas.

Instruction property accessors used by WMMA codegen:

WMMA Shapes and Types

sm_70 (Volta/Turing) -- 1st generation:

sm_72 (Turing) -- integer WMMA extension:

sm_80 (Ampere) -- 2nd generation extensions:

All WMMA shapes support three address spaces (generic, global, shared) and optional descriptor-based addressing (_desc variants). The binary contains separate intrinsic registrations for each combination, with the full prototype available in the string table.

MMA (mma.sync) Instructions

The mma.sync API uses a single PTX opcode mma dispatched to sub_5C10A0 (120KB). Unlike WMMA, it uses asymmetric matrix shapes with M and N decoupled from K, and operates at a single-warp granularity.

The numeric MMA hash table at a1+816 collapses the full variant space (shape + type + layout) into ~400 hash entries. Each entry maps a numeric string key (e.g., "2644314910") to a specific codegen handler function pointer, avoiding multi-dimensional dispatch.

sm_80+ MMA intrinsics (IDs 0x209--0x22F):

The 39 intrinsics registered under __cuda_sm_8x_mma_* cover:

Prototype examples from the binary:

.weak .func (.param .align 16 .b32 mma_dst[4])
    __cuda_sm_8x_mma_row_col_f16_f16_f16_f16
    (.reg .b32 a0, .reg .b32 a1, .reg .b32 b0, .reg .b32 b1,
     .reg .b32 c0, .reg .b32 c1, .reg .b32 c2, .reg .b32 c3);

.weak .func (.param .align 16 .b32 mma_dst[8])
    __cuda_sm_8x_mma_row_col_f32_f16_f16_f32
    (.reg .b32 a0, .reg .b32 a1, .reg .b32 b0, .reg .b32 b1,
     .reg .b32 c0, .reg .b32 c1, .reg .b32 c2, .reg .b32 c3,
     .reg .b32 c4, .reg .b32 c5, .reg .b32 c6, .reg .b32 c7);

Note the .param .align 16 .b32 mma_dst[N] return convention -- MMA results are returned through aligned parameter space, not registers, because the warp-cooperative nature of the operation means each thread holds only a fragment.

MMA Shape Summary Across Generations

MMA Validation

Three validator functions gate MMA features by SM version:

sub_4C2FD0 (12.2KB) -- WMMA/MMA master validator:

Performs three-way version checks:

• SM 75: base WMMA (f16)
• SM 80: extended types (BF16, TF32, FP64 -- "MMA with double types")
• SM 90: WGMMA features
• FP8: "mma with FP8 floating point type" (gated by sm_89+)

sub_49BBA0 (11.4KB) -- MMA type/scale validator:

Validates FP8 and block-scale configurations:

• "mma with FP8 floating point type" -- sm_89+ gate
• "mma with FP8 floating point type and FP16 accumulation" -- additional FP16 accum check
• "mma with FP8 floating point type and .m16n8k16 shape" -- shape/type cross-validation
• "Sparse mma with block scale" -- block-scale + sparsity interaction
• .block_scale modifier validation

sub_4BFED0 (10.3KB) -- WMMA shape validator:

Validates WMMA-specific shapes and the .aligned modifier:

• ".aligned modifier for wmma" -- alignment enforcement
• SM 75/80 version checks for shape legality

sub_490F90 -- Integer MMA validator:

Checks integer MMA shape validity: "Integer MMA with shape " -- validates m/n/k dimensions against the SM-level capability set.

sub_494210 (2.3KB) -- Sparse GMMA validator:

Validates sparse MMA metadata: "Sparse GMMA with " -- checks 2:4 sparsity pattern encoding.

sub_495900 -- WMMA floating-point validator:

Checks: "'wmma.mma' with floating point type" -- validates FP type compatibility with the target shape.

sub_4428E0 -- FP64 MMA validator:

Validates: "mma with .f64 type" -- gates double-precision MMA on sm_80+.

sm_90+ Sub-Byte MMA Intrinsics (IDs 0x23A--0x25F)

Hopper introduces 38 sub-byte MMA intrinsics covering s4/u4 sparse operations at warp granularity. These are distinct from the WGMMA API and provide backward-compatible sub-byte operations through the classical mma.sync interface.

Dense sub-byte (m8n8k32, m16n8k32, m16n8k64):

Sparse sub-byte (m16n8k64, m16n8k128):

The _0 and _1 suffixes on sparse m16n8k64 represent the two halves of a split operation -- the K dimension is decomposed into two steps for the sparsity pattern. The sparse variants take an additional e (metadata) operand encoding the 2:4 sparsity pattern.

Bit-operations (b1):

Prototype example (sparse m16n8k128):

.weak .func (.param .align 16 .b32 mma_dst[4])
    __cuda_sm_9x_mma_sub_byte_internal_sparse_m16n8k128_s4_s4
    (.reg .b32 a0, .reg .b32 a1, .reg .b32 a2, .reg .b32 a3,
     .reg .b32 b0, .reg .b32 b1, .reg .b32 b2, .reg .b32 b3,
     .reg .b32 c0, .reg .b32 c1, .reg .b32 c2, .reg .b32 c3,
     .reg .b32 e);

The sparse m16n8k128 variant takes 4 A-operands, 4 B-operands (double-width due to sparsity encoding), 4 C-operands, and 1 metadata operand e.

sm_100 Blackwell MMA (IDs 0x230--0x239)

Blackwell introduces 10 intrinsics for extended MMA operations:

HMMA/IMMA metadata helpers (__cuda_sm_10x_*_mdata_*):

These mdata functions compute sparse metadata for Blackwell's 5th-gen tensor cores. The e parameter is the sparsity selector, f_temp is a scratch register. The return array includes the transformed A-operands plus the computed metadata word.

Bit MMA (AND + XOR for sm_100):

Blackwell adds the AND reduction mode for 1-bit MMA (sm_90 only had XOR).

WGMMA -- Warpgroup MMA (SM 90+)

WGMMA (Warp Group Matrix Multiply-Accumulate) operates at warpgroup granularity (4 warps, 128 threads) and uses an asynchronous pipeline protocol. Four PTX instructions are registered:

WGMMA Pipeline Protocol

The hardware requires strict sequencing:

wgmma.fence              -- open pipeline stage
wgmma.mma_async  (1..N)  -- asynchronous MMA operations sharing accumulators
wgmma.commit_group       -- close pipeline stage
wgmma.wait_group N       -- wait for N outstanding groups to complete

Between fence and wait, strict register constraints apply:

1. No non-WGMMA definitions of accumulator registers
2. No non-WGMMA reads of accumulator registers
3. No non-WGMMA definitions of WGMMA input registers (including descriptors)

Violation of any constraint triggers pipeline serialization via sub_AE47B0, which collapses the pipeline to individual fence/mma/commit/wait per operation. The serialization reason is reported through warning codes 0x1D55--0x1D5E (see GMMA Pipeline).

WGMMA Descriptor Format

The wgmma.mma_async handler (sub_50AC70, 1,282 bytes) encodes the operation's matrix dimensions, data types, layout, and scale factors into the instruction. The A operand can be either a register operand or a descriptor -- a 64-bit value encoding the matrix base address, leading dimension, stride, and swizzle pattern. The B operand is always descriptor-based.

The descriptor format allows the hardware to fetch matrix data directly from shared memory via the TMA (Tensor Memory Accelerator), bypassing register file involvement for the B matrix operand entirely.

Ori Internal Encoding

Compiler-Inserted Warpgroup Instructions

The GMMA pipeline passes (phases 85 and 87) insert three compiler-internal pseudo-operations prefixed with underscore:

These are not directly written by the programmer. The compiler inserts them to manage register ownership transfer between the warpgroup's register file and the tensor core's accumulator pipeline.

TCGen05 -- 5th Generation Tensor Cores (SM 100+)

Blackwell introduces TCGen05 (Tensor Core Generation 5), which operates through tensor memory (TMEM) -- a dedicated on-chip memory visible only to the tensor core engine, separate from the register file and shared memory. Eleven PTX instructions are registered:

TCGen05 MMA Codegen

The tcgen05.mma handler (sub_5BBC30, 90KB) is the third-largest codegen handler in ptxas. It:

1. Allocates a 50,000-byte code generation buffer
2. Validates tcgen05 capability via sub_70FA00(*, 29)
3. Handles standard, sparse (.sp), and warp-shared (.ws) variants
4. Extracts sparse metadata via sub_70F0A0
5. Generates TMEM address computation code

The tcgen05.mma.ws formatter (sub_58FA20, 4,604B, also used for tcgen05.shift) handles the warp-shared variant where multiple warps contribute to a single MMA operation.

TCGen05 SASS Encoding

At the SASS level, TCGen05 operations are encoded by four specialized Mercury encoder functions:

The SASS encoder at sub_6D4350 references the tcmma operation namespace and validates block-scale configurations:

• "tcmma_*_o must be specified with blockscale" -- output operand requires block-scale modifier
• "uri width for tcmma_*_o must be 2" -- output URI width constraint
• "tcmma_*_q with blockscale must have uri width of 2" -- scale factor operand constraint
• "tcmma_*_mxq must be specified with blockscale" -- MX quantization operand constraint
• "For UTCHMMA, #scaleU4 must be 0 in SPA 10.1." -- SM 100 vs 103 compatibility

The string "UTCHMMA" (Unified Tensor Core HMMA) and "tcmma" (Tensor Core MMA) are the internal SASS-level names for Blackwell's tensor core operations.

TCGen05 Guardrails

Blackwell includes a debug/validation mode activated by --g-tensor-memory-access-check. When enabled, ptxas wraps TMEM operations with guardrail trap functions. Ten guardrail intrinsics are registered (IDs 0x20--0x2A):

Eight guardrail PTX instructions are registered in the opcode dispatch table:

The guardrail check functions are .FORCE_INLINE and return a boolean retVal:

.FORCE_INLINE .func (.reg .b32 retVal)
    __cuda_sm10x_tcgen05_guardrails_check_datapath_alignment
    (.reg .u32 tmemAddr, .reg .u32 iDesc, .reg .u32 cta_group,
     .reg .u32 hasWS, .reg .u32 hasSP, .reg .u32 matrix_kind);

The parameters reveal TMEM addressing structure: tmemAddr (TMEM base address), iDesc (instruction descriptor), cta_group (CTA group for cluster operations), hasWS (warp-shared flag), hasSP (sparse flag), matrix_kind (operand role).

Block-Scale MMA

Block-scale MMA allows per-block scaling factors for mixed-precision computation. In ptxas, this is gated by the .block_scale modifier on the PTX mma instruction:

• Validator sub_49BBA0 checks ".block_scale" and "Sparse mma with block scale" (sparsity + block-scale interaction)
• Additional intrinsic suffix __cuda_sm_100_tcgen05_ld_immhalfSplitOff and variants handle block-scale-aware loads
• The bf16x2.ue8m0x2 type string in the binary indicates UE8M0 (unsigned exponent-only) scale factor format for MX (microscaling) quantization

Intrinsic Registration Summary

Full Tensor Core Intrinsic ID Map

TCGen05 Intrinsics (Not in Master ID Table)

TCGen05 operations are dispatched through the named opcode table, not the numeric ID table. The 11 tcgen05.* instructions are registered directly in the opcode hash map at a1+808:

OCG-Level MMA Operations

The OCG system at sub_6C9EB0 registers additional SASS-level MMA operations for SM100+:

tcmma (2x64dp128bitlw02lw13, 2x64dp128bitlw01lw23, 4x32dp128bit)
tcbar, mmareadshma
16dp32bitt0t15, 16dp32bitt16t31
sparsify, spfactor2to4
tcshift, tcatomsws, tcldsws, tcstsws

These are internal Mercury-level operations that do not correspond 1:1 to PTX instructions. The tcmma variants encode the specific SASS datapath configuration: 2x64dp128bit means 2 datapaths, 64-element wide, 128-bit lane width.

Data Type Matrix

Complete data type support across tensor core generations as visible in the binary:

ELF Metadata

The cubin output includes tensor-core-specific EIATTR attributes:

Knobs and Options

Key Function Table

Cross-References

• Intrinsic Table Architecture -- Master registration, ID ranges, opcode dispatch
• GMMA/WGMMA Pipeline -- Phases 85/87, pipeline constraints, serialization
• Ada & Hopper Targets -- SM 89/90 feature gates, WGMMA details
• Turing & Ampere Targets -- SM 75--88 tensor core introduction
• Latency Model -- HMMA/IMMA/DMMA functional unit scheduling
• Register Model -- reg_type 6 (tensor/accumulator, allocator class 6)
• Mercury Encoder -- SASS encoding of MMA instructions
• ELF Output -- EIATTR_SPARSE_MMA_MASK, EIATTR_TCGEN05_*