Blackwell (SM 100–121)

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

ptxas v13.0.88 handles five Blackwell-era base targets — sm_100, sm_103, sm_110, sm_120, sm_121 — spanning datacenter, automotive, consumer, and DGX product lines.

• Shared identity. All share the internal profile code generation nibble 0x9 (stamped into profile+0x15c by each profile constructor) and the "Blackwell" family string internally.
• Distinct profile codes. The per-arch internal code in profile+0x15c is (9 << 12) | chip_ordinal: sm_100=0x9000, sm_101/sm_110=0x9001, sm_103=0x9003, sm_120=0x9004, sm_121=0x9005 (byte-verified at the movl $0x900N,0x15c(%rax) stamps 0x608E15/0x60AA45/0x60A725/0x60A935/0x60A505). This is not the legacy 16-bit hash-slot version-code (that table at 0x2020620 ends at 0x9005=sm_90f and structurally cannot encode sm_100+); the canonical user-visible code for these targets is the __CUDA_ARCH__ decimal (1000/1100/1030/1200/1210). See Version Codes.
• Defining feature. Capsule Mercury (capmerc) is the default binary output format, automatically enabled for SM numbers exceeding 99.
• tcgen05 split. sm_100, sm_103 and sm_110 (Jetson Thor) accept tcgen05 (5th-generation tensor cores with dedicated tensor memory) — empirically: ptxas v13.0.88 assembles a tcgen05.alloc for sm_100a/sm_103a/sm_110a; the consumer variants sm_120/sm_121 reject it (Instruction 'tcgen05.*' not supported).

SM Version Table

Profile-code encoding: (9 << 12) | chip_ordinal, stamped into profile+0x15c. sm_100 is ordinal 0 (generation base); sm_103 ordinal 3; sm_120 ordinal 4 (0x9004); sm_121 ordinal 5 (0x9005) — sm_120 and sm_121 are distinct (0x9004 vs 0x9005, byte-verified at the constructor stamps and confirmed by distinct e_flags SM bytes 0x78/0x79). sm_110 (Thor) reuses the sm_101 ordinal slot 0x9001. These internal codes are not the legacy hash-slot version code (which ends at sm_90f); the canonical code for these targets is the __CUDA_ARCH__ decimal.

Unreleased SM numbers referenced in the binary: The SASS formatter sub_583190 (rsqrt) checks for SM codes 102, 103, 107, 124, 130 in architecture-specific dispatch paths, suggesting internal/future variants beyond the five publicly exposed targets.

Sub-Variant System

Every Blackwell SM has three sub-variants. The base and a/f variants within an SM share all 7 dispatch table handler functions — they are identical silicon with different compatibility metadata and feature exposure.

Profile Object Fields

The profile constructor sub_6765E0 builds profile objects for each sub-variant with these fields:

CUDA_ARCH Macro Values

Suffix Bit Flags in Profile Objects

From the decompiled profile constructor, suffixed variants set specific byte flags:

The a flag is set on the SM and compute profile objects only. The f flag is set on all three (sm, compute, lto), reflecting the fact that f-compiled code must retain its feature-reduced metadata through linking.

isaClass Inheritance

Sub-variants reference their base SM's isaClass rather than defining a new one:

• sm_100a and sm_100f reference "(profile_sm_100)->isaClass"
• sm_103a and sm_103f reference "(profile_sm_103)->isaClass"
• sm_120a and sm_120f reference "(profile_sm_120)->isaClass"
• sm_121a and sm_121f reference "(profile_sm_121)->isaClass"

This confirms that sub-variants share the instruction set architecture class with their base. The a/f distinction is purely in compatibility metadata, not in the ISA or codegen.

Capability Dispatch

sub_607DB0 registers handler functions into 7 parallel hash maps. All sub-variants of a given SM register the same function pointers.

Handler Assignments (Maps 1–3)

Performance / Occupancy Handlers (Maps 6–7)

