Section Catalog & EIATTR

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

A CUDA cubin is a standard ELF container with NVIDIA-proprietary extensions. ptxas v13.0.88 populates it with approximately 4*(N+M) sections minimum for a program with N entry kernels and M device functions. Each section carries a specific kind of data — SASS instructions, constant bank contents, relocation entries, per-kernel resource metadata (EIATTR), shared memory layout, debug information, or Mercury-encoded streams for deferred finalization. This page catalogs every section type ptxas can emit, the NVIDIA-specific ELF section types used, the section ordering rules, and the complete EIATTR attribute encoding.

NVIDIA-Specific Section Types

Beyond standard ELF section types (SHT_PROGBITS, SHT_STRTAB, SHT_SYMTAB, SHT_RELA, SHT_NOTE), ptxas uses NVIDIA-defined types in the SHT_LOPROC--SHT_HIPROC range (0x70000000--0x7FFFFFFF). All of these were re-extracted from the binary by emitting both a final cubin and a -c relocatable object and reading the raw sh_type words.

Two output-mode vocabularies. In a -c relocatable object (ET_REL) the memory-space sections keep their SHT_CUDA_* LOPROC types as listed above. In a final cubin (ET_EXEC, default) the same sections are lowered to standard types — .nv.constant*→SHT_PROGBITS, .nv.global/.nv.shared→SHT_NOBITS. Only .nv.info, .nv.compat, and .nv.callgraph keep LOPROC types in the final cubin.

The section creator sub_1CB3570 contains a range check on CUDA-specific types:

// sub_1CB3570 -- section type validation
if (elf_mode != 1 && is_relocatable
    && ((sh_type - 0x70000064) <= 0x1A || sh_type == 0x70000006)) {
    // These CUDA section types require special handling in relocatable mode
}

0x70000064 (SHT_CUDA_CONSTANT_B0) is the base of the constant-bank range, not the callgraph type. The check covers 0x70000064--0x7000007E (the 18 numbered banks plus the seven Mercury constant banks) plus the legacy 0x70000006 — exactly the sections whose addresses are assigned by the constant-bank/shared-memory allocators rather than the generic ELF layout walk.

Complete Section Catalog

Standard ELF Infrastructure Sections

Created unconditionally by the ELFW constructor (sub_1CB53A0). These form the skeleton of every cubin.

NVIDIA Note Sections

Created unconditionally. Carry module-level metadata the CUDA driver reads before launching any kernel. Both notes use sh_type = SHT_NOTE (0x07), owner string "NVIDIA Corp\0" (namesz=12), and note version u16 = 2.

The full byte-level NOTE layout and program-header capture are in the ELF header & NOTE reference.

Per-Kernel Code Sections

Created by sub_1CB42D0, one set per kernel entry and device function:

The .rela companion is auto-created by the section creator when SHF_EXECINSTR is set. The assertion "adding function section after callgraph completed" fires if a code section is added after call graph analysis.

Per-Kernel Metadata Sections

The .nv.info.<func> section uses SHF_INFO_LINK (flag 0x40) to declare its association with the function's symbol, with sh_info holding the function's symbol index. Binary evidence from ptxas shows sub_1CC7FB0 consistently passes 1879048192 (0x70000000) as SHT_CUDA_INFO for both the global and per-entry .nv.info sections. (0x70000064 is the unrelated SHT_CUDA_CONSTANT_B0 constant-bank type — see the section-type table above.)

Global Metadata Sections

Constant Banks

CUDA supports up to 18 numbered constant banks (0–17) plus named constant sections:

The layout calculator sub_1C9DC60 skips .nv.constant0 sections during offset assignment because their addresses are managed by the OCG constant bank allocator, not the ELF layout engine.

Shared Memory Sections

Texture / Surface Descriptor Tables

Created when the program declares texture references, surface references, or samplers. Each section records the descriptor footprint for a specific binding mode so the driver can size its descriptor heap at load time.

The master section allocator sub_1CABD60 assigns addresses to shared, constant, and local memory sections. The layout calculator skips .nv.reservedSmem for the same reason it skips .nv.constant0 — its address comes from the shared memory master allocator.

Unified Function/Data Tables

