SM Architecture Map

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

ptxas validates the --gpu-name target against three sorted lookup tables, constructs a profile object with family metadata and a CUDA_ARCH macro value, then populates seven parallel dispatch tables that drive capability checks, code generation factory selection, performance modeling, and occupancy calculation throughout the compiler. The default target is sm_75 (Turing). Every downstream decision -- instruction legality, encoder selection, register file geometry, scheduling latencies -- routes through the profile object built here.

Per-SM Deep Dives:

• Turing & Ampere (SM 75--88) -- Baseline feature set, codegen factory 24577/28673
• Ada & Hopper (SM 89--90a) -- WGMMA, cluster operations, codegen factory 32768
• Blackwell (SM 100--121) -- tcgen05, arch/family gating, codegen factory 36864
• TCGen05 -- 5th Gen Tensor Cores -- Blackwell tensor core ISA detail

Complete SM Table

23 active SM base targets ship in ptxas v13.0.88 (plus 9 legacy and 2 internal/alias entries retained in the validation table for backward compatibility). Each base target optionally has a (accelerated) and/or f (feature-reduced) sub-variants. The CUDA_ARCH column shows the value the -D__CUDA_ARCH__ macro expands to.

The family name stored in the profile object (from sub_6765E0) uses NVIDIA's internal naming: "Turing", "Ampere", "Hopper", "Blackwell". Ada Lovelace (sm_89) is stored as Ampere-derived internally despite being a distinct microarchitecture. sm_120/121 use "Blackwell" internally despite being a different consumer microarchitecture from sm_100 datacenter Blackwell.

Suffix Semantics

ptxas uses three suffix modes to control forward compatibility. The distinction is critical: it determines which SASS binary a cubin can execute on.

The base variant (no suffix) produces SASS that runs on the named architecture and all later ones: sm_80 code runs on sm_86, sm_89, sm_90, sm_100, etc. The a suffix locks the binary to exact silicon: sm_90a code runs only on H100/H200 hardware and will not execute on Blackwell. The f suffix allows forward compatibility within the same family: sm_100f code runs on sm_100 and sm_103 (both Blackwell datacenter) but not on sm_120 (different family).

Compilation rules from help text:

• sm_90a PTX must be compiled to sm_90a SASS (no cross-arch compilation)
• sm_100f PTX can compile to sm_100f or sm_103f SASS (same family)
• sm_100a PTX must compile to sm_100a SASS only
• Base sm_100 PTX compiles to any sm_100+ SASS

Sub-Variant Expansion

sm_75 through sm_89 have no a or f variants. sm_90 has only the a variant (no f). All Blackwell-era targets (sm_100+) have both a and f. sm_101 is a legacy alias for sm_110 (Jetson Thor, original internal designation); it passes validation but is not registered as a profile object, so its CUDA_ARCH values are not populated.

SM Validation Tables

Target name validation uses three sorted arrays searched via bsearch(). The CLI parser extracts the SM string from --gpu-name, strips any suffix, and searches the appropriate table.

Base Table -- unk_1D16220 (32 entries)

Contains all valid base SM names without suffix, sorted by numeric SM ID. Includes legacy architectures no longer supported for active compilation but retained for validation, plus two internal/alias entries. Each entry is 12 bytes: {uint32 sm_id, uint32 ptx_major, uint32 ptx_minor}. The bsearch comparison (sub_484B70) compares the numeric sm_id extracted from the --gpu-name string via sscanf.

sm_10, sm_11, sm_12, sm_13,                    // Tesla (legacy, PTX 1.0--1.2)
sm_20, sm_21,                                  // Fermi (legacy, PTX 2.0)
sm_30, sm_32, sm_35, sm_37,                    // Kepler (legacy, PTX 3.0--4.1)
sm_50, sm_52, sm_53,                           // Maxwell (legacy, PTX 4.0--4.2)
sm_60, sm_61, sm_62,                           // Pascal (legacy, PTX 5.0)
sm_70, sm_72,                                  // Volta (legacy, PTX 6.0--6.1)
sm_75,                                         // Turing (active, PTX 6.3)
sm_80, sm_82, sm_86, sm_87, sm_88, sm_89,      // Ampere/Ada (active, PTX 6.2--7.8)
sm_90,                                         // Hopper (active, PTX 7.8)
sm_100, sm_101, sm_103, sm_110, sm_120, sm_121 // Blackwell (active, PTX 8.6--9.0)