Every Blackwell SM has unique handler functions in all 7 maps. This contrasts with Hopper where sm_90 and sm_90a share all handlers. Each Blackwell SM variant is architecturally distinct enough to warrant separate capability accessors, performance models, and intrinsic tables.

Warp Geometry

The warp geometry initializer at sub_8E4400 switches on the internal profile code (profile+0x15c, value 0x900N) to select dispatch parameters. All Blackwell targets (code > 0x8fff) fall into the maximum bucket:

code >  0x8fff  -> 16 warps, 240 dispatch slots

This is identical to Hopper (sm_90). The 16-warp / 240-slot geometry supports Blackwell's warpgroup execution model (4 warps per warpgroup, 4 warpgroups per SM partition).

Sub-Architecture Variant Table

sub_8E4400 is a switch on the profile code — cmp $0x9000/$0x9003/$0x9004/$0x9005 each branch to a distinct target (disassembly at 0x8E4438–0x8E445C). sm_120 (0x9004) and sm_121 (0x9005) take different branches (je 0x8E4670 vs je 0x8E46A0); they do not share a sub-variant. sm_110/sm_101 (0x9001) is handled by the jbe-low path (0x8E4444), not a 0x9004 case:

These codes select different entries within the per-SM latency tables, allowing the scheduler to use silicon-specific pipeline timing.

Hardware Resource Geometry

Per-SM hardware resource limits used by ptxas for register allocation, occupancy calculations, and scheduling decisions. Extracted from sub_8688F0 (universal baseline), sub_8E4400 (scheduler partition geometry), and sub_ABF250 (occupancy calculator). See targets/index.md — Per-SM Resource Geometry Table for the complete table across all architectures.

Column definitions:

• Regs/SM: Total 32-bit registers per streaming multiprocessor. 65,536 universally for sm_75+.
• Max Regs/Thread: Maximum registers a single thread can use. 255 universally (sub_8688F0 offset +612).
• Max Threads/CTA: Maximum threads per cooperative thread array (block).
• Warps/SM: Total concurrent warps per SM. Determines peak occupancy.
• Max CTAs/SM: Maximum concurrent CTAs per SM.
• Sched Partitions / Dispatch Slots: From sub_8E4400 offset +18 (packed DWORD) and offset +22 (WORD).
• Configurable Shared Memory: Valid shared memory sizes per CTA, selected by cudaFuncSetAttribute.

All Blackwell targets share the 16-partition / 240-slot geometry (identical to Hopper sm_90). The a and f sub-variants within each SM share the same geometry — differentiation is in compatibility metadata and feature exposure, not in resource limits. The primary distinction across Blackwell SMs is in the latency tables and tcgen05 availability, not in the scheduling partition structure.

Capsule Mercury (capmerc) — Default Output Format

Capsule Mercury is automatically enabled for all Blackwell targets. When the SM architecture number exceeds 99, ptxas sets the capmerc flag at offset+81 in the compilation context. This applies to sm_100, sm_103, sm_110, sm_120, and sm_121 uniformly.

Three Output Modes

Capsule Mercury vs Mercury

Both modes use the same Mercury encoder pipeline (phases 117–122). The capmerc distinction is at the ELF emission level:

• Mercury produces a fully-resolved SASS binary in .text.<funcname> sections
• Capsule Mercury wraps Mercury-encoded instructions in .nv.capmerc<funcname> sections with a 328-byte capsule descriptor, plus .nv.merc.* debug/metadata sections

The capsule descriptor (constructed by sub_1C9C300, 24KB) contains the Mercury instruction stream, relocation metadata (R_MERCURY_* types), KNOBS compilation configuration snapshot, and function-level metadata (register counts, barriers, shared memory usage).

Opportunistic Finalization

Capsule Mercury enables deferred finalization — compiling once for one SM and reconstituting SASS for a different SM at link or load time.

The compatibility checker sub_60F290 determines whether a capmerc binary compiled for SM X can be finalized for SM Y. On success, ptxas emits: "applied for off-target %u -> %u finalization".

Self-Check Mechanism

