.nv.info Metadata

The .nv.info section is NVIDIA's proprietary ELF metadata format that encodes per-kernel resource requirements and compilation artifacts. Every CUDA kernel carries a .nv.info section (or a per-function variant) that tells the GPU driver how many registers to allocate, how much shared memory to reserve, what barriers the kernel uses, and dozens of other resource descriptors. Without this metadata, the driver cannot launch the kernel -- it would have no way to know the kernel's hardware resource footprint.

nvlink both reads and writes .nv.info sections. During the merge phase it parses incoming .nv.info records to extract register counts for weak symbol resolution. During finalization it encodes computed properties (propagated register counts, barrier counts, stack sizes) back into .nv.info records in the output cubin. The embedded ptxas compiler produces .nv.info through a massive emission subsystem spanning ~190 functions across 1 MB of code.

Key Facts

Section Variants

A cubin contains two kinds of .nv.info sections, distinguished by name:

Global .nv.info

A single section named .nv.info with sh_link = 0 (no associated symbol). This contains attributes that apply to the entire compilation unit -- CUDA API version, compatibility flags, and shared metadata that is not specific to any one kernel.

sub_4504B0 creates this section when called with a2 = 0 (no symbol index). It calls sub_4411D0 to find an existing .nv.info section; if none exists, it creates one via sub_441AC0 with type 0x70000000, alignment 4, and sh_flags = 0.

Per-Function .nv.info.<name>

One section per kernel or device function, named .nv.info.<function_name> with sh_link pointing to the symbol table entry for that function. These sections carry per-kernel resource descriptors: register count, barrier count, stack sizes, parameter bank layout, and instruction-offset tables for various runtime patching needs.

sub_4504B0 creates per-function sections when called with a2 != 0. It constructs the name via sprintf(buf, "%s.%s", ".nv.info", func_name) where func_name is looked up from the symbol at offset +32. The section is created with sh_flags = 0x40 and sh_link pointing to the owning symbol. After creation, sub_4426D0 links the section to the function's symbol entry.

During the merge phase (sub_45E7D0), nvlink identifies .nv.info sections by checking sh_type == 0x70000000. The sh_link field determines whether a record is global (link=0) or per-function (link = symbol index). The merge function translates symbol indices from input-local to output-global using its mapping tables.

TLV Record Format

Each .nv.info section contains a flat sequence of 4-byte-aligned TLV (Type-Length-Value) records. There is no section header or record count -- the parser walks from byte 0 to sh_size, consuming records sequentially.

On-Disk Record Layout

Offset  Size  Field
------  ----  -----
0x00    1     format      Format byte (determines payload structure)
0x01    1     attr_code   EIATTR type code (identifies the attribute)
0x02    2     size        Payload size in bytes (little-endian uint16)
0x04    var   payload     Attribute-specific data (size bytes)

Total record size = 4 + size, padded to 4-byte alignment. The maximum theoretical payload is 65,535 bytes (16-bit size field), though in practice payloads rarely exceed a few hundred bytes.

In-Memory Node Layout (Linker Internal)

nvlink stores parsed TLV records as 16-byte linked-list nodes in an arena-allocated chain at elfw+392. Each node has:

Offset  Size  Field
------  ----  -----
0x00    1     format       Format byte (0x01--0x04)
0x01    1     attr_code    EIATTR type code
0x02    2     size         Payload size in bytes
0x04    4     secidx       Output section index (from sub_4504B0)
0x08    8     payload_ptr  Pointer to arena-allocated payload data

Nodes are linked via a separate linked-list structure managed by sub_4644C0 (list-append). The diagnostic function at sub_4478F0 iterates this list and prints each record as: "nvinfo <fmt=%d,attr=%d,size=%d>, secidx=%d" using fields at bytes 0, 1, 2, and 4 respectively.

Format Byte

The format byte at offset 0 controls how the payload is interpreted:

Format 0x04 (indexed) is the most common for per-function attributes. The 4-byte symbol index at payload offset 0 identifies which function the attribute applies to. The linker uses this index for symbol remapping during merge and for extracting per-function properties during finalization.

Format 0x03 (sized) encodes the value directly in the 16-bit size field of the TLV header, with no additional payload bytes. This is used for small integer attributes like MAXREG_COUNT and CBANK_PARAM_SIZE.

Format 0x01 (free) carries variable-length data. The size field gives the byte count. Offset tables use arrays of 4-byte instruction offsets; KPARAM_INFO uses structured records; EXTERNS uses arrays of symbol indices.

Parsing Pseudocode

From the decompiled merge path in sub_45E7D0 (lines 1900--2052):

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

while (ptr < end) {
    uint8_t  format    = ptr[0];
    uint8_t  attr_code = ptr[1];
    uint16_t size      = *(uint16_t *)(ptr + 2);

    if (format == 0x04) {
        /* Indexed format: first 4 bytes of payload = symbol index */
        uint32_t *payload = arena_alloc_copy(ptr + 4, size);

        /* Symbol index remapping per attr_code */
        switch (attr_code) {
        case 2: case 6: case 7: case 8: case 9:      /* texture/image */
        case 18: case 19: case 20: case 23: case 38:  /* param/cache */
        case 69:                                       /* kparam_info_v2 */
            payload[0] = sym_map[payload[0]];          /* remap first sym_idx */
            break;
        case 10:  /* PARAM_CBANK -- remap with lazy symbol creation */
            payload[0] = sym_map_or_create(payload[0]);
            break;
        case 15:  /* EXTERNS -- remap all uint32 entries */
            for (i = 0; i < size/4; i++)
                payload[i] = sym_map[payload[i]];
            break;
        case 17: case 35: case 47: case 59:  /* resource attrs */
            /* Check if symbol was deleted (weak-replaced) */
            if (is_deleted[payload[0]]) { ptr += 4 + size; continue; }
            payload[0] = sym_map[payload[0]];
            break;
        case 55:  /* CUDA_API_VERSION -- version compatibility check */
            validate_cuda_api_version(payload[0]);
            break;
        case 79:  /* AT_ENTRY_FRAGMENTS -- fragment type analysis */
            analyze_fragment_types(payload, size/4);
            break;
        default:
            break;
        }

        node = create_nvinfo_node(format, attr_code, nv_info_secidx);
        node->payload = payload;
        node->size = size;
    } else {
        /* Non-indexed: just create node with header, no payload copy */
        node = create_nvinfo_node(format, attr_code, nv_info_secidx);
        node->size = size;
        ptr += 4;  /* advance past header only, size already in node */
    }
    ptr = ptr + 4 + ALIGN_UP(size, 4);
}

EIATTR Attribute Catalog

nvlink v13.0.88 defines 97 EIATTR (ELF Info ATTRibute) codes, numbered 0 through 96. The name-to-code mapping was extracted directly from the pointer table at VA 0x1D37D60 in the nvlink binary (16-byte entries: 8-byte string pointer + 8-byte metadata, indexed by code). The string names reside at 0x1D36819--0x1D37170. Codes were verified against cubin TLV records produced by ptxas/nvcc v13.1 and cross-checked against the compute_entry_properties (sub_451D80) dispatch table. The validation function sub_42F760 confirms 97 codes by rejecting any attr_code > 0x60.

Name Table Structure

The EIATTR name table is a contiguous array of 97 x 16-byte entries at VA 0x1D37D60. Each entry:

Offset  Size  Field
------  ----  -----
0x00    8     name_ptr          Pointer to null-terminated string (e.g., "EIATTR_REGCOUNT")
0x08    4     min_toolkit_ver   Minimum CUDA toolkit version that supports this attribute
0x0C    4     usage_policy      0=warn, 1=error, 2=silently drop when version too old