sm_82 (PTX 6.2): Undocumented internal Ampere target. Not registered in sub_6765E0 (no profile object). Serves as the SASS opcode generation boundary (SM82_FIRST/SM82_LAST, opcode indices 172--193). The anomalously low PTX version requirement (6.2 vs sm_80's 7.0) suggests it was an early development target added before PTX ISA versioning was finalized.

sm_101 (PTX 8.6): Original internal designation for Jetson Thor, renamed to sm_110 in a later CUDA release. Both entries coexist in the validation table for backward compatibility with PTX files referencing the old name. sub_6765E0 registers only sm_110; sm_101 is validation-only.

Accelerated Table -- unk_1D161C0 (7 entries)

sm_90a, sm_100a, sm_101a, sm_103a, sm_110a, sm_120a, sm_121a

One Hopper entry, six Blackwell entries. sm_101a is the legacy alias for sm_110a (Jetson Thor, original internal designation).

Feature-Reduced Table -- unk_1D16160 (6 entries)

sm_100f, sm_101f, sm_103f, sm_110f, sm_120f, sm_121f

No Hopper entry (sm_90 has no f variant). All Blackwell-era. sm_101f is the legacy alias for sm_110f.

Architecture Registration -- sub_6765E0

This 54KB function constructs profile objects for every SM version. Each profile contains:

The function registers each profile into three hash maps indexed by sm_XX, compute_XX, and lto_XX strings. This allows lookup by any of the three naming conventions used in different contexts (CLI, PTX .target directive, LTO linking).

Family assignment from sub_6765E0:

All sm_100 through sm_121 share the "Blackwell" family string internally, even though sm_110 (Jetson Thor) and sm_120 (consumer RTX 50xx) are distinct microarchitectures. The compiler distinguishes them through the capability dispatch tables, not through family name.

Capability Dispatch -- sub_607DB0

The capability dispatch initializer builds 7 parallel hash maps at initialization time, protected by a once-guard (byte_29FE1D8). Each map indexes sm_XX / compute_XX strings to per-architecture values or handler functions. Error recovery uses setjmp/longjmp.

Handler Function Assignments

Each SM version registers its own handler functions into these maps. Functions within the same suffix group (e.g., sm_100/100a/100f) share all handlers -- they are the same silicon with different feature exposure.

Map 1 -- Handler A (per SM):

Map 2 -- Handler B (per SM):

Map 3 -- Intrinsic table initializer (per SM):

Shared Handler Groups

Sub-variants within a base SM share all handler functions, confirming they are identical silicon:

Codegen Factory Values

The profile object stores an encoded architecture identifier at a known offset (visible as field +348 on the profile pointer chain, e.g., *(_QWORD *)(a1+1584)+348). This value is compared throughout the compiler to gate features:

These values appear in feature-gating checks. For example, FMA/DFMA combining in the peephole optimizer checks profile[+372] > 28673 to require sm_70+ capability. The exact encoding formula is (isa_generation << 12) | variant, where the high bits identify the SASS instruction set generation.

Related encoded values seen in the binary:

12288  = sm_30 (Kepler)       // 3 << 12
16385  = sm_50 (Maxwell)      // 4 << 12 | 1
20481  = sm_50 alt (Maxwell)  // 5 << 12 | 1
24576  = sm_60 (Pascal)       // 6 << 12
24577  = sm_75 (Turing)       // 6 << 12 | 1
28673  = sm_80 (Ampere)       // 7 << 12 | 1
28674-28677 = sm_86/87/88/89  // 7 << 12 | 2..5
32768  = sm_90 (Hopper)       // 8 << 12
36864  = sm_100 (Blackwell)   // 9 << 12
36865-36869 = sm_103..121     // 9 << 12 | 1..5

Hardware Resource Geometry

ptxas assembles per-SM hardware parameters from three data sources: sub_8688F0 (universal baseline), sub_8E4400 (scheduler partition geometry), and sub_ABF250 (occupancy calculator properties). These parameters control register allocation limits, shared memory partitioning, occupancy calculations, and scheduling decisions throughout the compiler.

Universal Constants (sub_8688F0)

sub_8688F0 sets the baseline hardware profile shared by all SM 75+ targets. These values are architecture-invariant within the ptxas v13.0.88 binary:

The register file size of 65,536 registers is confirmed by the EIATTR_REGCOUNT formula (code 0x2F): max_warps_per_SM = total_registers / (regcount * warp_size), and by explicit reference in codegen/templates.md ("the entire physical register file is 65,536 32-bit registers shared across all active warps").

Per-SM Resource Geometry Table

Combines binary evidence (sub_8E4400 scheduling profile, sub_8688F0 baseline, sub_ABF250 occupancy properties, sub_60AXXX per-SM initializers) with NVIDIA public specifications for parameters not stored as scalar constants in the binary. Confidence column rates how directly the value was extracted from the binary vs. inferred from public documentation.

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). Not stored as a ptxas constant; derived from warps_per_SM * warp_size / max_CTAs.
• 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). The scheduler partition count is the number of warp scheduler units; dispatch slots is the total scheduling capacity.
• Configurable Shared Memory: Valid shared memory sizes per CTA, selected by cudaFuncSetAttribute. Stored as pointer-to-table at profile offsets +1488/+1496; sm_75 has 3 entries, later architectures have more.