The --self-check CLI option performs roundtrip verification:

1. Generate capmerc output (Mercury encoding + metadata)
2. Reconstitute SASS from the capmerc data
3. Compare section-by-section; report error codes 17 (content mismatch), 18 (count mismatch), 19 (metadata mismatch)

The reconstituted SASS can be dumped with --out-sass for debugging self-check failures.

TCGen05 — 5th Generation Tensor Cores

TCGen05 is the defining hardware feature of Blackwell datacenter parts. It introduces tensor memory (TMEM) as a dedicated register-like storage directly connected to the tensor core, eliminating the shared-memory bottleneck of previous WGMMA designs.

SM Availability

The tcgen05 ISA is gated by SM version checks (visible as sub_70FA00(*, 29) capability queries). sm_120 and sm_121 have no warpgroup or tensor-memory MMA at all — the tensor path is warp-level mma.sync:

• The HMMA/IMMA lineage from sm_70–sm_89.
• Extended with the block-scaled narrow-precision family (mma.sync.aligned.kind::f8f6f4 and kind::mxf4nvf4 for FP4/FP6/FP8).

CORRECTION (WGMMA-01) — earlier revisions of this page listed WGMMA (wgmma.mma_async) among the tensor paths inherited by Blackwell. That is wrong: WGMMA is Hopper-exclusive.

• ptxas v13.1.115 assembles wgmma.fence/wgmma.commit_group only for .target sm_90a; every Blackwell target — datacenter and consumer — rejects it with error: Instruction 'wgmma.fence' not supported on .target 'sm_100a' (and identically for sm_103a/sm_110a/sm_120a/sm_121a).
• Datacenter Blackwell replaced the warpgroup matmul with tcgen05; consumer Blackwell fell back to warp-level mma.sync.
• No software knob re-enables a hardware instruction the silicon lacks.

PTX Instructions

Registered in the opcode dispatch table at sub_5D4190:

The tcgen05.mma codegen handler at sub_5BBC30 is 90KB — the largest single-instruction handler in ptxas — reflecting the complexity of 5th-gen tensor core MMA with tensor memory operands, scale factors, sparsity, and accumulator management.

TCGen05 Guardrail Functions

Eight debug/validation functions provide runtime instrumentation for tensor memory operations when compiled with --g-tensor-memory-access-check (or -g-tmem-access-check):

The bounds checker (sub_70E0E0, 296 lines decompiled) generates inline PTX code to validate tensor memory column counts, extracting bitfield positions from tcgen05 descriptors (e.g., and.b32 %s, 0x7E0000, %s; shr.u32 %s, %s, 17; mul.lo.u32 %s, %s, 8).

TCGen05 Intrinsics

Additional tcgen05 helper intrinsics observed in decompiled code:

• __cuda_sm_100_tcgen05_ld_red_immhalfSplitOff — load-reduce with immediate half-split offset
• __cuda_sm_100_tcgen05_ld_immhalfSplitOff — load with immediate half-split offset
• __cuda_sm_100_tcgen05_st_immhalfSplitOff — store with immediate half-split offset
• __cuda_sm_100_tcgen05_ld_red_funcRetArr — load-reduce returning array
• __cuda_sm_100_tcgen05_ld_funcRetArr — load returning array

These helpers (decompiled in sub_70D910 and sub_70DDB0) generate inline PTX for complex tensor memory access patterns including array returns via ld.param.b32 sequences.

Bulk Copy Intrinsics (sm_1xx)

18 intrinsics in the __cuda_sm1xx_* namespace cover cp.async.bulk.tensor 1D–5D in tile and im2col modes. These extend the Hopper TMA infrastructure with Blackwell-specific enhancements.

SM 100 / SM 100a / SM 100f — Blackwell Datacenter

sm_100 is the reference Blackwell architecture. Internal profile code 0x9000 (profile+0x15c), CUDA_ARCH 1000, e_flags SM byte 0x64, min PTX .version 8.6.

Products