The error "Number of .nv.uft jump slots != Number of entries" fires when the UFT and entry tables are inconsistent. "missing nv.uft.entry" fires when the required entry table section was never created.

DWARF Debug Sections

Generated when --device-debug or --generate-line-info is active:

NVIDIA Debug Extensions

Mercury / Capsule Mercury Sections (SM 100+)

For Capsule Mercury output (Blackwell and later), the cubin contains a parallel set of .nv.merc.* sections carrying Mercury-encoded instruction streams plus all metadata needed for deferred finalization:

The Mercury section cloner (sub_1CA2E40) iterates all sections and duplicates constant, global, shared, and local sections into the .nv.merc.* namespace, creating corresponding .nv.merc.rela sections for relocations.

Global vs Per-Kernel Sections

The .nv.info / .nv.info.<func> split is the primary distinction between global and per-kernel metadata:

Global .nv.info (one per cubin):

• sh_link = 0 (no associated symbol)
• Contains module-wide EIATTR records: EIATTR_CUDA_API_VERSION, EIATTR_STATISTICS, EIATTR_HAS_PRE_V10_OBJECT, EIATTR_MERCURY_ISA_VERSION
• Created by sub_1CC7FB0(elfw, 0) — the zero argument selects global mode

Per-kernel .nv.info.<func> (one per kernel):

• Section name: sprintf(".nv.info.%s", func_name) (visible in sub_1CC7FB0)
• sh_link points to the symbol table entry for the function
• sh_flags includes SHF_LINK_ORDER (0x40) to declare its association
• Contains per-kernel EIATTR records: EIATTR_REGCOUNT, EIATTR_NUM_BARRIERS, EIATTR_FRAME_SIZE, etc.
• Created by sub_1CC7FB0(elfw, sym_idx) when sym_idx != 0

The .nv.info section creator (sub_1CC7FB0) first searches for an existing section of type 0x70000000 with the appropriate name. If none exists, it creates one. The per-function variant links the new section to the function's .text section via sub_1CB4180.

Section Ordering

During finalization, sections are sorted into 8 priority buckets that determine their order in the output ELF. The bucket assignment is performed against the well-known section name set seeded from the static table at off_2403A60 (22 entries, terminator dword_2403B70); unknown names fall through to a kind-based default. Authoritative table is in elf-emitter.md → Section Ordering; the abbreviated form below covers the names this catalog discusses.

Within each bucket, sections appear in creation order. Section file offsets are assigned by sub_1C9DC60 walking the sorted list with alignment padding. The .debug_line section receives special alignment padding for DWARF line table requirements. .nv.constant0.<func> and .nv.reservedSmem are present in the bucket-4 sort order but are skipped during offset assignment — their file offsets are written by the OCG constant-bank allocator and the shared-memory master allocator (sub_1CABD60) respectively.

Offset Assignment

// sub_1C9DC60 -- simplified layout algorithm
uint64_t offset = elf_header_size;
for (int i = 0; i < section_count; i++) {
    section_t* sec = sorted_sections[i];
    if (is_virtual(sec))          continue;  // flag & 4 -> no file data
    if (is_nv_constant0(sec))     continue;  // OCG allocator manages these
    if (is_nv_reservedSmem(sec))  continue;  // shared memory allocator manages these

    if (sec->sh_addralign > 1)
        offset = (offset + sec->sh_addralign - 1) & ~(sec->sh_addralign - 1);
    sec->sh_offset = offset;
    offset += sec->sh_size;
}

Three section types are skipped during offset assignment:

1. Virtual sections (flag bit 2 set) — have no file data, only metadata
2. .nv.constant0 — address assigned by the OCG constant bank allocator
3. .nv.reservedSmem — address assigned by the shared memory master allocator sub_1CABD60

EIATTR Encoding

Each .nv.info section contains a flat sequence of EIATTR (Entry Information Attribute) records. There is no section header or record count — the parser walks from byte 0 to sh_size, consuming records sequentially. The EIATTR builder is sub_1CC9800 (14,764 B native / 86 KB decomp) — one of the three largest functions in the output pipeline.

Element Wire Format

Each element begins with a 1-byte EIFMT format tag and a 1-byte EIATTR attribute code; the format tag — not the attribute — determines the payload shape. The four formats and their on-wire sizes:

The writer picks the EIFMT per emit call from the payload width: no value → NVAL, one byte → BVAL, half → HVAL, sized blob → SVAL. Elements are 4-byte aligned; an 8-byte SVAL struct that would straddle an odd boundary is preceded by an EIATTR_PAD element. Note the format tag is independent of the attribute — the same EIATTR code may be emitted as different EIFMTs in different contexts (e.g. EIATTR_MAXREG_COUNT is HVAL, EIATTR_REGCOUNT is SVAL carrying a {symidx, regCount} record).

Worked Wire Examples

Global .nv.info for a leaf kernel (sm_90), one triple per function symbol:

04 2f 0800 0e000000 0a000000   EIFMT_SVAL EIATTR_REGCOUNT   {symidx=0xe, regCount=0xa}
04 11 0800 0e000000 00000000   EIFMT_SVAL EIATTR_FRAME_SIZE {symidx=0xe, frameSize=0}
04 12 0800 0e000000 00000000   EIFMT_SVAL EIATTR_MIN_STACK_SIZE {symidx, stackSize}

Per-kernel .nv.info.addvec (sm_90), in emission order:

04 37 ...   EIATTR_CUDA_API_VERSION    0x82 (=130)
04 17 ...   EIATTR_KPARAM_INFO         ordinal 2  (12-byte record, descending ordinal)
04 17 ...   EIATTR_KPARAM_INFO         ordinal 1
04 17 ...   EIATTR_KPARAM_INFO         ordinal 0
03 50 0000  EIATTR_SPARSE_MMA_MASK     0x0  (HVAL; default on sm_90, absent on sm_75)
03 1b ff00  EIATTR_MAXREG_COUNT        0xff (255 = no maxregcount limit)
03 5f 0101  EIATTR_MERCURY_ISA_VERSION 0x101 (Mercury ISA v1.1; =0x0 on sm_75)
04 1c ...   EIATTR_EXIT_INSTR_OFFSETS  0x90 (byte offset of EXIT in .text)
03 19 1400  EIATTR_CBANK_PARAM_SIZE    0x14 (=20 bytes total param size)
04 0a ...   EIATTR_PARAM_CBANK         {symidx=0xf, offset=0x210, bytes=0x14}
04 36 ...   EIATTR_SW_WAR              0x8

When a kernel actually uses named barriers (e.g. an sm_75 kernel with three bar.sync), an extra element appears — EIFMT_BVAL EIATTR_NUM_BARRIERS = 0x03. The barrier count lives in EIATTR_NUM_BARRIERS (code 76); the SHF_BARRIER_MASK section-flag field that earlier ptxas builds populated (count = sh_flags >> 20) is no longer written.

EIATTR_KPARAM_INFO Bitfield

The 12-byte EIATTR_KPARAM_INFO (code 23) record describes one kernel parameter. Decoded byte-exact from the emitted payload (no reliance on cuobjdump labels):

bytes [0:4]   u32 index      symbol index (0 for the common case)
bytes [4:6]   u16 ordinal    parameter position
bytes [6:8]   u16 offset     byte offset within the param constant bank
bytes [8:12]  u32 packed, LSB-first:
                logAlign  : 8
                space     : 4
                cbankid   : 5
                isSmemParam : 1
                size      : 14

Worked example — ordinal-2 record bytes 02 00 10 00 00 f0 11 00: index=0, ordinal=2, offset=0x10, packed=0x0011f000 → logAlign=0, space=0, cbankid=(0x11f000>>12)&0x1f=0x1f (=31, param cbank), isSmem=0, size=(0x11f000>>18)&0x3fff=4. Parameters are emitted in descending ordinal.

Parsing Pseudocode

uint8_t *ptr = section_data;
uint8_t *end = section_data + section_size;

while (ptr < end) {
    uint8_t eifmt = ptr[0];
    uint8_t attr  = ptr[1];
    switch (eifmt) {
    case 0x01:  // NVAL: no payload
        process_flag(attr);
        ptr += 2;
        break;
    case 0x02:  // BVAL: one byte
        process_byte(attr, ptr[2]);
        ptr += 4;          // 1 value byte + 1 pad
        break;
    case 0x03:  // HVAL: u16
        process_half(attr, *(uint16_t *)(ptr + 2));
        ptr += 4;
        break;
    case 0x04: {  // SVAL: u16 size + payload
        uint16_t size = *(uint16_t *)(ptr + 2);
        process_blob(attr, ptr + 4, size);
        ptr += 4 + ALIGN_UP(size, 4);
        break;
    }
    }
}