The table is indexed by EIATTR code number. sub_42F760 reads dword_1D37D68[4 * attr_code] (the min_toolkit_ver field at offset +8) to check whether the current toolkit version (elfw+624) supports the attribute. If the toolkit version is too old, the usage_policy determines whether to emit a warning, error, or silently omit the record.

Complete Code Table (Sequential)

All 97 codes in numeric order. Use this as the authoritative reference when parsing .nv.info TLV records.

Most Important EIATTR Entries

This section provides detailed format information for the EIATTR entries that are most critical to GPU kernel launch, driver resource allocation, and the linker's own processing.

EIATTR_REGCOUNT (0x2F) -- Register Count

The single most important occupancy-determining attribute. The GPU driver computes max_warps_per_SM = total_registers / (regcount * warp_size) to determine how many warps can execute concurrently.

Format: 0x04 (Indexed)
On-disk layout (12 bytes total):
  Byte 0:     0x04    (indexed format)
  Byte 1:     0x2F    (EIATTR_REGCOUNT)
  Bytes 2-3:  0x0008  (size = 8 bytes)
  Bytes 4-7:  sym_idx (uint32, function symbol index)
  Bytes 8-11: regcount (uint32, physical register count per thread)

nvlink creates REGCOUNT records via sub_450B70(elfw, 0x2F, 8, payload_ptr, sym_idx, 0). The payload is an 8-byte pair: [sym_idx:4][regcount:4].

Linker operations on REGCOUNT:

1. Weak resolution (sub_45D180): When competing weak definitions exist, REGCOUNT is extracted to choose the winner. The definition with fewer registers wins, maximizing occupancy.

2. Weak stripping (sub_45D180): Bitmask 0x800800020000 marks REGCOUNT (bit 47), FRAME_SIZE (bit 17), and MAX_STACK_SIZE (bit 35) as resource attributes. When a weak symbol is replaced, records matching these codes for the discarded definition are zeroed (*(_BYTE *)(record + 1) = 0, overwriting attr_code to ERROR). The debug message "remove weak nvinfo" is printed when zeroing the entire nv.info section, and "remove weak frame_size" when zeroing individual resource records.

3. Propagation (sub_450ED0): propagate_register_counts walks the callgraph and propagates the maximum register count from callees to each entry kernel. If a callee uses more registers than the entry kernel, the entry kernel's REGCOUNT is raised to match. The verbose trace "regcount %d for %s propagated to entry %s" logs each propagation.

4. Validation: If no REGCOUNT is found for an entry function, the linker emits "no regcount?" via the error system. If a max-regcount-limited entry function calls a callee with a higher register count, it prints: "entry function '%s' with max regcount of %d calls function '%s' with regcount of %d".

EIATTR_NUM_BARRIERS (0x4C) -- Named Barrier Count

Controls how many named barrier slots the CTA hardware allocates. Most architectures support up to 16 barriers per CTA.

Format: 0x04 (Indexed) when read from input cubins
        0x02 (Value) when synthesized by the linker during finalization
On-disk layout (indexed, 12 bytes):
  Byte 0:     0x04    (indexed format)
  Byte 1:     0x4C    (EIATTR_NUM_BARRIERS)
  Bytes 2-3:  size
  Bytes 4-7:  sym_idx
  Bytes 8-11: barrier_count
Linker-synthesized layout (internal node):
  Byte 0:     0x02    (value format)
  Byte 1:     0x4C    (EIATTR_NUM_BARRIERS)
  Byte 2:     barrier_count (stored in the size field low byte)
  Byte 3:     0x00
  Bytes 4-7:  secidx  (output section index)

Barrier migration from section flags: The barrier count is also encoded in the section flags of .text sections as bits 26:20 (7 bits, mask 0x07F00000). During finalization, if a kernel has no EIATTR_NUM_BARRIERS record but its section flags carry a non-zero barrier count, the linker synthesizes one. The verbose message is:

"Creating new EIATTR_NUM_BARRIERS and moving barcount %d
 from section flags of %s to nvinfo for entry symbol %s"