B100, B200 (datacenter GPU), paired as GB200 NVL72 superchips.

Key Features

• TCGen05: Full 5th-gen tensor core with TMEM (alloc, dealloc, ld, st, commit, cp, shift, mma, mma.ws)
• Capsule Mercury: Default output format (auto-enabled for SM > 99)
• No WGMMA: the Hopper wgmma.mma_async warpgroup matmul does not carry forward (ptxas rejects wgmma.* on sm_100a) — datacenter Blackwell's warpgroup-class matmul is tcgen05 instead
• Cluster operations: Thread-block clusters, distributed shared memory (from Hopper)
• setmaxnreg: Dynamic register allocation (from Hopper)
• Uniform register ALU: UFADD, UFFMA, UFSEL, UFSETP, UVIADDR (Blackwell uniform register ISA additions)

Handler Functions

HW Latency Table

sub_8E8A90 (3.0KB) — the base Blackwell latency table. Two-part structure: a 3.0KB base table for standard instructions plus a ~949-byte TCGEN05 supplement covering tensor core scheduling classes 745–772+.

Profile Object

From sub_6765E0:

SM name:       "sm_100"
Compute name:  "compute_100"
LTO name:      "lto_100"
Family:        "Blackwell"
CUDA_ARCH:     "-D__CUDA_ARCH__=1000"

The profile constructor stores dword_29FE2C4 = 100 after constructing all sm_100 sub-variants, likely recording the current highest-registered base SM number.

SM 103 / SM 103a / SM 103f — Blackwell Ultra

sm_103 is Blackwell Ultra, targeting the GB300 NVL72 platform. Internal profile code 0x9003 (profile+0x15c), CUDA_ARCH 1030, e_flags SM byte 0x67, min PTX .version 8.8.

Products

GB300 (datacenter, Blackwell Ultra). Incremental silicon revision over sm_100.

Differentiation from sm_100

The SASS formatter sub_583190 (rsqrt instruction) explicitly checks for "sm_103" and applies a Blackwell Ultra-specific operand layout. The formatter dispatch first tests raw SM codes (102, 103, 107, 130, 124), then uses check_target_sm(instr, 0, "sm_103") for string-based validation.

From sweep data on the SM103-specific encoding path:

• Sets encoding flags XOR 0x10, XOR 0x40 for sm_103-specific instruction variants
• New operand layout for transcendental instructions (rsqrt, likely also rcp/sqrt)

sm_103 has a separate 618-byte supplementary latency table — the smallest in the binary — suggesting minimal scheduling parameter changes from sm_100.

Handler Functions

All unique from sm_100:

Profile Object

SM name:       "sm_103"
Compute name:  "compute_103"
LTO name:      "lto_103"
Family:        "Blackwell"
CUDA_ARCH:     "-D__CUDA_ARCH__=1030"

SM 110 / SM 110a / SM 110f — Jetson Thor

sm_110 targets the Jetson Thor SoC for automotive and robotics applications. Internal profile code 0x9001 (profile+0x15c, shares the sm_101 chip-ordinal slot), CUDA_ARCH 1100, e_flags SM byte 0x6e. Minimum PTX .version 9.0 (the only modern family that requires PTX 9.0; verified by binary-searching -arch=sm_110 and by the base min-PTX table at 0x1D161C0).

Products

Jetson Thor (automotive-grade SoC with integrated GPU). Originally internally designated sm_101 before rename.

sm_101 Legacy Alias

sm_101 (with variants sm_101a and sm_101f) was the original internal name for Jetson Thor. It was renamed to sm_110 in a later CUDA release. sm_101 is not a usable target in this build — -arch=sm_101 (and sm_101a/sm_101f) is rejected by the gpu-name CLI parser with Value 'sm_101' is not defined for option 'gpu-name' (verified on v13.0.88 and v13.1.115). It survives only as an internal enum slot: it occupies the sm-id enumeration array, owns a register-class arch selector (0x5001, vtable thunk 0xb081c0), and keeps its {sm, major, minor} rows in the min-PTX .version tables:

(These are the three {sm, major, minor} min-PTX tables; each is a sorted run beginning at {90,8,0}/{100,8,…} and queried by bsearch. The sm_101 rows are real, but they are unreachable from the CLI because the gpu-name parser rejects the string before any min-PTX check runs.)

sub_6765E0 (the profile constructor) registers only sm_110 / sm_110a / sm_110f — there is no separate profile object for sm_101; the Mercury normalize rule folds sm_101 → sm_110. So sm_101 has no __CUDA_ARCH__ descriptor string and is a phantom slot.

The PTX ISA version is notable: the sm_101 rows carry PTX 8.6 (same as sm_100), while sm_110 requires PTX 9.0 ({110, 9, 0} in every table) — the only modern family gated to PTX 9.0. This reflects the timeline: sm_101 was named when the Jetson Thor target was first added alongside sm_100, before the sm_110 numbering and PTX 9.0 specification existed.

Key Characteristics

• Full tcgen05 hardware: despite being an automotive/robotics SoC (not a datacenter part), sm_110 carries the Blackwell tensor-memory + 5th-gen tensor-core path — tcgen05.* instructions assemble for sm_110a (verified), the same as sm_100/sm_103 and unlike consumer sm_120/sm_121
• Cluster operations inherited: thread-block clusters and distributed shared memory carry forward from Hopper; the WGMMA matmul does not (replaced by tcgen05, as on sm_100)
• SoC-specific constraints: Differentiated from sm_100 through capability flags, not through missing features — the capability accessor functions (sub_609F30, sub_608F50) return SoC-appropriate resource limits

Handler Functions

Profile Object

SM name:       "sm_110"
Compute name:  "compute_110"
LTO name:      "lto_110"
Family:        "Blackwell"
CUDA_ARCH:     "-D__CUDA_ARCH__=1100"

Note: sm_110 uses different xmmword constants for profile fields [5]--[7] compared to sm_100, visible in the profile constructor where v98[5] = xmmword_2027D70; v98[6] = v103 (from v209); v98[7] = v104 (from v207). This encodes different hardware resource parameters.

SM 120 / SM 120a / SM 120f — Blackwell Consumer

sm_120 targets consumer and enterprise workstation GPUs. Internal profile code 0x9004 (profile+0x15c), CUDA_ARCH 1200, e_flags SM byte 0x78, min PTX .version 8.7.

Products

RTX 5090, RTX 5080, RTX 5070 Ti, RTX 5070, RTX 5060 (consumer). RTX Blackwell Pro (enterprise workstation).

The tcgen05 Gap

sm_120 is architecturally distinct from sm_100 despite sharing the "Blackwell" family string. The critical difference: no tcgen05 support. The entire tensor memory subsystem (alloc, dealloc, ld, st, commit, cp, shift, mma) is absent. The tensor path is warp-level only:

• HMMA/IMMA inherited from sm_70–sm_89 — the warp-level mma.sync matmul (no warpgroup; accumulators live in RMEM, not TMEM)
• Block-scaled narrow precision — the sm_120-native addition: mma.sync.aligned.kind::f8f6f4.m16n8k32 (FP4 e2m1, FP6 e2m3/e3m2, FP8 e4m3/e5m2) and the scale-vector kind::mxf4nvf4 form

There is no WGMMA and no tcgen05 — both the Hopper warpgroup path and the datacenter tensor-memory path are absent. This is gated by SM version checks in the capability accessor functions; the intrinsic table initializer (sub_608DF0) registers neither tcgen05 nor wgmma.* handlers for sm_120.

GOTCHA — what sm_120 does carry forward from Hopper is the data-movement and scheduling model — TMA bulk-async tensor copies, thread-block clusters, and warp-specialized producer/consumer pipelines — not the matmul instruction. A kernel chasing MFU on sm_120 pairs warp-level mma.sync compute with TMA + warp-specialized async staging; it cannot issue wgmma.mma_async (ptxas v13.1.115 rejects it for sm_120a).