EIATTR Record Emitter — sub_1CC85F0

The low-level function that writes one EIATTR TLV record. Called from the builder and propagator with parameters:

// sub_1CC85F0 -- emit one EIATTR record
void emit_eiattr(
    ELFW*    elfw,       // a1: ELFW object
    uint8_t  attr_code,  // a2: EIATTR type code (e.g., 0x2F for REGCOUNT)
    int16_t  size,       // a3: payload size in bytes
    void*    payload,    // a4: pointer to payload data
    uint32_t sym_idx     // a5: symbol index (0 = global)
);

Before emitting, it calls sub_1C97840 to check whether the attribute code is supported on the current SM architecture. If not supported, the record is silently skipped. It then calls sub_1CC7FB0 to obtain or create the appropriate .nv.info section, allocates a 16-byte record descriptor, fills the format byte and attribute code, and appends it to the section's linked list (offset +392 in the ELFW object).

EIATTR Attribute Catalog

ptxas v13.0.88 defines 97 EIATTR codes numbered 0 through 96 (plus the EIATTR_ERROR_LAST sentinel at 96). The complete catalog below is cross-referenced against the nvlink v13.0.88 name table (extracted from pointer table at VA 0x1D37D60) and verified against EIATTR codes observed in the ptxas EIATTR builder (sub_1CC9800 switch cases and sub_1CC85F0 call sites).

Complete Code Table

Fmt column: Idx = format 0x04 (indexed, per-symbol), Free = format 0x01 (raw bytes), Val = format 0x02 (single 32-bit value), Sized = format 0x03 (16-bit value).

EIATTR Codes Confirmed in ptxas Builder

The following codes appear as explicit case labels in the sub_1CC9800 switch statement or as arguments to sub_1CC85F0:

The builder's first pass uses a switch with cases 0x04, 0x0D, 0x0F, 0x11, 0x12, 0x1B, 0x1E, 0x23, 0x26, 0x2F, 0x38, 0x3B, 0x41, 0x4A, 0x4C, 0x4F, 0x50, 0x51, 0x52, 0x54 to sort EIATTR records into per-entry arrays. A second pass emits the final records via sub_1CC85F0 and sub_1CC86D0.

Symbol Index Resolution Pass

Before the main builder runs, the EIATTR builder performs a symbol index resolution pass (lines 700–884 in the decompiled builder). This pass walks all pre-existing EIATTR records and resolves symbol indices through the linker's mapping tables:

// Simplified from sub_1CC9800 lines ~716-824
for (record in eiattr_list) {
    switch (record->attr_code) {
    case 0x02: case 0x06: case 0x07: case 0x08: case 0x09:
    case 0x0A: case 0x11: case 0x12: case 0x13: case 0x14:
    case 0x17: case 0x23: case 0x26: case 0x2F: case 0x3B:
    case 0x45:
        // Indexed format: resolve sym_idx through mapping table
        int32_t *sym_ptr = (int32_t *)record->payload;
        if (mapping_table && *sym_ptr != 0) {
            if (*sym_ptr < 0)
                *sym_ptr = negative_mapping[-(*sym_ptr)];
            else
                *sym_ptr = mapping_table[*sym_ptr];
        }
        if (*sym_ptr == 0 && attr_code != 0x45 && attr_code != 0x17) {
            record->attr_code = 0;  // disable record
        }
        break;
    case 0x0F:
        // EXTERNS: resolve each 4-byte symbol index in the array
        int count = record->size / 4;
        for (int i = 0; i < count; i++) {
            resolve_sym(&payload[i], mapping_table, negative_mapping);
        }
        break;
    }
}

The bitmask 0x800800060000 (seen at line 716) encodes which attribute codes use the simple indexed-resolve path: it selects codes 2, 6, 7, 8, 9, 10, 17, 18, 19, 20, 23, 38, 47, 59, 69.