After creating the nv.info record, the section flags are cleared: sh_flags &= 0xF80FFFFF.

Barrier propagation: If a callee requires more barriers than the entry kernel's current count, the linker propagates the higher count: "Propagating higher barcount %d to the section flags of %s of entry symbol %s".

EIATTR_FRAME_SIZE (0x11) -- Per-Thread Local Memory Frame

Format: 0x04 (Indexed)
Payload: [sym_idx:4][frame_size:4]
frame_size: bytes of local memory per thread (register spills + local arrays)

One of the three resource attributes stripped during weak symbol replacement (bitmask bit 17). The frame size determines the .nv.local.<funcname> section size visible to the driver.

EIATTR_MAX_STACK_SIZE (0x23) -- Maximum Stack Size

Format: 0x04 (Indexed)
Payload: [sym_idx:4][max_stack_bytes:4]

The worst-case per-thread stack allocation, computed by propagating CRS (Call-Return Stack) sizes through the callgraph. If the driver allocates less than this, the kernel will corrupt memory. If it allocates more, occupancy drops unnecessarily.

One of the three resource attributes stripped during weak replacement (bitmask bit 35).

EIATTR_MIN_STACK_SIZE (0x12) -- Minimum Stack Size

Format: 0x04 (Indexed)
Payload: [sym_idx:4][min_stack_bytes:4]

The non-recursive stack size minimum. Used when the callgraph has no recursion and the exact stack depth is statically known. For recursive kernels, the linker emits "Stack size for entry function '%s' cannot be statically determined".

EIATTR_CRS_STACK_SIZE (0x1E) -- Call-Return Stack Size

Format: 0x04 (Indexed)
Payload: [sym_idx:4][crs_bytes:4]

Size of the call-return stack for nested function calls. Propagated through the callgraph to compute per-entry worst-case stack requirements.

EIATTR_SAM_REGION_STACK_SIZE (0x3B) -- SAM Region Stack

Format: 0x04 (Indexed)
Payload: [sym_idx:4][sam_stack_bytes:4]

SAM (Streaming Asynchronous Memory) region stack size. Processed by sub_44C880 during compute_entry_properties (case 0x3B at line 1322 in sub_451D80). One of the resource attributes subject to symbol index remapping during merge (same case group as 17, 35, 47).

EIATTR_MAX_THREADS (0x05) -- Maximum Threads Per Block

Format: 0x04 (Indexed)
Payload: [sym_idx:4][max_threads:4]
max_threads: maximum threads per block (from .maxntid PTX directive)

EIATTR_REQNTID (0x10) -- Required Thread Count

Format: 0x04 (Indexed)
Payload: [sym_idx:4][reqntid:4]
reqntid: required thread count per dimension (from .reqntid PTX directive)

EIATTR_MAXREG_COUNT (0x1B) -- Maximum Register Hint

Format: 0x03 (Sized)
On-disk layout (4 bytes total, no payload):
  Byte 0:     0x03    (sized format)
  Byte 1:     0x1B    (EIATTR_MAXREG_COUNT)
  Bytes 2-3:  maxreg  (uint16, from --maxrregcount or .maxnreg)

This is a compiler hint, not an absolute limit. The value comes from --maxrregcount or the .maxnreg PTX directive. Uses sized format (0x03) -- the value is encoded directly in the 16-bit size field with no additional payload.

EIATTR_NUM_MBARRIERS (0x38) -- Memory Barrier Count

Format: 0x04 (Indexed)
Payload: [sym_idx:4][num_mbarriers:4]

Number of mbarrier objects used by the kernel. Processed in compute_entry_properties at case 0x38.

EIATTR_PARAM_CBANK (0x0A) -- Parameter Constant Bank

Format: 0x04 (Indexed)
Payload: [sym_idx:4][bank_offset:4]
bank_offset: packed {bank_number:16, offset:16}

