SM80-88 Ampere

The Ampere family in nvlink v13.0.88 covers four compute capabilities -- sm_80, sm_86, sm_87, and sm_88 -- all sharing the ISA class string "Ampere", the same ISel backend at 0xCA0000--0xDA0000 (1 MB), and the same 128-bit SASS instruction encoding introduced by Turing. The sm_88 variant is new in CUDA 13.0. None of the four Ampere targets carry 'a' or 'f' suffix variants (unlike Blackwell targets), so the profile database contains exactly 12 Ampere profile entries: 4 real, 4 virtual, and 4 LTO.

For general Ampere architecture details (hardware specs, PTX ISA requirements, codegen factory encoding, scheduler parameters, latency tables), see the ptxas wiki: Turing/Ampere. For cicc-level feature gates and optimizer flag configuration, see the cicc wiki: SM70-89. This page focuses on the nvlink-internal per-sub-architecture data: profile registration, capability vectors, dispatch table function pointers, and the ISel backend shared by all four.

Per-Architecture Profiles

Profile Registration in sub_484F50

The profile database initializer sub_484F50 (53,974 bytes) registers the four Ampere architectures in numeric order. Each architecture produces three profile objects via sub_484DB0 (real sm_, virtual compute_, LTO lto_), which are inserted into the global hash map at qword_2A5F8D8. The Ampere block spans lines 288--462 of the decompiled source.

The registration pattern for each is identical:

// 1. Create real profile (is_virtual=0, is_lto=0)
sm_80 = ArchProfile::create(0, 0, "sm_80", "sm_80", "Ampere",
                            "-D__CUDA_ARCH__=800", "sm_80");

// 2. Create virtual profile (is_virtual=1, is_lto=0)
compute_80 = ArchProfile::create(1, 0, "compute_80", "compute_80", "Ampere",
                                 "-D__CUDA_ARCH__=800", "compute_80");

// 3. Link real <-> virtual
sm_80->virtual_ptr = compute_80;       // offset +72
compute_80->virtual_ptr = compute_80;  // self-ref

// 4. Insert into hash map
LinkerHash::insert(qword_2A5F8D8, "sm_80", sm_80);
LinkerHash::insert(qword_2A5F8D8, "compute_80", compute_80);

// 5. Create LTO profile (is_virtual=1, is_lto=1)
lto_80 = ArchProfile::create(1, 1, "lto_80", "compute_80", NULL,
                             "-D__CUDA_ARCH__=800", "lto_80");
lto_80->virtual_ptr = compute_80;
LinkerHash::insert(qword_2A5F8D8, "lto_80", lto_80);

// 6. Build compatibility lists
list_append(compute_80->compat_list_2, sm_80);    // offset +64
list_append(sm_80->compat_list_2, compute_80);
list_append(sm_80->compat_list_1, sm_80);          // offset +56: family
list_append(sm_80->compat_list_0, sm_80);          // offset +48: cross-variant

// 7. Copy capability vectors from rodata
sm_80->capability[0] = xmmword_1D40F10;   // offset +80
sm_80->capability[1] = xmmword_1D40F40;   // offset +96
sm_80->capability[2] = xmmword_1D40F30;   // offset +112

After registering each sub-architecture, sub_484F50 links it into the Ampere family chain. sm_86 is appended to sm_80's family list (compat_list_0 and compat_list_1), sm_87 and sm_88 follow the same chaining to sm_86's lists (which transitively connect back to sm_80).

The dword_2A5F8CC = 80 assignment (line 326) sets the default minimum architecture to sm_80 -- any link operation that does not specify an explicit --arch flag defaults to this value.

Architecture Identity Matrix

All four share same-decade group 8, which also includes sm_89 (Ada). This means code compiled for sm_80 is compatible with any target in the 80--89 range via the same-decade rule (see the Compatibility page).

Capability Vectors

Each profile stores three 128-bit XMM vectors at offsets +80, +96, and +112. These encode hardware capability bitmasks checked by the finalization pipeline (sub_470DA0). The vectors are loaded from rodata constants in sub_484F50:

The rodata symbols encode these capability tiers:

Key observations:

• Vec 0 is constant across all architectures. It encodes the universal baseline capabilities.
• Vec 1 differentiates sub-architectures. sm_80 gets xmmword_1D40F40 (Ampere base), while sm_86/87/88 share xmmword_1D40F50 (extended Ampere). The sm_86 value propagates to sm_87 and sm_88 through copy chains (v211 = v29 at line 375; v210 = v37 at line 418), not through independent rodata loads. This means sm_86, sm_87, and sm_88 are capability-identical from the finalization pipeline's perspective.
• Vec 2 is constant within the pre-Blackwell group. All Ampere targets use xmmword_1D40F30.
• The Ampere-to-Ada boundary is marked by the Vec 1 change from xmmword_1D40F50 (sm_86/87/88) to xmmword_1D40F60 (sm_89).

What Distinguishes Each Sub-Architecture

sm_80 (GA100 -- datacenter Ampere). The base Ampere target and generational anchor. It uses xmmword_1D40F40 for Vec 1, which differs from the sm_86+ value xmmword_1D40F50. This means sm_80's capability mask is a strict subset of sm_86's -- code finalized for sm_86 is not guaranteed to be re-finalizable for sm_80 without capability mask verification. sm_80's codegen factory is 28673 (0x7001), placing it at sub-variant 1 within generation 7. In the ptxas scheduler, sm_80 falls through to the default variant (0 or 1) and uses baseline latency tables.

sm_86 (GA10x -- consumer/enterprise Ampere). Extends sm_80 with xmmword_1D40F50 capability bits. The codegen factory is 28674 (0x7002, sub-variant 2). In the ptxas scheduler, sm_86 maps to sub-architecture variant 2, receiving tuned scheduling parameters distinct from sm_80's baseline. The ptxas latency table for sm_86 (sub_8E7D80, 4.4 KB) is the largest in the Ampere family, reflecting the different pipeline characteristics of the RTX 30xx consumer die versus the A100 datacenter die.