sm_88 note: No known product ships on sm_88. It shares all handler functions with sm_86. Listed parameters are inherited; actual hardware behavior is unverifiable.

Scheduler Partition Geometry (sub_8E4400 Detail)

The packed DWORD at offset +18 of the warp-level profile encodes scheduler partition counts. The WORD at offset +22 is the dispatch slot count -- a scheduling capacity value distinct from the raw warp count.

The dispatch slot count increases monotonically across generations, reflecting wider scheduling capacity. All sm_75 through sm_89 targets (Turing, Ampere, Ada Lovelace) share identical scheduling partition geometry despite their hardware differences -- the differentiation occurs in the per-SM latency tables, not in the partition structure.

Shared Memory Configuration Tables

ptxas stores configurable shared memory sizes as a pointer + count pair at profile offsets +1488 and +1496. The driver uses this table to validate cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize, ...) calls.

sub_8688F0 sets the sm_75 configuration:

*(a1+1488) = &unk_21D9168    // pointer to shared memory size table
*(a1+1496) = 3               // 3 valid configurations

For sm_75 (Turing), the 3 entries correspond to 32 KB, 48 KB, and 64 KB configurable shared memory. The L1/shared partitioning on Turing splits the 96 KB unified data cache between L1 and shared memory.

For sm_80 (Ampere), the configurable shared memory extends to 164 KB, reflecting the larger shared memory/L1 combined capacity. sub_ABF250 records the maximum as 167,936 bytes (163.8 KB) for the base sm_60 path and 233,472 bytes (228 KB) for sm_70+ paths, though these values encode via xmmword constants that depend on the specific SM variant.

For sm_90+ (Hopper, Blackwell), sub_ABF250 populates a maximum configurable value of 233,472 bytes (228 KB), supporting the opt-in extended shared memory mode added in Hopper.

Register Allocation Mechanics

ptxas allocates registers in units determined by the register allocation granularity stored in sub_ABF250:

The register allocation unit directly affects occupancy. With 256-register granularity on sm_75+, a kernel using 33 registers effectively consumes 40 (rounded up to the next multiple of 8), which means each warp uses 40 * 32 = 1280 of the 65,536 available registers, allowing up to 51 warps -- but capped by the hardware limit of 32 warps on sm_75.

The formula the GPU driver uses (from EIATTR_REGCOUNT documentation):

effective_regs = ceil(regcount / alloc_granularity) * alloc_granularity
regs_per_warp  = effective_regs * warp_size
max_warps      = min(registers_per_SM / regs_per_warp, hw_max_warps)
max_CTAs       = min(max_warps / warps_per_CTA, hw_max_CTAs)

SM Version Encoding

The raw SM version number stored in profile objects and code object headers uses a packed integer format. This is the value at v4[93] in the code object builder (sub_A465F0):

The code object builder (sub_A465F0 at 0xA465F0) maps these encoded SM versions to ELF code object version fields and thread-per-CTA limits. The magic number 0x16375564E is written at offset 0 of every code object header, with the SM version at offset +8.

Per-SM Capability Accessors -- sub_609XXX

The 24 functions in the sub_609XXX cluster (range 0x609280--0x609F60, ~1.2KB each) are the per-SM-version capability accessor functions. They are registered into Maps 1 and 2 of the dispatch tables and return architecture-specific values: register file sizes, feature flags, warp geometry, shared memory limits, and similar hardware parameters.