Identifies which constant bank and at what offset kernel parameters begin. During merge, the symbol index is remapped with lazy symbol creation (case 10 in the merge switch). The linker may need to create a new symbol table entry if the constant bank symbol doesn't yet exist in the output.

EIATTR_CBANK_PARAM_SIZE (0x19) -- Constant Bank Parameter Size

Format: 0x03 (Sized)
Value: uint16 in the size field = parameter constant bank size in bytes

EIATTR_KPARAM_INFO (0x17) -- Kernel Parameter Info

Format: 0x01 (Free)
Payload: variable-length array of parameter descriptors
  Each descriptor: [ordinal:2][offset:2][size:2][?:2] (8 bytes per param)

Describes the type, size, and alignment of each kernel parameter. The v2 variant (EIATTR_KPARAM_INFO_V2, code 69/0x45) carries additional fields.

EIATTR_EXTERNS (0x0F) -- External Symbol References

Format: 0x01 (Free)
Payload: array of uint32 symbol indices
  Each 4-byte entry is one external symbol reference

During merge, every 4-byte entry in the payload is remapped through the symbol index translation table (case 15 in the merge switch). During finalization, compute_entry_properties creates new EXTERNS records via sub_450B70(elfw, 0x0F, 4*count, payload, root_kernel_sym, ...).

EIATTR_CUDA_API_VERSION (0x37) -- CUDA API Version

Format: 0x04 (Indexed)
Payload: [sym_idx:4][api_version:4]
api_version: encoded CUDA version (e.g., 0x83 = CUDA 13.1)

During merge, this attribute triggers a version compatibility check (case 55). If the input's CUDA API version exceeds the linker's current maximum (elfw+628), the linker emits an error. This prevents linking objects compiled for a newer CUDA version than the linker supports.

EIATTR_AT_ENTRY_FRAGMENTS (0x4F) -- Entry Fragments (Blackwell)

Format: 0x01 (Free)
Payload: array of uint32 fragment type descriptors

During merge, each 4-byte entry is analyzed for fragment types: values 4--5 indicate type 1, values 6--7 indicate type 2. The detected fragment type is stored at elfw+664. If conflicting fragment types are found across inputs, an error is emitted. During finalization, EXTERNS-style record creation via sub_450B70(elfw, 0x4F, size, payload, root_sym, ...).

EIATTR_RESERVED_SMEM_USED (0x41) -- Reserved Shared Memory

Format: 0x04 (Indexed)
Payload: [sym_idx:4][flags:4]

Processed in compute_entry_properties at case 0x41. The linker resolves the owning section through sub_442270 and checks if the referenced .nv.reservedSmem.* section has content. Related sections in the binary: .nv.reservedSmem.begin, .nv.reservedSmem.cap, .nv.reservedSmem.offset0, .nv.reservedSmem.offset1, .nv.reservedSmem.end.

EIATTR_SHARED_SCRATCH (0x32) -- Shared Scratch Space

Format: 0x04 (Indexed)
Payload: [sym_idx:4][scratch_size:4]

Shared memory scratch space used for register spilling when the register file is exhausted.

Attributes by Category

Resource Allocation (GPU Driver Critical)

These attributes directly control how the GPU driver allocates hardware resources for kernel launch. Incorrect values cause silent performance degradation or launch failure.

Parameter Bank Layout

These describe how kernel parameters are laid out in constant memory bank 0 (c[0x0]).

Instruction Offset Tables

These record byte offsets of specific instruction types within the kernel's .text section, enabling the driver and tools to locate and patch instructions at load time.

Texture and Surface Binding

Cluster and Cooperative Launch (sm_90+)

Shared Memory and Reserved Resources

Software Workarounds

Hardware errata requiring instruction-level patching by the driver.

Compilation Metadata

Graphics-Specific

Blackwell+ Features (sm_100+)

Mercury-Specific

Sentinel and Error

How nvlink Processes .nv.info