sm_87 (GA10B -- Jetson Orin). Capability-identical to sm_86 (same Vec 1 xmmword_1D40F50, copied at line 418 via v37 = _mm_load_si128(&v211) where v211 holds sm_86's vector). The codegen factory is 28675 (0x7003, sub-variant 3). The ptxas scheduler maps it to variant 3, giving Orin its own latency profile (sub_8E8070, 3.5 KB). This is the only material binary difference between sm_86 and sm_87 in nvlink -- the capability mask, ISel patterns, and encoding pipeline are identical.

sm_88 (undocumented -- CUDA 13.0). Capability-identical to sm_86 and sm_87 (same Vec 1, copied at line 458 via the sm_87 chain). The codegen factory is 28676 (0x7004, sub-variant 4). The ptxas scheduler maps it to variant 4, shared with sm_110 (Jetson Thor). No separate latency table was found in the ptxas sweep data -- sm_88 may share sm_86's or sm_87's table. No public product ships on sm_88; it may represent an unreleased Ampere derivative or internal validation target.

Dispatch Table (sub_15C0CE0)

The dispatch table initializer sub_15C0CE0 registers seven callback function pointers per architecture into hash maps (qword_2A644B8 through qword_2A64488). Each callback serves a specific role in the embedded compilation pipeline.

Slot Assignments

Each compute_ prefix variant receives the same internal version constant as its real counterpart (e.g., compute_80 = byte_2A5EE3C, compute_86 = byte_2A5EE38).

Backend Init Functions (Slot A8)

The A8 slot handlers are the only callbacks with architecturally significant differences between Ampere sub-architectures. Each allocates a per-function codegen context via sub_189F230(a3) and writes a per-SM codegen factory value at offset +348:

The codegen factory encodes (isa_generation << 12) | sub_variant. All four Ampere targets are generation 7 (bits 12--15 = 0x7), differing only in the low 12-bit sub-variant field. The factory value controls:

1. Scheduler profile selection in sub_8E4400: all four fall into the 7-warp / 208-dispatch-slot bucket (factory range 24577--28676, threshold <= 32767).
2. Sub-architecture variant for latency tuning: sm_80 gets default variant, sm_86 gets variant 2, sm_87 gets variant 3, sm_88 gets variant 4.
3. Instruction encoding table selection: each sub-variant can enable/disable specific instruction forms.

All other fields at +344 (shared memory config base, 0x100000 = 1 MB) and the remaining context fields are identical across all four Ampere handlers.

Internal Version Numbers

The internal version numbers at qword_2A644A0 are stored as 4-byte integers at decreasing addresses, each 4 bytes apart. From the version mapping table in the SM89 Ada page:

The specific internal versions for sm_80--sm_88 are inferred from the address spacing pattern and the known threshold at *(a1+376) > 26 (sm_86+, confirmed in the compilation driver sub_1112F30 line 991). The sm_89 value of 29 is confirmed from decompiled code. The gap between sm_75 (internal 14) and sm_80 (inferred 25) reflects deprecated architectures in the 15--24 range that were removed from the profile database but whose internal version slots remain allocated.

Ampere vs Turing: Generational Differences

The transition from Turing (sm_75, generation 6) to Ampere (sm_80, generation 7) manifests in nvlink at several levels. For hardware and ISA-level details, see ptxas wiki: Turing/Ampere. Within nvlink specifically:

Profile-Level Changes

Capability Boundary

The Vec 1 change from xmmword_1D40F20 (Turing) to xmmword_1D40F40 (Ampere) marks a hard capability boundary. Code finalized for sm_80 cannot be re-finalized for sm_75 -- the capability mask comparison in sub_470DA0 will fail because Ampere capabilities are a superset of Turing capabilities. The reverse direction (sm_75 code on sm_80 targets) is permitted by the same-decade rule only if both are in the same family linked list, which they are not (different decades: 7 vs 8).

ISel Backend Independence

Despite sharing the same 128-bit instruction encoding format, sm_75 and sm_80 use completely separate ISel backends:

The Ampere ISel is slightly smaller than Turing's despite being a later generation. The reduction in pattern matchers (276 to 259) reflects instruction set refinement rather than feature removal -- Ampere reorganized and consolidated some multi-variant patterns.

Feature Gates in cicc and ptxas

From the cicc wiki: SM70-89, the Ampere-specific features enabled at the sm_80 threshold include:

• C++20 __VA_OPT__ support via unk_4D041B8
• llvm.nvvm.branch.if.convergent intrinsic (sm_80+, distinct from the sm_70+ branch.if.all.convergent)
• L2 cache hint atomics (PTX 7.3+)
• cp.async.bulk patterns for asynchronous memory copy

From the ptxas wiki: Turing/Ampere, the sm_82 SASS opcode boundary defines 22 Ampere-era opcode slots (indices 172--193) covering sparse MMA, binary tensor core shapes, FP64 tensor MMA, async copy infrastructure, and warp-wide reduction.

Family Linkage

The Ampere family linked list built by sub_484F50 chains all four sub-architectures through compat_list_1 (offset +56) and compat_list_0 (offset +48):

sm_80.compat_list_1 -> { sm_80, sm_86, sm_87, sm_88 }
sm_86.compat_list_1 -> { sm_86 }
    also: sm_86 appended to sm_80's compat_list_0 and compat_list_1
sm_87.compat_list_1 -> { sm_87 }
    also: sm_87 appended to sm_86's compat_list_0 and compat_list_1
sm_88.compat_list_1 -> { sm_88 }
    also: sm_88 appended to sm_87's compat_list_0 and compat_list_1

Additionally, sm_89 (Ada) links into the Ampere family chain:

// sub_484F50 lines 507--510:
list_append(sm_80->compat_list_0, sm_89);       // v14[3].m128i_i64[0] = sm_80's list
list_append(sm_80->compat_list_1, sm_89);       // v14[3].m128i_i64[1] = sm_80's list
list_append(sm_86->compat_list_0, sm_89);       // v22[3].m128i_i64[0] = sm_86's list
list_append(sm_86->compat_list_1, sm_89);       // v22[3].m128i_i64[1] = sm_86's list

This means sm_89 (Ada) is in the same compatibility family as the Ampere targets. The same-decade rule (80/10 = 89/10 = 8) already ensures this, but the explicit linkage provides direct traversal without arithmetic.

Worked Compatibility Example: sm_80 vs sm_86

Consider linking a cubin compiled for sm_80 into a final binary targeting sm_86.

1. sub_4878A0 parses both architecture strings: record_a = {arch_number=80, is_virtual=0, has_suffix=0}, record_b = {arch_number=86, is_virtual=0, has_suffix=0}.
2. Neither record is virtual, so the function does not reject on the "both virtual" check.
3. record_a.has_suffix is 0, so the function enters Case 1: profile_family_match(sm_80.family_list, sm_86).
4. sub_465450 traverses sm_80's compat_list_1: { sm_80, sm_86, sm_87, sm_88 }. sm_86 is found in the list.
5. Result: compatible. The linker accepts the sm_80 cubin into the sm_86 link.

The reverse direction (sm_86 cubin targeting sm_80) follows a different path:

1. Parse: record_a = {86, ...}, record_b = {80, ...}.
2. Case 1: profile_family_match(sm_86.family_list, sm_80).
3. sm_86's compat_list_1 contains { sm_86 } -- it does not contain sm_80.
4. Result: incompatible. sm_86 code cannot be linked for sm_80.

This asymmetry reflects the forward-compatibility guarantee: code compiled for a lower SM within a family runs on higher SMs, but not the reverse.

However, capability mask verification at finalization time (sub_470DA0) adds a second check: even if the linker accepts the cubin, re-finalization will compare the capability vectors. Since sm_80 uses xmmword_1D40F40 (Vec 1) and sm_86 uses xmmword_1D40F50 (Vec 1), the sm_86 target has a superset of sm_80's capabilities. The finalization check passes for sm_80-compiled code targeting sm_86 but would fail for sm_86-compiled code targeting sm_80 if any sm_86-specific capability bits were used.

ISel Backend (0xCA0000--0xDA0000)

All four Ampere sub-architectures share a single ISel backend. The backend is variant-agnostic -- it produces identical SASS encoding for sm_80, sm_86, sm_87, and sm_88. Any differences between sub-architectures are resolved upstream in the dispatch table (slot A8 codegen factory) and downstream in the scheduler (latency tables), not within the ISel code itself.

Address Map

Total analyzed functions (>3 KB): 413. The mega-hub at sub_D5FD70 is 239 KB (55,985 instructions, 1,340 callees) and too large for Hex-Rays to decompile. It is the third-largest function in the binary after the SM75 mega-hub (sub_FBB810, 280 KB) and the cuda_builtin_prototype_generator (sub_15B86A0, 345 KB).

Three-Phase Pipeline

Every IR instruction processed by this backend passes through three phases in sequence: instruction selection, operand emission, and binary encoding. The output is a 128-bit SASS instruction word ready for insertion into the device ELF .text section.

Phase 1: Instruction Selection. The mega-hub sub_D5FD70 iterates all 259 pattern matcher functions. Each matcher receives (ctx, ir_node, &pattern_id, &priority) and tests a single candidate encoding. The highest-priority match wins.

for each pattern_matcher in sm80_pattern_table[0..258]:
    pattern_matcher(ctx, ir_node, &pattern_id, &priority)
    if priority > best_priority:
        best_priority = priority
        best_id = pattern_id
sm80_emitter_table[best_id](ctx, ir_node)

Each pattern matcher queries: instruction attributes via sub_A49150(ctx, ir_node, slot_id) (up to 12 attribute checks per pattern), definition/use operand counts via sub_530FD0/sub_530FC0, individual operands via sub_530FB0(ir_node, index), and operand type predicates: isGPR (sub_CDD600), isPredicate (sub_CDD610), isUniformReg (sub_CDD630), isImmediate (sub_CDD670), isConstBuf (sub_CDD680).

Priority values range from 14 (least specific, fallback patterns) to 34 (most specific, heavily constrained patterns). Higher priority means more attribute checks and narrower operand constraints.

Phase 2: Operand Emission. The winning pattern dispatches to an operand emission function from Zone 1 (0xCA0000--0xCDC000). Each emitter handles one (opcode, format) combination and populates a structured instruction descriptor: opcode ID at *(a2+12), encoding format at *(a2+14), max operand slot count at *(a2+15), modifier flags via setter functions, and decoded register operands.

Phase 3: Binary Encoding. Zone 3 functions (0xD9A400--0xDA0000, 17 functions) produce the final 128-bit SASS binary instruction word using SSE2 intrinsics (_mm_or_si128) for efficient 128-bit bulk operations.

Instruction Set Coverage

The SM80 backend handles 19 distinct SASS opcodes with a total of 80 (opcode, format) emission variants:

These 19 opcodes represent the core compute-intensive instructions that the linker's embedded ptxas must emit during LTO compilation. The full Ampere ISA is substantially larger; instructions that never appear in LTO-generated code (control flow, barriers, texture, surface, etc.) are handled by separate codec tables at 0xC00070--0xC50970.

ISel Pattern Classes

The 259 ISel pattern matchers divide into 9 functional classes:

Priority levels range from 14 to 34. The priority system ensures more specific patterns win over generic fallbacks. Within each class, patterns form a lattice from general to specific:

Encoding Formats

The format field at *(a2+14) selects the operand encoding layout. Formats observed across all 19 opcodes:

Emission Function Catalog

The following tables list all 80 operand emission functions in Zone 1, organized by opcode. Each row is one (opcode, format) combination.

HMMA (Tensor Core) -- 11 variants

FFMA -- 12 variants

LDG (Global Memory Load) -- 9 variants

The RR/RI/RC/RR.ALT base forms are substantially larger (~11 KB each) than the predicated (.P) and shuffle forms (~3.5 KB), reflecting the complex cache hierarchy modifiers (strongOrder, eviction, scope, cacheOp, memoryType) that the base forms must encode.

IMAD.WIDE (64-bit Multiply-Add) -- 9 variants

IMAD.WIDE has the largest base-form emitters at ~12.7 KB each. The wide result format requires handling both halves of the 64-bit product, with separate register pair allocation.

Remaining Opcodes

Notable size outliers: sm80_emit_FADD_RC at 44,490 bytes is the single largest emission function, followed by sm80_emit_FADD_RI at 27,061 bytes. Both require extensive fixup tables for the constant-buffer and immediate operand forms of FP32 addition with negation, absolute value, saturation, and data type modifiers.

Bitfield Packing Functions

Zone 1 contains 75 bitfield packing functions. These translate the instruction descriptor's register IDs, opcode fields, and modifiers into bit positions within the 128-bit SASS instruction word:

Operand Type Predicates

Fifteen small predicate functions at 0xCDD5F0--0xCDD690 classify operands by type:

Instruction Modifier Setters

External Dependencies

The SM80 backend calls into shared infrastructure functions:

Relationship to Other Backends

Despite having fewer ISel patterns than SM75, the SM80 mega-hub matches the SM50-7x hub at 239 KB. The SM89/90 backend at 1.9 MB is substantially larger overall, though its mega-hub is smaller, because it includes 750 instruction encoder template instantiations that the SM80 backend does not require.

Confidence Assessment

Cross-References

nvlink Internal

• Architecture Profiles -- SM80 family profiles in the linker database, struct layout, capability vector table
• SM75 Turing -- predecessor ISel backend
• SM89 Ada -- successor backend, dispatch table comparison
• Compatibility Checking -- same-decade rule, family linkage, capability mask verification
• Architecture Dispatch -- per-arch function pointer dispatch mechanism
• ISel Hubs -- SM80 mega-hub sub_D5FD70 (239 KB)

Sibling Wikis

• ptxas: Turing/Ampere -- standalone ptxas SM80 target documentation (codegen factory encoding, scheduler profiles, latency tables, SASS encoding format)
• cicc: SM70-89 -- cicc compiler SM80 through SM88 feature gates (__VA_OPT__, convergent branches, L2 cache hint atomics)