HW Latency Table

sm_120 has a distinct latency model, split into two parts:

The 5.5KB extended table is larger than any other individual latency table, suggesting that sm_120's consumer pipeline has significantly different scheduling characteristics from datacenter sm_100. The consumer pipeline likely has different functional unit counts, memory latencies, and tensor core throughput profiles.

Handler Functions

Profile Object

SM name:       "sm_120"
Compute name:  "compute_120"
LTO name:      "lto_120"
Family:        "Blackwell"
CUDA_ARCH:     "-D__CUDA_ARCH__=1200"

sm_120 uses a third distinct set of xmmword constants for profile fields, including xmmword_2027DC0 at field [6], confirming different hardware resource parameters from both sm_100 and sm_110.

SM 121 / SM 121a / SM 121f — DGX Spark

sm_121 targets the DGX Spark desktop AI workstation. Internal profile code 0x9005 (profile+0x15c), CUDA_ARCH 1210, e_flags SM byte 0x79, min PTX .version 8.8.

Products

NVIDIA DGX Spark (desktop AI workstation with Grace CPU + Blackwell GPU).

Relationship to sm_120

sm_121 has its own profile code 0x9005 (chip ordinal 5), distinct from sm_120's 0x9004 (ordinal 4) — they take separate branches in the sub_8E4400 switch (je 0x8E46A0 vs je 0x8E4670). But sm_121 inherits sm_120's xmmword profile constants and its SASS code generation is byte-identical to sm_120 (nvdisasm diff = 0; the two differ only in the e_flags SM byte 0x79 vs 0x78 and the .nv.compat CAN_FASTPATH_FINALIZE mask 0x00 vs 0x50). This suggests sm_121 is a binned or slightly modified sm_120 die, similar to how sm_86 relates to sm_80 in the Ampere generation.

Like sm_120, sm_121 has no tcgen05 support and no WGMMA — tensor operations use the warp-level mma.sync path (HMMA/IMMA plus block-scaled FP4/FP6/FP8 kind::f8f6f4/mxf4nvf4), identical to sm_120.

Handler Functions

All unique from sm_120:

Profile Object

SM name:       "sm_121"
Compute name:  "compute_121"
LTO name:      "lto_121"
Family:        "Blackwell"
CUDA_ARCH:     "-D__CUDA_ARCH__=1210"

Blackwell Uniform Register ISA

Blackwell extends the uniform register (UR) ISA introduced in Turing/Ampere with dedicated uniform ALU instructions:

These instructions execute on the uniform datapath (UDP, functional unit index 9), allowing floating-point uniform computations to stay in the UR file without round-tripping through the R file. Mercury encoding assigns major opcode 0x0E with 6 variants (sub_10C0550) for uniform ALU.

Architecture Version Threshold Checks

All Blackwell targets share profile code 0x9000+ (profile+0x15c ≥ 36864). The binary uses these thresholds to gate Blackwell-specific features:

The tcgen05 gating is not a simple threshold — it uses a per-SM capability query (sub_70FA00 with argument 29) that returns true for sm_100/103/110 and false for sm_120/121.

BB Initialization (Secondary Encoding)

The basic block initializer sub_6E8EB0 uses a secondary encoding space for Blackwell:

sub_6E8EB0 reads the register-class selector at profile+0xD4 and matches $0x5000 (sm_100) and $0x5004 (sm_120) exactly (disasm at 0x6E915E/0x6E9169). These are the same selector codes the register allocator dispatches on: 0x5000=sm_100, 0x5003=sm_103, 0x5004=sm_120, 0x5005=sm_121 (see Register Allocation). The XOR 0x10/0x40 flag toggles are applied on the sm_120 (consumer) path, not sm_103. The selector's generation nibble is 0x5 (5 << 12), separate from the profile code's generation nibble 0x9.

Scheduling Profile Differences

Blackwell targets share the 16-warp / 240-slot geometry with Hopper but have distinct latency tables:

The 5.5KB extended table for sm_120 is the largest individual latency table in the binary, reflecting the consumer microarchitecture's distinct pipeline design. The sm_100 table uses a supplement mechanism for TCGEN05-specific scheduling classes that the consumer sm_120 table does not need (since sm_120 lacks tcgen05).

Scheduling Class Assignments

From the opcode-to-scheduling-class mapper sub_89FBA0 (85KB), Blackwell-era opcodes use high-numbered scheduling classes:

Intrinsic Table

Blackwell intrinsic availability is cumulative — all sm_70, sm_80, sm_8x, and sm_9x intrinsics carry forward. Blackwell adds two new intrinsic groups:

sm_10x Intrinsics (21 entries)

sm_1xx Intrinsics (18 entries)

OCG (On-Chip Generated) Intrinsics

The OCG builtin name table at sub_6C9EB0 (13KB) contains the master list of Blackwell+ runtime-generated intrinsic names:

cp_async_bulk, cp_red_async_bulk, cp_async_tensor,
cp_async_prefetch_tensor, fence_view_async,
viaddmax, viaddmin, viadd, vimax, vimin, vimax3, vimin3,
write_async, tcbar, mmareadshma, tcmma,
tcshift, tcatomsws, tcldsws, tcstsws,
gdesc, breuse, bkeep, virtcount,
memclear, acqshminit, sparsify, spfactor2to4,
2x64dp128bitlw02lw13, 2x64dp128bitlw01lw23, 4x32dp128bit,
16dp32bitt0t15, 16dp32bitt16t31,
selfcast, broadcast, findandset, align

These names represent the Blackwell SASS operations exposed through the OCG intrinsic interface, covering tensor core scheduling (tcbar, tcmma, tcshift), sparse operations (sparsify, spfactor2to4), integer min/max variants (viaddmax, viaddmin), and async memory operations.

New PTX Instructions (Blackwell-Specific)

Beyond tcgen05, Blackwell introduces or extends several instruction families visible in the opcode dispatch:

CLI Options — Tensor Memory Checks

Two CLI options control tcgen05 guardrail instrumentation:

When enabled, the compiler inserts inline validation code (the 8 guardrail functions) around tcgen05 operations. These emit trap instructions if tensor memory invariants are violated at runtime — useful for debugging TMEM allocation errors, bounds violations, and ownership conflicts.

Feature Comparison

SASS Instruction Encoding

Blackwell continues the 128-bit per-instruction format introduced in Turing. The Mercury encoder handles the SM100+ instruction set through a dedicated encoding subsystem spanning approximately 851KB in the address range 0xDFC000--0x107B000.

The encoding subsystem covers 16 instruction format groups with full Blackwell ISA support including:

• Standard ALU/FPU/memory operations (inherited from earlier architectures)
• TCGEN05 tensor memory operations (new encoding classes)
• BGMMA/QMMA block-scale and quantized MMA variants
• Extended bulk copy operations (UBLKCP variants)
• Sparse tensor operations

Function Map

Cross-References

• SM Architecture Map — Validation tables, codegen factory encoding, suffix semantics
• Turing & Ampere (SM 75–88) — Baseline features that Blackwell inherits
• Ada & Hopper (SM 89–90a) — WGMMA, clusters, TMA — carried forward to Blackwell
• TCGen05 — 5th Gen Tensor Cores — Detailed tcgen05 ISA, TMEM model, instruction encoding
• Capsule Mercury & Finalization — Capmerc format, ELF structure, finalization levels
• Mercury Encoder — Shared encoding pipeline for all Mercury/capmerc output
• Intrinsic Table (608 Entries) — Full intrinsic catalog with sm_10x and sm_1xx ranges
• Latency Model & HW Profiles — Per-SM scheduling parameters and functional units
• Uniform Register Optimization — UR-file passes, Blackwell uniform ALU additions
• CLI Options — --gpu-name, --binary-kind, tensor memory check flags
• SASS Instruction Encoding — 128-bit format, Mercury encoder selection