Phase 1: Merge (Input Processing)

When merge_elf (sub_45E7D0) encounters a section with sh_type == 0x70000000 (line 1854 of decompiled code), it enters the .nv.info processing path. This path:

1. Resolves the target section: If sh_link != 0, the symbol is looked up via the input symbol table, mapped to the output symbol index, and the per-function .nv.info.<name> section is created or found via sub_4504B0.

2. Walks TLV records sequentially: The parser runs while (ptr < end) where end = section_data + sh_size. For each record, it reads format = ptr[0], attr_code = ptr[1], size = *(uint16*)(ptr+2).

3. Dispatches on format byte: If format == 0x04 (indexed), the payload is copied to arena memory and symbol indices within the payload are remapped. If format != 0x04, the record header is stored directly (no payload copy needed for non-indexed records).

4. Per-attribute symbol remapping (format 0x04 only): A switch on attr_code determines how symbol indices within the payload are translated:

5. Weak symbol handling: For already-processed weak symbols, the entire .nv.info section is skipped with verbose message "weak %s already processed".

Phase 2: Weak Symbol Resolution

merge_weak_function (sub_45D180) strips resource attributes from the discarded weak definition. It walks the output's nv.info linked list at elfw+392 and for each node:

int64_t bitmask = 0x800800020000LL;  // bits 17, 35, 47
for (node = list_head; node != NULL; node = node->next) {
    if (node->secidx == target_nv_info_secidx) {
        // Section matches the discarded weak -- zero the entire record
        if (verbose) fprintf(stderr, "remove weak nvinfo\n");
    } else {
        uint8_t code = node->attr_code;
        if (code <= 0x2F && _bittest64(&bitmask, code)
            && payload[0] == target_sym_idx) {
            // Resource attribute for the discarded symbol -- zero it
            if (verbose) fprintf(stderr, "remove weak frame_size\n");
        } else {
            continue;  // keep this record
        }
    }
    node->attr_code = 0;  // overwrite to EIATTR_ERROR (suppressed)
}

The three codes matching bitmask 0x800800020000:

• Bit 17 = EIATTR_FRAME_SIZE (0x11)
• Bit 35 = EIATTR_MAX_STACK_SIZE (0x23)
• Bit 47 = EIATTR_REGCOUNT (0x2F)

The rationale: when a weak function is replaced, the replacement's resource descriptors are the ones that matter. The discarded definition's REGCOUNT, FRAME_SIZE, and MAX_STACK_SIZE would be incorrect for the replacement and must not contaminate the output.

Phase 3: Pre-Finalization (Callgraph Processing)

sub_44C030 (the EIATTR serialization builder, called from sub_44DB00 in the pre-finalization fixup) processes the callgraph to prepare nv.info data:

1. Callgraph traversal: Iterates all callgraph entries starting at elfw+408. For each function pair (caller, callee), resolves the callee's callgraph node and links it to the caller's callee list.

2. Recursion detection: Uses a two-flag approach (bytes at offsets +48 and +49 in callgraph nodes). If recursion is detected, the function is marked as recursive ("recursion at function %d"). Recursive functions cannot have statically-determined stack sizes.

3. Callee list propagation: For each entry function, the transitive closure of callees is computed. Each callee's entry is added to the entry function's "reachable" list via sub_4644C0.

Phase 4: Finalization (Output Generation)

compute_entry_properties (sub_451D80, 97,969 bytes -- the largest function in the linker) runs during the finalization phase. It operates on all nv.info records and computes derived properties for each kernel entry point:

4a. Symbol index fixup: For indexed records with codes 2, 6--10, 17--20, 23, 35, 38, 47, 59, 69 (the same set as in merge remapping), the function re-resolves symbol indices through the positive/negative symbol mapping tables (elfw+456/elfw+464). If a symbol was deleted but the record still references it, attr_code is set to 0 (suppressed).

4b. Register count propagation (sub_450ED0): Walks the callgraph and propagates the maximum register count from callees to each entry kernel.