These are the functions that downstream code calls (through the dispatch table) to answer questions like "how many registers does this SM have?" or "is feature X available on this target?"

Profile Layering

Three levels of SM profile information cooperate:

Level 1: sub_607DB0                    // Capability dispatch (7 hash maps)
    |                                  //   -> feature flags, handler functions
    v
Level 2: sub_6765E0                    // Profile objects (name, family, CUDA_ARCH, lto)
    |                                  //   -> identity metadata, isaClass
    v
Level 3: sub_609XXX / sub_60AXXX       // Per-SM accessor functions
                                       //   -> concrete hardware parameter values

Level 1 provides the dispatch infrastructure. Level 2 provides identity metadata for diagnostics and linking. Level 3 provides the actual numeric values that drive register allocation, scheduling, and instruction legality.

Generation-Specific Features

Turing (sm_75)

sm_75 is the default architecture for ptxas v13.0.88, returned by sub_6784B0 when no --gpu-name is specified. Codegen factory value: 24577.

• Base tensor core (WMMA m16n16k16 f16/f32, m32n8k16, m8n32k16)
• Integer MMA (IMMA int8/int4), binary MMA (BMMA b1)
• Base warp-level operations (shfl.sync, vote.sync, match.sync, redux.sync)
• Named barrier support (bar 0-15 with arrive/red/sync variants)

Ampere (sm_80 -- sm_88)

Codegen factory value: 28673. Shared with Ada Lovelace (sm_89).

• Extended tensor core: TF32, BF16, FP64 MMA shapes
• cp.async for asynchronous shared memory copies
• L2 cache hints on atomic operations
• createpolicy instructions for cache management
• 14 additional intrinsics (__cuda_sm80_*: bf16/tf32/s4/s8/b1 MMA, createpolicy)

Ada Lovelace (sm_89)

Codegen factory value: 28673 (same as Ampere). Stored as "Ampere" internally despite being a distinct Ada Lovelace microarchitecture.

• Same codegen path as Ampere; differentiated through capability flags, not codegen factory
• 39 additional MMA intrinsics (__cuda_sm_8x_mma_*)

Hopper (sm_90 / sm_90a)

Codegen factory value: 32768. sm_90a is architecture-locked (H100/H200 only).

• WGMMA (warpgroup MMA async): wgmma.mma_async, wgmma.fence, wgmma.commit_group, wgmma.wait_group
• Cluster operations: barrier.cluster.arrive/wait, distributed shared memory
• setmaxnreg: Dynamic register allocation limit
• Cluster special registers: %clusterid, %cluster_ctaid, %cluster_ctarank, etc.
• 38 sub-byte MMA intrinsics (__cuda_sm_9x_mma_sub_byte_internal_*: s4/u4 sparse)

Blackwell Datacenter (sm_100, sm_103)

Codegen factory value: 36864. Both a and f sub-variants available.

• tcgen05: 5th-generation tensor core ISA (alloc, dealloc, ld, st, commit, cp, shift, mma) -- a/f sub-variants only
• tcgen05 guardrails: 8 debug validation functions (phase validity, column allocation, bounds checking)
• Extended MMA: 10 Blackwell-specific hmma/imma + bit MMA intrinsics (__cuda_sm_10x_*)
• 11 tcgen05 guardrail trap intrinsics (__cuda_sm10x_tcgen05_guardrail_trap_*)
• 18 sm_1xx bulk copy intrinsics (__cuda_sm1xx_*: cp.async.bulk.tensor 1D-5D tile/im2col)

Jetson Thor (sm_110)

Codegen factory value: 36864 (same as sm_100). Originally sm_101 before rename. Automotive/robotics SoC.

• Retains full tcgen05/TMEM hardware on a/f sub-variants
• Same Blackwell datacenter feature set for tensor operations
• Differentiated through capability flags for SoC-specific constraints

Blackwell Consumer (sm_120, sm_121)

Codegen factory value: 36864 (same as sm_100). Architecturally a distinct consumer microarchitecture despite sharing the "Blackwell" family string.

• No tcgen05: The entire tcgen05 ISA is absent on sm_120/121 -- gated by SM version checks
• Tensor core falls back to HMMA/IMMA/WGMMA inherited from sm_70--sm_90 path
• sm_120 = RTX 50xx consumer / RTX Blackwell Pro (enterprise)
• sm_121 = DGX Spark