Barrier and Register Propagation — sub_1CC8950

When a device function uses barriers or a high register count, those requirements must propagate upward through the call graph to each entry kernel. The propagator sub_1CC8950 handles this:

"Creating new EIATTR_NUM_BARRIERS and moving barcount %d
    from section flags of %s to nvinfo for entry symbol %s"
"Propagating higher barcount %d to the section flags
    of %s of entry symbol %s"
"regcount %d for %s propagated to entry %s"

The propagator emits EIATTR_REGCOUNT (0x2F) records via sub_1CC85F0(_, 0x2F, 8, ...) and handles EIATTR_NUM_BARRIERS (0x4C) through the sub_1CC7FB0 path. Barrier counts are extracted from the section flags field at bit offset 20 (7-bit field, mask 0x7F), then cleared from the section flags (&= 0xF80FFFFF) after being moved into an EIATTR record.

EIATTR Categories by Function

Resource allocation (GPU driver reads these to configure hardware at launch):

Parameter bank:

Instruction offset tables (driver/tools locate and patch instructions at load time):

Texture and surface:

Cluster and cooperative launch (sm_90+):

Shared memory:

Hardware workarounds:

Blackwell+ (sm_100+):

Mercury:

Graphics-specific:

.nv.compat Section

The .nv.compat section (SHT_CUDA_COMPAT = 0x70000086) stores forward-compatibility attributes. Its records use a different format from EIATTR — each is a small TLV with:

Offset  Size  Field
------  ----  -----
0x00    1     format (always 0x02 = value)
0x01    1     compat_code
0x02    1     value

The sub_1CC93A0 handler processes these with a switch over compat codes 2–6:

The guard *(_DWORD *)(a1 + 72) <= 0x59 (SM version <= 89 decimal) means compat processing only applies to SM 90 (Hopper) and later. Unknown compat codes trigger: "unknown .nv.compat attribute (%x) encoutered with value %x." (note the typo "encoutered" in the binary string).

Architecture-Gated EIATTR Emission

Not all EIATTR codes are valid on all SM architectures. The function sub_1C97840 performs architecture checks before emitting a record. Observed gates:

The sub_1C97840 function takes an EIATTR code and the SM version from the ELFW object's field at offset 624, returning a boolean. This prevents older EIATTR codes from appearing in Blackwell cubins and prevents Blackwell-only codes from appearing in Hopper cubins.

Constant Bank Optimization

The master section allocator sub_1CABD60 (11,856 B native / 66 KB decomp) performs two major space optimizations during address assignment: constant value deduplication within .nv.constant0 banks, and shared memory interference-graph coloring for extern shared variables. Both run before final offset assignment.

Constant Value Deduplication — sub_1CA6890

When multiple kernels in the same compilation unit use identical constant values, the OCG constant bank can contain duplicates. sub_1CA6890 (454 lines decompiled) eliminates them by value-matching, reducing .nv.constant0 section size.

The algorithm dispatches on constant value width:

For each constant data node in the section's linked list (at section+72):

1. Extract the value bytes (node+0), alignment (node+16), and size (node+24).
2. Look up the value in the appropriate dedup structure.
3. If duplicate found: alias the current symbol's offset to the existing symbol's offset. Debug output: "found duplicate value 0x%x, alias %s to %s" (32-bit) or "found duplicate 64bit value 0x%llx, alias %s to %s" (64-bit) or "found duplicate %d byte value, alias %s to %s" (N-byte via sub_1CA6760).
4. If not found: align the section cursor to the required alignment, append the data via sub_1CA6650, and insert into the dedup structure.

After aliasing, the function rewrites pending relocations that targeted the now-eliminated range:

// Simplified relocation rewriting after dedup alias
for (reloc in pending_relocs) {
    if (reloc.section == target_section
        && reloc.offset >= old_data_offset
        && reloc.offset <  old_data_offset + old_data_size) {
        reloc.offset = reloc.offset + alias_target_offset - old_data_offset;
        // "optimize ocg constant reloc offset from %lld to %lld"
        unlink(reloc);  // remove from pending list
    }
}

Special cases:

• Zero-valued constants: A "seen set" (parameter a15) prevents distinct zero-valued symbols from being aliased to each other, since different __constant__ variables may legitimately hold zero but need separate addresses.
• Redirect mode: When parameter a13 is set and sub_1CB15C0 returns true for a symbol, the constant is redirected to its defining section rather than deduplicated.

The caller sub_1CABD60 wraps this in an optimization check: "optimize OCG constants for %s, old size = %lld". If dedup does not reduce the section size, it reverts: "ocg const optimization didn't help so give up".

Shared Memory Interference Graph — sub_1CA92F0

When a CUDA program declares multiple extern __shared__ variables used by different kernels, they can potentially share the same memory if no single kernel uses both simultaneously. sub_1CA92F0 (585 lines decompiled) builds an interference graph and performs greedy graph coloring to pack shared objects into minimum total space.

Phase 1 — Build usage sets (which kernels reference each shared object):

For each global shared object, walk all referencing functions. A kernel "uses" a shared object if it directly references it or transitively calls a device function that does (traced via sub_1CBD800). Objects used by exactly one kernel are privatized — moved into a per-entry .nv.shared.<func> section. Unused objects are removed entirely (symbol flags set to mark deleted).

"global shared %s only used in entry %d"    -- privatize
"remove unused global shared %s"             -- delete

Phase 2 — Build interference edges:

For each pair of remaining shared objects (i, j), test whether their usage sets intersect (via sub_42E460 set membership). If any kernel uses both, they interfere — they cannot overlap in memory. Edges are stored as linked lists per object.

Phase 3 — Greedy graph coloring:

Objects are processed in sorted order. For each object:

1. Mark all colors used by interfering neighbors as unavailable.
2. Assign the lowest available color (starting from 1).
3. Update the color's alignment requirement (max of all objects in that color group).
4. Update the color's size requirement (max of all objects in that color group).

"  allocate to group %d"    -- color assignment

Phase 4 — Compute group offsets:

Groups are laid out sequentially with alignment padding:

group_offset[1] = align_up(base, group_align[1]);
for (g = 2; g <= num_groups; g++)
    group_offset[g] = align_up(group_offset[g-1] + group_size[g-1], group_align[g]);
total_size = group_offset[last] + group_size[last];

Each shared object's final offset is group_offset[its_color]. The total extern shared size is written to the section descriptor. Per-entry shared sections are expanded if a referenced object's offset + size exceeds their current size.

"esh %s size = %lld"
"for shared object (%d) %s:"
"  offset = 0x%llx, size = 0x%llx"
"  edge to %d"
"  allocate to group %d"

Constant Bank Optimization Functions

Section Attribute Builder — sub_60FBF0

The per-kernel section attribute builder sub_60FBF0 (76 KB decompiled, 2,541 lines, VA 0x60FBF0) runs once for each kernel entry point and device function. It assembles the full per-function section configuration object (648 bytes), parses compile option overrides, remaps PTX memory space codes to ELF section type IDs, conditionally creates three section types, then invokes the Mercury codegen pipeline (sub_6F52F0, DecodePipeline::RunStages).

Inputs

The function takes three parameters:

Option Parsing

The function iterates the key-value list at a1+38 and matches five string keys by character-by-character comparison:

// Simplified from sub_60FBF0 lines ~638-812
for (int i = 0; i < list_length(a1->options); i++) {
    const char** kv = list_get(a1->options, i);
    if (strcmp(kv[0], "deviceDebug") == 0)
        config->deviceDebug = 1;                // config+24
    else if (strcmp(kv[0], "lineInfo") == 0)
        config->lineInfo = 1;                   // config+25
    else if (strcmp(kv[0], "optLevel") == 0) {
        if (!config->optLevel_locked)            // config+108
            config->optLevel = strtol(kv[1], ...);  // config+104
    }
    else if (strcmp(kv[0], "IsCompute") == 0)
        config->isCompute = (strcmp(kv[1], "True") == 0);  // config+184
    else if (strcmp(kv[0], "IsPIC") == 0)
        config->isPIC = (strcmp(kv[1], "True") == 0);      // config+191
}

Additional inputs read from a1 and a2:

Memory Space Type Remapping

PTX internal memory space type codes in the 0x10000 range are remapped to compact ELF section type IDs by the helper sub_60DA40 (and inline copies). This remapping is applied to every symbol in the global, extern, texture, surface, and per-function symbol lists:

Special handling: when the space code is 0x10003 (constant0) and the compilation mode is relocatable (*(a3+48) == 12), the descriptor's needs_reloc flag (byte 33) is set to 1, indicating the constant0 section requires special relocation handling during linking.

The value 65596 (0x1003C) serves as a threshold — symbols with (space_type - 0x1003C) < 2 are counted into the texture/surface allocation arrays.

Conditional Section Creation

Three per-kernel sections are conditionally created:

.sass_map.<func> — created when *(a2+232) != 0 (sass_map generation enabled):

if (context->sass_map_enabled) {                 // a2+232
    descriptor = alloc(64);                       // 64-byte section descriptor
    memset(descriptor, 0, 64);
    pthread_mutex_lock(context->mutex);           // a2+240
    // Allocate instruction tree node and connect to codegen state
    name = sprintf(".sass_map%s", func_name);     // ".sass_map" + func_name
    descriptor->name = name;
    pthread_mutex_unlock(context->mutex);
}

.nv.local.<func> — created when the register spill size (config+112) is non-zero:

if (config->spill_size != 0) {                   // config+112
    descriptor = alloc(64);
    descriptor->size = config->spill_size;
    // Name: ".nv.local." + bare_func_name (skip ".text." prefix)
    name = sprintf(".nv.local.%s", func_name + 6);
}

The spill size at config+112 is set from the sum of the register spill count and frame size when the spill flag is non-zero.

.nv.constant<N>.<func> — created when:

1. The compilation mode field equals 2 (*(a1->target+48) == 2)
2. No pre-existing constant section exists (*(a1+172) == 0)
3. The function's symbol list is empty

if (mode == 2 && !has_constant_section && no_symbols) {
    descriptor = alloc(64);
    int bank = target->get_section_type() - 0x70000064;
    int size;
    if (func_const_size <= target->get_min_const_size())
        size = target->get_min_const_size();     // vtable+80
    else
        size = target->get_max_const_size();     // vtable+88
    descriptor->size = size + func_const_size;
    name = sprintf(".nv.constant%d.%s", bank, bare_func_name);
    descriptor->data = calloc(descriptor->size);
}

Assembler Flag Processing

The assembler flag list at a1+39 is iterated. Each entry's value string (at offset +8) is split on spaces via strtok_r. Each token is validated by sub_60F790, which constructs a temporary 656-byte object to test the flag. Valid tokens are concatenated with spaces and appended to config+48 (the toolkit info string that ends up in .note.nv.tkinfo).

Codegen Pipeline Invocation

After configuration, the function calls sub_6F52F0 (DecodePipeline::RunStages) with 18 parameters including the configuration object, all 7 descriptor arrays, the ELFW context, and the function name. The return code is mapped:

Post-Pipeline Section Registration

After the Mercury pipeline returns successfully:

1. Calls sub_60DD30 twice for pre/post code region finalization
2. Calls sub_60DBE0 for each optional symbol table (texture, surface, global) to register their sections with the ELFW emitter
3. Calls sub_1CB9C30 on the ELFW object (a2+32) to commit all sections
4. If SM version <= 0x45 (SM 69): creates UFT/UDT entries (section types 68/69) for each resolved symbol
5. Under mutex lock, ORs the per-function WAR bitmask (config+232..240) into the global accumulator at a2+504

Thread Safety

All shared state modifications are protected by the mutex at a2+240:

• String length accumulator updates (a2+296, a2+304) for string table pre-allocation
• WAR bitmask accumulation (a2+504)
• .sass_map section setup (instruction tree access)
• Instruction merge from secondary codegen contexts (a2+80, a2+88)

Key Functions

Cross-References

• Custom ELF Emitter — ELFW object, section ordering, ELF header
• Relocations & Symbols — relocation resolution, UFT/UDT management
• Debug Information — DWARF generation and .debug_* sections
• Mercury Encoder — Mercury encoding that feeds .nv.merc.* sections
• Capsule Mercury — SM 100+ capsule and .nv.capmerc sections
• Pipeline Overview — where section emission fits in the pipeline
• ELF header & NOTE reference — byte-exact header, program-header, NOTE, and .nv.compat capture