for each entry_function in callgraph:
    regcount = entry_function.regcount;
    for each callee in transitive_callees(entry_function):
        if callee.regcount > regcount:
            /* Raise entry's count to match */
            entry_function.regcount = callee.regcount;
            if (verbose)
                fprintf(stderr, "regcount %d for %s propagated to entry %s\n",
                    callee.regcount, callee.name, entry_function.name);

If no REGCOUNT record exists for an entry function (e.g., it was a leaf function with the count stored in section flags), the linker creates one via sub_450B70(elfw, 0x2F, 8, payload, 0, 0). The payload is [sym_idx:4][regcount:4].

4c. Barrier count creation and propagation: When a kernel's section flags contain a barrier count (bits 26:20 of sh_flags) but no EIATTR_NUM_BARRIERS record exists, the function creates one:

"Creating new EIATTR_NUM_BARRIERS and moving barcount %d
 from section flags of %s to nvinfo for entry symbol %s"

The linker then clears the barrier bits from section flags (sh_flags &= 0xF80FFFFF) because the nv.info record is now the authoritative source.

If a callee's barrier count exceeds the entry kernel's:

"Propagating higher barcount %d to the section flags of %s of entry symbol %s"
"Propagating higher barcount %d to the section %s of entry symbol %s"

4d. EXTERNS and AT_ENTRY_FRAGMENTS creation: For entry functions, the finalization phase creates EIATTR_EXTERNS (0x0F) and EIATTR_AT_ENTRY_FRAGMENTS (0x4F) records for the transitive closure of external references and fragment descriptors.

