TCGen05 -- 5th Generation Tensor Cores

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

TCGen05 is the Blackwell-generation tensor core instruction family introduced with SM 100. It replaces Hopper's WGMMA with a descriptor-based programming model that operates on Tensor Memory (TMEM) -- a dedicated register-file-like storage visible only to the tensor core. ptxas implements TCGen05 as 13 PTX instruction mnemonics (plus 8 debug guardrails), backed by a 90KB MMA codegen function, 11 SASS opcode groups (28 encoding variants), and a set of compiler-inserted validation hooks. TCGen05 is absent on sm_120/sm_121 (consumer Blackwell).

Architecture Overview

Descriptor-Based Model

TCGen05 abandons the register-operand model of previous tensor core generations (WMMA, HMMA, WGMMA) in favor of descriptors. The instruction descriptor (idesc) encodes the matrix operation configuration -- dimensions, data types, data path width, sparsity, and layout. The descriptor is passed as an operand to tcgen05.mma rather than encoded in the instruction mnemonic.

This design decouples the instruction encoding from the operation specification. Where WGMMA required hundreds of distinct intrinsic hash entries to cover every shape/type/layout combination, tcgen05 uses a single instruction with different descriptor values. The ~400 numeric MMA hash entries in the intrinsic dispatch table (at a1+816 in sub_5D4190) map WGMMA variants; tcgen05 replaces that complexity with descriptor-driven dispatch.

Tensor Memory (TMEM)

TMEM is a dedicated storage region private to the tensor core unit. It is not part of the general register file and is not directly addressable by non-tensor-core instructions. TMEM is organized into columns that are allocated, used, and deallocated explicitly by the programmer.

Key properties from binary analysis:

• Column-based allocation: tcgen05.alloc reserves columns; tcgen05.dealloc releases them
• Two CTA granularities: Operations execute at .cta_group::1 (single CTA) or .cta_group::2 (CTA pair) granularity. A function cannot mix both -- ptxas enforces: "Function '%s' uses single CTA(.cta_group::1) and CTA pair granularity(.cta_group::2) and that is not allowed."
• Allocation tracking: The compiler inserts reserved shared memory variables to track allocation state:
  • __nv_reservedSMEM_tcgen05_partition -- partition identifier
  • __nv_reservedSMEM_allocation_phase -- current allocation phase
  • __nv_reservedSMEM_allocation_mask -- bitmask of allocated columns
  • __nv_reservedSMEM_tmem_allocation_pipeline_mbarrier -- mbarrier for allocation pipeline
  • __nv_reservedSMEM_tmem_allocation_pipeline_mbarrier_parity -- parity tracking

TMEM Address Computation

Tensor memory addresses are computed through a standardized pattern visible in the TMEM address generator functions (sub_70E740, sub_70E940, sub_70EB00):

cvt.u32.u64 __cuda_sm_100_tcgen05_tmem_addr_base, %s;
add.u32 %s, __cuda_sm_100_tcgen05_tmem_addr_base, %s;

Five named TMEM address roles exist for MMA operations:

Constraint from the binary: "URa must be uint32 when URa is TMEM" -- uniform registers addressing TMEM must use 32-bit unsigned integers. When addressing a global descriptor: "URa must be uint64 when URa is GDESC".

PTX Instruction Set

Lifecycle Instructions

The alloc instruction has two CTA-granularity variants visible in the prototype strings:

• __cuda_sm10x_tcgen05_alloc_one_sm -- single-SM allocation (.cta_group::1)
• __cuda_sm10x_tcgen05_alloc_two_sm -- two-SM allocation (.cta_group::2)

Both take a destination pointer argument (__cuda_sm10x_tc_alloc_dst_ptr_arg) and a column count (__cuda_sm10x_tc_alloc_num_cols_arg).

Data Movement Instructions

Three intrinsic helper arrays support the ld/st/ld.red operations:

Each has a corresponding immhalfSplitOff parameter controlling split behavior:

• __cuda_sm_100_tcgen05_ld_immhalfSplitOff
• __cuda_sm_100_tcgen05_ld_red_immhalfSplitOff
• __cuda_sm_100_tcgen05_st_immhalfSplitOff

Synchronization Instructions

Compute Instructions

TCGen05.MMA -- Matrix Multiply-Accumulate

Codegen Function: sub_5BBC30 (90KB)

The largest per-instruction codegen function for TCGen05. Registered as the "tcgen05.mma" handler in sub_5D4190 (the intrinsic dispatch table builder). The function:

1. Allocates a 50,000-byte working buffer
2. Queries sub_70FA00(*, 29) to validate tcgen05 capability on the current target
3. Processes the instruction descriptor to determine operation parameters
4. Generates tensor memory addressing code for all operands (D, A, scaleA, scaleB, sparsity meta)
5. Emits the final MMA instruction encoding

MMA Modifiers

The binary reveals a rich set of MMA modifiers extracted by functions in the sub_70D1F0--sub_70D410 cluster:

Data Path Configurations

The MMA data path width determines the number of elements processed per cycle and the accumulator layout. Six configurations exist:

Constraint: "fused and l16dp32bit must be specified together" -- the fused mode requires the 16dp32bit data path.

Block Scaling (MX Format)

TCGen05 adds native block scaling support for microscaling (MX) format operations, visible through the tcmma prefix strings:

• "tcmma_*_o must be specified with blockscale" -- output modifier requires blockscale
• "uri width for tcmma_*_o must be 2" -- output uniform register index width must be 2
• "tcmma_*_q with blockscale must have uri width of 2" -- quantization with blockscale
• "tcmma_*_mxq must be specified with blockscale" -- MX quantization requires blockscale

Warp-Specialized MMA (.ws)

The .ws modifier enables warp-specialized execution where different warps in a warpgroup contribute to different phases of the MMA pipeline. Constraints from the binary:

• "When using buffer1-3, WS modifier must be specified" -- triple buffering requires .ws
• "ws opcode modifier not allowed with .2CTA" -- warp specialization is single-CTA only
• "ws opcode modifier not allowed with areuse or akeep" -- .ws incompatible with A-matrix reuse
• "ws opcode modifier not allowed with ashift" -- .ws incompatible with A-matrix shift

Triple-buffer register reuse strings for .ws mode:

Sparsity Support

TCGen05 supports structured sparsity through the sparsity metadata TMEM address (spMetaTmem). The _ashift modifier is constrained: "Ashift can only be specific when URa is in TMEM".

SASS Encoding

Opcode Map

TCGen05 SASS instructions span three opcode regions in the SM 100 SASS ISA. The encoding information comes from the latency model tables (sub_8E8A90 for sm_100) and the master instruction encoder (sub_6D9690, 94KB).

TMEM Operations (Opcodes 122--139)

TCGEN05 MMA/FENCE (Opcodes 213--221)

TCGEN05 Control (Opcodes 342--372)

28 encoding variants across 10 opcodes. These are the primary tensor core pipeline control instructions:

Opcode 344 (TCGEN05 execute) has the most variants (14), spanning encoding classes from F1F08 to F3008 with 2 to 7 operands. This is the actual MMA dispatch instruction -- the wide encoding range reflects the variety of descriptor configurations, operand modes, and data path widths.

Encoding Class Distribution

The encoding classes used by TCGen05 SASS instructions:

Latency Model

