GPU Target Architecture

45 SM variants across 6 generations. Processor table at qword_502A920 (stride-2 layout: name + PTX version). Architecture gating throughout the binary controls feature availability.

Per-SM Deep Dives:

• SM 70-89 (Volta through Ada Lovelace) -- Feature configuration call order, complete sub_60E7C0 flag table, atomic lowering, cumulative flag profiles
• SM 90 -- Hopper -- Thread block clusters, TMA descriptor format and lowering, WGMMA, setmaxnreg, distributed shared memory
• SM 100 -- Blackwell Datacenter -- tcgen05 tensor core ISA, arch-conditional vs. family-conditional gating, cvt_packfloat FP4/FP6/MX formats
• SM 120 -- Blackwell Consumer -- No tcgen05, .offset.bindless texture intrinsics, f16 texture support, mma.sync.block_scale (future)

Complete SM Table

Legacy architectures also present in the table but not in the CLI mapping: sm_20, sm_21, sm_30, sm_32, sm_35, sm_37, sm_50, sm_52, sm_53, sm_60, sm_61, sm_62, sm_70, sm_72, sm_73.

Suffix Meanings

PTX Version Mapping

Architecture Gating

Four subsystems cooperate to configure feature flags from the SM version. The master configurator sub_60E7C0 runs last and has the highest non-CLI priority. For the complete flag table per tier, see SM 70-89 Complete sub_60E7C0 Flag Table.

Feature Configuration Pipeline

CLI parser (sub_617BD0)              Sets byte_4CF8* override flags
    |
sub_60DFC0 (secondary)               Sets unk_4D041B8 at sm_80+ (__VA_OPT__)
    |
sub_60D650 (optimization level)      ~109 flags from -O level
    |
sub_60E7C0 (master SM configurator)  ~60 flags via SM threshold comparisons
    |--- sub_60E530 (tertiary)        Supplementary progressive unlocks
    |
sub_982C80 (NVPTX subtarget)         224-byte bitfield for LLVM backend

Override priority: CLI flag > SM version > Optimization level > C++ standard version > CUDA mode > Virtual arch flag. See CLI Flag Inventory for the complete CLI flag-to-pipeline routing and Optimization Levels for per-level flag differences.

EDG-Level Gates -- sub_60E7C0

Sets ~60 unk_4D04* feature flags based on SM version thresholds. Each flag is gated by a byte_4CF8* user-override check.

Backend Subtarget Feature Offsets (NVPTXSubtarget)

For the complete NVPTXSubtarget analysis, see NVPTX Target Infrastructure.

Intrinsic Verifier Architecture Gates -- sub_2C7B6A0

The NVVMIntrinsicVerifier (143KB) gates intrinsics by SM version. For the complete three-layer verification architecture, see NVVM IR Verifier.

Feature Gate Matrix

This matrix shows which major compiler features are available at each SM tier. Each cell links to the detailed discussion in the per-SM deep-dive page.

Tensor Core / MMA Instructions

See Tensor / MMA Builtins for the per-builtin ID reference and Tensor / MMA Codegen for the code generation pipeline.

Memory and Synchronization

See Atomics Builtins for atomic PTX generation detail and Barriers & Sync for barrier builtins.

Thread Block Clusters

Numeric Formats

Texture and Surface

See Surface & Texture Builtins for the tex_surf_handler dispatch algorithm.

EDG Frontend Feature Flags

See EDG 6.6 Frontend for the 737-define configuration system.

tcgen05 Sub-Variant Access Table

The tcgen05 instruction family uses a two-tier gating system unique to Blackwell. Base variants (sm_100, sm_103, sm_110) are excluded; only a and f sub-variants pass the bitmask check.

Generation-Specific Features

Turing (sm_75)

sm_75 is the default architecture for cicc v13.0, hardcoded as "compute_75" in sub_900130 and sub_125FB30.

• Base tensor core (HMMA m16n16k16) -- see Tensor / MMA Builtins
• Conversion intrinsics
• Baseline for cicc v13.0 (default architecture) -- see CLI Flag Inventory

Full detail: SM 70-89 (Volta through Ada)

Ampere (sm_80-sm_89)

• L2::cache_hint on atomic operations (sub_21E6420) -- see Atomics Builtins
• Extended tensor core shapes (tf32, bf16) -- see Tensor / MMA Builtins
• Async copy (cp.async) -- see SM 70-89: Ampere
• C++20 __VA_OPT__ support -- the sole differentiator between sm_75 and sm_80+ in sub_60E7C0/sub_60DFC0

Full detail: SM 70-89 (Volta through Ada)

Hopper (sm_90/90a)