4e. Encoding: nvinfo_encode (sub_468760, 14,322 bytes) serializes the in-memory linked-list nodes into on-disk TLV records using SSE2/AVX intrinsics for efficient bitfield packing. The encoder processes a descriptor array of 4 x uint32 entries per field, with the third uint32 serving as a type code that dispatches to one of ~15 encoding strategies (cases 0 through 0x12 in the encoder's switch).

Emission Subsystem (Embedded ptxas)

When nvlink performs LTO compilation, the embedded ptxas compiler generates .nv.info attributes through a dedicated emission subsystem at 0x15C5000--0x160FFFF. This subsystem is the single largest code region dedicated to .nv.info processing in the entire binary.

Architecture

The emission pipeline has three layers:

Layer 1: SM dispatch (sub_15C0CE0). A singleton initialization function registers per-SM callback tables for 12 architecture families (sm_75 through sm_121). Each SM gets an nv.info emitter callback looked up through map A8 via sub_15C3DB0. The callback creates a ~1,936-byte codegen state with architecture-specific constants at offsets 344 and 348 (compute capability encoding).

Layer 2: Master emitters (4 functions, 78--55 KB each). These are the top-level attribute-lowering functions that read compilation state and dispatch to per-attribute-type handlers:

The master emitters use an FNV-1a hash table (offset basis 0x811C9DC5, prime 16,777,619) at object+488 for O(1) function-ID-to-attribute lookup using 24-byte entries.

Layer 3: Per-attribute handlers (~190 functions at 0x15CF070--0x160FFFF). Each function is 4--8 KB and handles exactly one EIATTR type. They follow a uniform template:

1. Read attribute descriptor pointer from a2
2. Read sub-attribute fields from known offsets (m128i-based, 32-byte descriptors)
3. Call sub_A4CBB0 to create attribute IR node
4. Call sub_A49120 to set EIATTR type code
5. Call sub_A49190/sub_A49140 for type validation
6. Write output via sub_4A3D60 (operand builder)
7. Return constructed attribute list

The uniformity across ~190 functions suggests they are generated from a data-driven table or macro expansion in NVIDIA's source code, one handler per EIATTR type.

Supporting Functions

How .nv.info Drives GPU Resource Allocation

The .nv.info section is not just metadata for tools -- it is the primary input to the GPU driver's kernel launch resource allocator. The relationship is:

1. Register allocation: EIATTR_REGCOUNT tells the driver how many registers each thread needs. The driver computes: max_warps_per_SM = total_registers / (regcount * warp_size). This is the single most important occupancy-determining attribute.

2. Shared memory reservation: EIATTR_SMEM_PARAM_SIZE and EIATTR_RESERVED_SMEM_0_SIZE determine how much shared memory to carve out before the kernel's dynamic shared memory allocation.

3. Stack allocation: EIATTR_CRS_STACK_SIZE and EIATTR_MAX_STACK_SIZE determine per-thread stack allocation. If the driver gets this wrong (too small), the kernel will corrupt memory; if too large, occupancy drops.

4. Barrier reservation: EIATTR_NUM_BARRIERS reserves named barrier slots. On most architectures the hardware supports 16 barriers per CTA. The driver must configure the barrier hardware before launch.

5. Instruction patching: The offset tables (EIATTR_EXIT_INSTR_OFFSETS, EIATTR_S2RCTAID_INSTR_OFFSETS, EIATTR_SW*_WAR) tell the driver which instruction words to patch. This enables hardware workarounds and CTA-ID remapping for cluster launch without recompilation.

6. Cluster configuration: EIATTR_CTA_PER_CLUSTER and EIATTR_EXPLICIT_CLUSTER (sm_90+) control the cluster launch hardware, determining how many CTAs share distributed shared memory.

Binary Artifacts

Typos Preserved in the Binary

A corrected variant EIATTR_AT_ENTRY_FRAGMENTS also exists at 0x245E8D9, suggesting awareness of the typo but preservation of the original for backward compatibility. Similarly, EIATTR_COROUTINE_RESUME_ID_OFFSETS at 0x245F010 is an alternate name for code 58 (EIATTR_COROUTINE_RESUME_OFFSETS), used in the embedded ptxas compiler.

Diagnostic Strings

"nvinfo <fmt=%d,attr=%d,size=%d>, secidx=%d"            (sub_4478F0 debug dump)
"no new register count found for %s, checking .nv.info"  (0x1D3BA68)
"no original register count found for %s, checking .nv.info" (0x1D3BAA0)
"regcount %d for %s propagated to entry %s"              (0x1D3B070)
"no regcount?"                                           (0x1D3AF6C)
"entry function '%s' with max regcount of %d calls function
 '%s' with regcount of %d"                               (0x1D39930)
"Creating new EIATTR_NUM_BARRIERS and moving barcount %d
 from section flags of %s to nvinfo for entry symbol %s" (0x1D3AF88)
"Propagating higher barcount %d to the section flags of
 %s of entry symbol %s"                                  (sub_450ED0)
"remove weak nvinfo"                                     (sub_45D180 debug)
"remove weak frame_size"                                 (sub_45D180 debug)
"weak %s already processed"                              (sub_45E7D0 merge)
"error while updating global nvinfo section"             (embedded ptxas)
"Stack size for entry function '%s' cannot be
 statically determined"                                  (sub_451D80)

Cross-References

• NVIDIA Section Types -- SHT_CUDA_INFO definition and section flag encoding
• Section Catalog -- .nv.info and .nv.info.<funcname> catalog entries
• Weak Symbol Handling -- Register count extraction from .nv.info during weak resolution
• Section Merging -- .nv.info section merge with symbol index remapping
• Finalization Phase -- compute_entry_properties and propagate_register_counts
• Architecture Dispatch -- Per-SM nv.info emitter callback registration
• Constant Banks -- EIATTR_PARAM_CBANK and EIATTR_CBANK_PARAM_SIZE interaction

Sibling wikis (ptxas):

• ptxas: Section Catalog & EIATTR -- ptxas-side EIATTR emission pipeline and section type constants
• ptxas: EIATTR Reference -- EIATTR attribute code reference from the ptxas perspective

Confidence Assessment