Turing & Ampere (SM 75--88)

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

SM 75 through SM 88 span two microarchitecture generations that ptxas treats as a contiguous feature band. sm_75 (Turing) is the default target for ptxas v13.0.88 -- when --gpu-name is omitted, sub_6784B0 returns sm_75. The Ampere targets (sm_80, sm_86, sm_87, sm_88) share generation-7 SASS encoding and add incremental tensor core and async-copy capabilities. sm_89 (Ada Lovelace) is architecturally Ampere-derived internally but is covered in Ada & Hopper because it bridges to sm_90 features.

SM Version Table

Codegen factory encoding: (isa_generation << 12) | sub_variant. Turing is generation 6; Ampere is generation 7. The sub-variant distinguishes silicon cut within a generation. sm_75 and Pascal sm_60 share generation 6 (sm_60 = 24576 = 0x6000), differentiated by sub-variant 0 vs 1.

sm_88 note: Registered in ptxas with CUDA_ARCH=880 and codegen factory 28676, but no public product ships on this SM. It may represent an unreleased Ampere derivative or internal test target.

SM 82 -- Internal Ampere Target

sm_82 is an undocumented internal Ampere target present in the base validation table (unk_1D16220, entry [20]) but not registered in the profile constructor sub_6765E0. It has no capability dispatch handler, no profile object, and no handler functions in any of the 7 hash maps. It exists in ptxas solely as a validation table entry and as the SASS opcode generation boundary.

PTX 6.2 Anomaly

sm_82 requires PTX ISA version 6.2, which is lower than both its neighbors:

PTX 6.2 corresponds to CUDA 10.1 (Turing era). This backward version number strongly suggests sm_82 was created as an early Ampere development target -- a PTX-level placeholder added before the Ampere PTX ISA (7.0) was defined. The validation table entry was never removed, but no profile object was ever created for it.

SASS Opcode Boundary Role

sm_82's primary significance in ptxas is as the opcode generation boundary for Ampere SASS instructions. The opcode hierarchy uses SM-number-based range labels:

SM82_FIRST  = index 172   (first Ampere-era SASS opcode)
SM82_LAST   = index 193   (last opcode in the sm_82 range)

These 22 opcode slots (indices 172--193) cover the core Ampere SASS additions:

The name SM82_FIRST/SM82_LAST is used as the boundary label even though these instructions are available on sm_80+ (any codegen factory >= 28673). The "82" in the label refers to the internal target used during Ampere development, not to a minimum SM requirement for the opcodes themselves.

Why sm_82 Matters

sm_82 is a ghost target: it occupies a validation table slot and lends its name to an opcode range, but cannot be compiled for. Passing --gpu-name sm_82 to ptxas would pass the initial validation check (bsearch succeeds in the base table) but fail during profile construction because sub_6765E0 has no case for SM 82. The practical consequence is that sm_82 is a naming artifact preserved from Ampere development, not a usable compilation target.

Profile Object Construction

Each SM's intrinsic table initializer (Map 3 handler) calls sub_917990 to allocate a 1,936-byte profile object, then populates architecture-specific fields.

Source: decompiled sub_60A2E0 (sm_75), sub_60A3E0 (sm_80), sub_60AC30 (sm_86), sub_60AD30 (sm_87), sub_60AB30 (sm_88).

All five initializers are structurally identical. The only field that differs between them is offset +348 (codegen factory). Common initialization:

v6 = sub_917990(a3);              // allocate 1936-byte profile object
*(_DWORD *)(v6 + 344) = 0x100000; // shared memory config base (1 MB)
*(_DWORD *)(v6 + 348) = FACTORY;  // codegen factory (per-SM)
v7[147] = 0;                      // cleared
v7[421] = cli_option_value;       // from a1+108
v7[460] = 1;                      // enable flag
v7[482] = 1;                      // enable flag
v7[470] = 1;                      // enable flag
v7[177] = 1;                      // enable flag

The sub_917990 base constructor sets the default codegen factory to 0x2000 (8192), which every per-SM initializer immediately overwrites. Key base constructor fields:

Handler Dispatch

ptxas registers per-SM handler functions into 7 parallel hash maps via sub_607DB0. For sm_75--88, all handlers are thin wrappers around shared codegen infrastructure.

Map 1 (Handler A) and Map 2 (Handler B)

Every Handler A function is identical:

bool handler_A(int64_t a1, int64_t a2) {
    *(int32_t*)(a2 + 100) = read_option(a2, "cpf_optx");
    return sub_663C30(a2, 0) != 0;  // a2=0: sets *(a1+96) = 1
}

Every Handler B function is identical except the second argument:

bool handler_B(int64_t a1, int64_t a2) {
    *(int32_t*)(a2 + 100) = read_option(a2, "cpf_optx");
    return sub_663C30(a2, 1) != 0;  // a2=1: skips the flag set
}

The "cpf_optx" option controls OptiX IR compilation mode. sub_663C30 is the core codegen-factory-aware driver that delegates to sub_662920 (ELF section iteration, 26KB) and sub_7FBB70 (actual compilation pass). The only behavioral difference between Handler A and Handler B: when called as Handler A (arg=0), sub_663C30 sets *(a1 + 96) = 1 before processing -- likely a "primary pass" flag.

Map 3 (Intrinsic Table Initializer)

Maps 6--7 (Performance / Occupancy)

These handlers return per-SM occupancy parameters used by the driver API's occupancy calculator. Other Ampere SMs have their own handlers registered into the same maps (addresses not fully traced in sweep data).

Scheduler Profile

sub_8E4400 (InitHWProfile_Warp) uses the codegen factory value to select scheduling parameters. The dispatch is a linear threshold cascade:

All SM 75--88 targets fall into the 7-warp / 208-slot bucket. After the warp count, a secondary switch maps specific codegen factory values to sub-architecture variants (stored at a1+26):

sm_75 (24577) and sm_80 (28673) are absent from the variant table and fall through to the default variant (0 or 1). This means sm_75 and sm_80 use the baseline latency model, while sm_86--88 get tuned sub-architecture parameters.

HW Latency Tables

Each SM has a dedicated latency table function containing per-opcode pipeline stage assignments, stall cycle costs, and functional unit mappings. These are called from the scheduling infrastructure to drive instruction ordering.

sm_80 uniquely has two latency tables: a "base" table (sub_8E7940, 2.9KB) and a full table (sub_8E7B40, 3.3KB), suggesting a layered lookup where the base provides defaults and the full table overrides specific entries. sm_86's table is the largest at 4.4KB, likely because RTX 30xx consumer GPUs have different pipeline characteristics from A100 datacenter parts.

No separate table entry was found for sm_88 in the sweep data. It may share sm_86 or sm_87's latency profile, or be registered through a path not captured in the sweep.

SASS Instruction Encoding

SM 75--88 all use the 128-bit per-instruction encoding format introduced with Turing. This replaced the Volta/Pascal scheme where scheduling control was packed into a separate 64-bit control header shared by 3 instructions.

Control Word Layout

Each 128-bit SASS instruction carries a 23-bit control word encoding scheduling decisions. The control word is generated by sub_A36360 (52KB, the control-word / scoreboard encoder).

bits [0:3]   stall count       (4 bits, max 15 cycles)
bit  [4]     yield flag        (1 bit, warp scheduler hint)
bits [5:7]   write barrier idx (3 bits, selects 1 of 6 barriers)
bits [8:13]  read barrier mask (6 bits, one per barrier)
bits [14:19] wait barrier mask (6 bits, one per barrier)
bits [20:25] reuse flags       (6 bits, per source operand)

Instruction Word Structure

128-bit instruction word:
  bits [0:3]     = 0x2       (format code: 128-bit)
  bits [4:6]     = slot      (scheduling group slot)
  bits [8:16]    = MAJOR     (9-bit major opcode, 0x00-0x171)
  bits [17:24]   = MINOR     (8-bit minor opcode / variant)
  bits [25:31]   = SUBOP     (7-bit sub-opcode / format ID)
  bits [48+]     = MODIFIERS (format-dependent modifier fields)
  bits [132:134] = 0x0       (extended opcode flag, at offset 0x84)

The 3-level opcode hierarchy (major/minor/subop) allows up to 102 major opcodes with 48 sub-operations each. Maximum observed variant value: 0x2F (47 decimal).

Scoreboard / Dependency Barriers

SM 75--88 provide 6 hardware dependency barriers per warp. The scoreboard tracker is managed by sub_8E4920 (BuildScoreboardEntries, 6.9KB) and encoded by sub_A36360.

The scoreboard tracker allocates 952 bytes per function when bit 4 of the flag byte at offset +1385 is set. An additional 856-byte bitset is allocated when bit 8 is also set (for barrier register tracking in writeback mode).

The scoreboard infrastructure is shared across all SM 75--88 targets. The barrier count (6) is constant for this entire range. sm_90 (Hopper) potentially increases this, and sm_100+ (Blackwell) changes the barrier model further.

Intrinsic Table

Intrinsic availability is cumulative. Each generation adds to the previous.

sm_75 Baseline (IDs 0x89--0x1FA, 370 intrinsics)

sm_75 inherits the full sm_70 (Volta) intrinsic set labeled __cuda_sm70_*:

WMMA intrinsics cover all combinations of:

• Shapes: m16n16k16, m32n8k16, m8n32k16
• Operations: load_a, load_b, load_c, store_d, mma
• Layouts: row/col combinations
• Types: f16, f32, with/without satfinite
• Address spaces: generic, global, shared

sm_80 Additions (IDs 0x1FB--0x22F, 53 intrinsics)

sm_80 adds two intrinsic groups:

14 __cuda_sm80_* intrinsics (IDs 0x1FB--0x208):

39 __cuda_sm_8x_mma_* intrinsics (IDs 0x209--0x22F):

Extended MMA shapes and sparse variants for the 2nd/3rd generation tensor core.

Intrinsic Gate Mechanism

The per-SM intrinsic table initializer (Map 3 handler) controls which intrinsics are available. sm_75 registers only the sm_70 intrinsic block. sm_80+ additionally registers the sm_80 and sm_8x blocks. The gate is not a runtime check -- it is a registration-time decision: if the intrinsic's handler function is not registered for a given SM, the PTX parser emits an error when it encounters a call to that intrinsic.

Peephole Optimizer Gates

The SASS-level peephole optimizer (sub_83EF00, 4,858 lines decompiled) applies pattern-matching transformations to the instruction stream. Several transformations are gated by codegen factory thresholds.

FMA/DFMA Combining

// sub_83EF00, case 0x50 (opcode 80 = FMA/DFMA)
if (*(_QWORD *)(*(_QWORD *)(a1 + 1584) + 372) > 28673) {
    // FMA combining enabled: sm_86+ (codegen factory 28674+)
    // Look for CVT -> FMA patterns and combine
}

Gate: codegen_factory > 28673. This means:

• sm_75 (24577): FMA combining disabled
• sm_80 (28673): FMA combining disabled (equality fails the > check)
• sm_86 (28674): FMA combining enabled
• sm_87 (28675): FMA combining enabled
• sm_88 (28676): FMA combining enabled

This is a significant compiler optimization difference between A100 (sm_80) and RTX 30xx (sm_86). A100 code does not get FMA/DFMA combining, possibly because A100's pipeline already has hardware FMA fusion or because the combining transformation is not profitable on GA100 silicon.

Master Gate

All peephole passes check capability 487 ("enable peephole") and capability 350 (per-instruction gate) before applying any transformation. These capabilities are controlled by optimization level and user options, not by SM target.

BB Initialization Flags

The basic block initializer sub_6E8EB0 sets architecture-specific flags using a secondary SM version encoding (distinct from the codegen factory):

This secondary encoding uses (gen << 12) with generation 5 for Ampere in the BB init context. The specific bit flags control opcode descriptor table population -- each BB gets a set of 40+ (opcode_id, encoding_word) pairs that define which SASS instructions are legal in that basic block.

Feature Comparison

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 Turing/Ampere targets share the 7-partition / 208-slot scheduling geometry. The major resource difference is sm_80 (A100 datacenter) with 2,048 max threads and 64 warps vs. the consumer/embedded parts (sm_86--88) with 1,536 max threads and 48 warps. sm_75 (Turing) is the most constrained with 1,024 max threads and 32 warps.

Codegen Factory Gating Patterns

Throughout ptxas, the codegen factory value at profile offset +348 (or equivalently *(*(QWORD*)(a1+1584)+372) from the function context) is used to gate features. Common patterns for the Turing/Ampere range:

// Pattern 1: Generation check (factory >> 12)
int gen = codegen_factory >> 12;
if (gen >= 7) { /* Ampere+ feature */ }

// Pattern 2: Exact threshold
if (codegen_factory > 28673) { /* sm_86+ only */ }
if (codegen_factory >= 28673) { /* sm_80+ */ }

// Pattern 3: Range check
if (codegen_factory >= 24577 && codegen_factory < 32768) {
    /* Turing/Ampere only, not Hopper+ */
}

The >> 12 shift extracts the ISA generation, allowing coarse checks (Turing = 6, Ampere = 7). Fine-grained checks compare against specific factory values to distinguish sm_80 from sm_86, etc.

Function Map

Cross-References

• SM Architecture Map -- Overview of all 23 SM targets and the 3-level profile system
• Ada & Hopper (SM 89--90a) -- sm_89 (Ada) shares Ampere codegen factory range but bridges to Hopper features
• Blackwell (SM 100--121) -- Next-generation targets with codegen factory 36864+
• Intrinsic Table (608 Entries) -- Full intrinsic catalog with per-SM generation ranges
• SASS Instruction Encoding -- 128-bit encoding format, bitfield packer, opcode hierarchy
• Peephole Optimization -- FMA combining and other post-scheduling SASS transforms
• 3-Phase Scheduler Architecture -- Scheduler infrastructure that consumes HW latency tables
• CLI Options -- --gpu-name parsing, sm_75 default