Diagnostic Strings

Function Map

Profile Object Layout (1936 bytes)

Every SM's intrinsic table initializer (Map 3 handler) calls sub_917990 to allocate a 1,936-byte profile object that carries target-specific parameters throughout the compiler. This is the compilation unit's target descriptor -- the single structure that downstream code reads to answer "what hardware am I compiling for?"

Construction Sequence

1. sub_71BDE0(1936, a1)     heap allocate 1936 bytes
2. sub_C1B7A0(profile)      zero-fill + structural defaults (8 SSE blocks, 5 scalars)
3. sub_917990(a3)           overlay: codegen factory default, tail constants
4. sub_60AXXX(a1,a2,a3,a4)  per-SM: codegen factory, shmem base, capability flags

Key Fields -- Explicitly Initialized

These fields receive non-zero values during construction. Offsets are byte offsets from the profile object base pointer. Type column: D=DWORD(4B), Q=QWORD(8B), O=OWORD(16B), B=BYTE.

SSE Constant Blocks

10 blocks of 16 bytes each are loaded from .rodata segments via _mm_load_si128. These likely contain per-register-class sizing parameters, pipeline configuration constants, or default opcode table pointers. Exact values require .rodata dump.

Scheduling Feature Flags (+1385 bitfield)

The byte at offset +1385 is the most heavily accessed bitfield on the profile object (21+ consumer sites in 15+ decompiled functions). Each bit gates a scheduling or codegen behavior.

Per-SM Initializer Differences

All 12 per-SM initializers (one per SM family) are structurally identical. Only two fields differ between families.

All other fields (+344, +432, +588, +708, +1684, +1840, +1880, +1912, +1928) are set identically across all SM families.

Critical Distinction: Profile Object vs SM Backend

The profile object (1936 bytes, this layout) is the compilation unit's target descriptor, stored somewhere in the compilation context. The pointer at context+1584 (sm_backend_vtable_ptr) points to a separate polymorphic SM backend object -- not to another profile object. Fields accessed as *(*(ctx+1584) + N) are on the SM backend, not on this profile object.

Commonly accessed SM backend fields (NOT on the 1936-byte profile):

The profile's own +348 stores the codegen factory value at construction time. The SM backend's +372 is where all downstream code reads it. These are the same numeric value stored in two different objects.

Object Region Map

Offset Range   Size    Content
-----------    ----    -------
[0..111]       112B    Object header, vtable chain, zeroed bulk
[112..159]     48B     Config scalars (opt level, register limit)
[160..343]     184B    Zeroed + 3 SSE constant blocks (184, 280, 312)
[344..587]     244B    Target identity (shmem base, codegen factory, enable flags)
[588..879]     292B    Capability fields, 4 SSE constant blocks (680, 696, 740, 788)
[880..1063]    184B    Scheduling config, barrier config (+944=4 default)
[1064..1279]   216B    Extended config, vendor metadata (+1200="NVIDIA")
[1280..1535]   256B    Architecture feature flags (+1385 bitfield, +1417/+1418)
[1536..1663]   128B    Dynamic region pointer, phase counter, sm_backend ptr (+1584)
[1664..1831]   168B    CLI passthrough, ELF section data, barrier tracking ptr
[1832..1935]   104B    Tail region: 2 SSE constant blocks, mode value (+1892=2)

Cross-References

• Turing & Ampere (SM 75--88) -- Detailed feature flags for sm_75 through sm_89
• Ada & Hopper (SM 89--90a) -- WGMMA, cluster operations, sm_90a arch-lock
• Blackwell (SM 100--121) -- tcgen05, arch-conditional vs family-conditional gating
• TCGen05 -- 5th Gen Tensor Cores -- tcgen05 instruction set detail
• Intrinsic Table (608 Entries) -- Master intrinsic catalog with per-SM generation ranges
• Tensor Core Intrinsics -- WMMA/MMA/tcgen05 intrinsic lowering
• Latency Model & HW Profiles -- Per-SM scheduling parameters
• SASS Instruction Encoding -- Codegen factory -> encoder selection
• Mercury Encoder -- SM-dependent SASS encoding
• CLI Options -- --gpu-name parsing and default target
• Data Structures -- context+1584 polymorphic pointer documentation