The sm_100 latency table (sub_8E8A90) uses a two-part structure: a 3.0KB base table covering standard instructions and a 949-byte supplement dedicated to TCGEN05 operations. The sm_120 consumer Blackwell table (sub_8E9000 + sub_8E92E0, 5.5KB) is the largest individual table and does not include TCGEN05 entries (confirming the feature's absence on consumer silicon).

CTA Granularity

TCGen05 operations specify whether they execute at single-CTA or CTA-pair granularity through the .cta_group modifier:

The compiler emits the appropriate EIATTR marker into the output cubin based on which granularity the kernel uses. The CUDA runtime uses this to configure the CTA launch parameters.

The binary enforces exclusivity: a single function cannot mix .cta_group::1 and .cta_group::2 operations. The error message is explicit: "Function '%s' uses single CTA(.cta_group::1) and CTA pair granularity(.cta_group::2) and that is not allowed."

ELF/Cubin Markers

EIATTR Entries

EICOMPAT Attributes

Entry Fragment Markers

TMEM usage per-CTA is recorded in entry fragment markers:

Guardrail Debug Instrumentation

When compiling with -g (debug mode), ptxas inserts runtime validation checks around tcgen05 operations. These are controlled by the --g-tensor-memory-access-check / --gno-tensor-memory-access-check CLI options.

Guardrail Check Functions

Eight _tcgen05.guardrails.* pseudo-instructions insert inline validation code:

Guardrail Trap Functions

When a guardrail check fails, it calls a trap function that reports the violation and terminates:

These are intrinsic IDs 0x20--0x2A (11 entries total including a mask creation helper) in the intrinsic table.

Guardrail Check Wrappers

The compiler also generates .FORCE_INLINE wrapper functions that combine multiple checks:

Bulk Copy Operations (cp.async.bulk.tensor)

TCGen05 is complemented by asynchronous bulk copy operations for loading data into tensor memory. These are registered as separate intrinsic IDs (0x2B--0x3C, 18 entries) under the __cuda_sm1xx_* naming convention:

The cp.async.bulk.tensor handler is 45KB and covers all dimensionality variants (1D through 5D), both tile and im2col access patterns, and unicast/multicast delivery modes.

SM Availability Gating

Capability Check

TCGen05 availability is gated by sub_70FA00(*, 29), which checks the target SM version. The check returns true for sm_100, sm_103, and sm_110 (and their a/f sub-variants) and false for sm_120/sm_121.

OCG Builtin Names

The OCG (Optimized Code Generation) layer uses short mnemonic names for tcgen05 operations visible in the builtin name lookup (sub_6C9EB0):

The .tcgen05op string identifies an Ori IR instruction as belonging to the tcgen05 family during the optimizer pipeline.

Version Compatibility

CUDA 12.x to 13.0 Breaking Change

ptxas v13.0.88 includes a linker-level version check for tcgen05 objects:

"Object '%s' cannot be linked due to version mismatch. Objects using tcgen05 in 12.x cannot be linked with 13.0 or later, they must be rebuilt with latest compiler"

The EICOMPAT_ATTR_INST_TCGEN05_MMA_DEPRECATED attribute tags objects compiled with the 12.x-era tcgen05 encoding, which is binary-incompatible with the 13.0 encoding. The SASS instruction encoding for tcgen05.mma changed between CUDA 12.x and 13.0 -- objects must be recompiled.

SM 100 vs SM 103 Differences

Both sm_100 and sm_103 share the same tcgen05 instruction set and codegen factory (36864). They share all 7 dispatch-table handler functions. The differences between sm_100 and sm_103 are:

• Different Handler A and Handler B capability accessor functions (sm_100: sub_609C30/sub_609BD0; sm_103: sub_608F20/sub_609D20)
• Different intrinsic table initializers (sm_100: sub_60A910; sm_103: sub_60A700)
• sm_103 may expose additional capability flags for GB300-specific features

Both targets produce identical SASS for tcgen05 instructions. The f sub-variants (sm_100f, sm_103f) allow cross-compilation within the family: sm_100f code can run on sm_103 hardware.

Compiler Pipeline

PTX Parsing and Validation

1. Lexer (sub_720F00, 64KB): Recognizes tcgen05.* tokens during lexical analysis
2. Validator (sub_4C5FB0, 28KB): Shared MMA/WMMA/tcgen05 validation function. Checks instruction legality for the current SM target, validates operand types, descriptor fields, and modifier combinations
3. Instruction table (sub_46E000, 93KB): Registers tcgen05 instruction variants with their type combinations (e.g., .tcgen05op)

Intrinsic Dispatch

The intrinsic dispatch table builder (sub_5D4190, 41KB) registers tcgen05 handlers:

Intrinsic Lowering

The TCGen05 MMA handler (sub_6D7AF0, 19KB) and validator (sub_6D69B0, 12KB) in the encoding zone handle the lowering from abstract intrinsic operations to concrete SASS encoding. The handler checks modifier consistency:

• "fused and l16dp32bit must be specified together"
• "Inputs vector length is inconsistent with layout and num modifiers"

TMEM Address Generation

The TMEM address generator cluster (sub_70E740, sub_70E940, sub_70EB00) generates PTX parameter passing code for tensor memory addresses:

st.param.b32 [%s + %d], %s;
ld.param.b32 %s, [%s + %d];

SASS Encoding

The master instruction encoder (sub_6D9690, 94KB) handles the final binary encoding. TCGen05 instructions use the Mercury encoding pipeline (encoder factory 36864) with Blackwell-specific opcode tables.

Function Map

Cross-References

• Blackwell (SM 100--121) -- Target-level architecture gating, codegen factory 36864
• SM Architecture Map -- Complete SM table, capability dispatch infrastructure
• GMMA/WGMMA Pipeline -- Predecessor tensor core pipeline (sm_90), same warpgroup execution model
• Intrinsic Table (608 Entries) -- IDs 0x20--0x3C (tcgen05 guardrails + bulk copy)
• Tensor Core Intrinsics -- WMMA/MMA/tcgen05 intrinsic lowering detail
• Late Expansion & Legalization -- tcgen05 guardrail insertion during late expansion
• SASS Instruction Encoding -- Mercury encoder, opcode tables
• Latency Model & HW Profiles -- SM 100 TCGEN05 supplement table
• SASS Text Generation -- TCGen05 instruction formatters
• CLI Options -- --g-tensor-memory-access-check option