• Cluster operations: barrier.cluster.arrive/wait, fence.sc.cluster -- see Cluster Barriers
• Cluster registers: %cluster_ctarank, %clusterid.x/y/z, %is_explicit_cluster -- see Cluster Special Registers
• Kernel attributes: .blocksareclusters, .maxclusterrank, .reqnctapercluster, .cluster_dim -- see PTX Directives
• setmaxnreg: Dynamic register allocation limit (sub_21EA5F0) -- see setmaxnreg
• TMA: Tensor Memory Access with Im2Col, dimension validation, 2CTA mode -- see TMA
• WGMMA: Warpgroup MMA async (f16, bf16, tf32, f8) -- see WGMMA
• Distributed shared memory: .shared::cluster qualifier for cross-CTA access -- see DSMEM
• Mbarrier extensions: DMA fence/arrive/wait for TMA coordination -- see Mbarrier

Full detail: SM 90 -- Hopper

Blackwell Datacenter (sm_100-sm_103)

• tcgen05: Next-gen tensor core instruction set (scaleD, transA, negA, negB at sub_21E8CD0) -- see tcgen05
• Arch-conditional vs. family-conditional gating: Two-tier feature system for tcgen05 sub-instructions -- see Gating
• match instruction: Architecture-gated ("match instruction not supported on this architecture!") -- see sm100
• Extended MMA shapes: m16n8k256 with MX format support
• .offset.bindless intrinsics -- gated at sm_120+, NOT sm_100 (see sm120 .offset.bindless)
• cvt_packfloat extended types: FP4, FP6, MX formats -- see cvt_packfloat

Full detail: SM 100 -- Blackwell Datacenter

Jetson Thor (sm_110)

sm_110 is architecturally a datacenter Blackwell derivative (originally sm_101 before rename). It retains full tcgen05/TMEM hardware on a/f sub-variants. The sm_110 section is documented on the sm_120 page because the two are often compared.

Full detail: SM 120 -- Jetson Thor section

Blackwell Consumer (sm_120, sm_121)

• No tcgen05: The entire tcgen05 ISA is rejected by cicc for all sm_120/121 variants -- see No tcgen05
• .offset.bindless texture intrinsics (68 variants) -- see .offset.bindless
• 16-bit texture element types -- see f16 Texture
• mma.sync.block_scale: Present in upstream LLVM 22 but NOT emitted by cicc v13.0 -- see block_scale
• Tensor core falls back to HMMA/IMMA inherited from sm_70-sm_90 path

Full detail: SM 120 -- Blackwell Consumer

NVVM Container Architecture Enum -- sub_CD09E0

The NVVM container format uses an architecture enumeration. See NVVM Container for the complete tag inventory.

Notable: NVVM_ARCH_HW_SM_10_4 (sm_104) and NVVM_ARCH_BLACKWELL_11_0 are not publicly documented. NVIDIA's internal naming uses "BLACKWELL" for all sm_100-sm_121 variants, even though sm_110 is marketed as Jetson Thor and sm_120/121 are a distinct consumer microarchitecture (RTX 50xx). See SM 120: Architecture Identity for the "SM 10.4" internal designation.

Target Triples

The nvsass-nvidia-directx and nvsass-nvidia-spirv triples (discovered in sub_2C80C90) reveal that NVIDIA's SASS-level backend supports DirectX and SPIR-V targets alongside traditional CUDA and OpenCL.

Data Layout Strings

Address space 3 (shared memory) uses 32-bit pointers even in 64-bit mode, controlled by nvptx-short-ptr and nvptx-32-bit-smem flags. See Address Spaces for the complete address space reference.

SM Version Encoding

Two parallel version tracking systems coexist in the binary:

• qword_4F077A8 -- Encodes SM_MAJOR * 10000 + SM_MINOR * 100. Used in approximately 309 decompiled files, primarily in the NVVM frontend and optimizer. Boundary thresholds use the XX99 pattern (e.g., 69999 for pre-Volta, 89999 for pre-Hopper). See SM 70-89: SM Version Encoding for full detail.

• unk_4D045E8 -- Stores the raw SM number as a decimal (e.g., 75 for sm_75, 89 for sm_89). Used in approximately 12 decompiled files, primarily in the builtin checker and atomic lowering logic. See SM 70-89: unk_4D045E8 Frontend Gates for the complete gate table.

Cross-References

• NVPTX Target Infrastructure -- NVPTXTargetMachine, NVPTXSubtarget, TTI hooks
• Tensor / MMA Builtins -- Per-builtin-ID reference for all MMA generations
• Tensor / MMA Codegen -- Code generation pipeline for tensor core operations
• Atomics Builtins -- Atomic PTX generation and scope validation
• Surface & Texture Builtins -- Texture intrinsic dispatch algorithm
• NVVM IR Verifier -- SM-gated intrinsic verification
• NVVM Container -- Architecture enum and tag inventory
• CLI Flag Inventory -- -arch=compute_XX parsing and flag routing
• Optimization Levels -- Per-level flag differences that interact with SM gates
• EDG 6.6 Frontend -- 737-define configuration, CUDA keyword handling
• Address Spaces -- Address space 3 shared memory and data layout strings
• GPU Execution Model -- CTA, warp, and cluster execution model context