Debug Info Verification

cicc includes a custom debug info verification pass (sub_29C8000) that validates DWARF-like debug metadata after each optimization pass in the pipeline. This is not the upstream LLVM IR Verifier (llvm::Verifier::verify(Module)); it is an NVIDIA-specific implementation derived from LLVM's CheckDebugInfoPass (in Debugify.cpp) with two significant extensions: a structured JSON reporting mechanism that tracks exactly which optimization passes degrade debug info quality, and a configurable verbosity system that allows the verification overhead to be tuned from silent to exhaustive. The pass lives in a self-contained module of approximately 93 functions in the 0x29C0000--0x29FFFFF address range, alongside the Debugify synthetic debug info injector and general pass infrastructure utilities. Its purpose is to ensure that when a developer compiles with -g or -generate-line-info, the debug metadata that cuda-gdb and Nsight Compute rely on survives the aggressive optimization pipeline intact.

Three Verification Modes

cicc supports three independent verification protocols, each activated by a different set of knobs. Understanding which protocol is active determines what diagnostic output to expect and how much overhead the verification adds.

Mode 1: Post-Pass Debug Info Verification (verify-each)

The default verification mode, activated by the verify-each (or its alias verify-after-all) LLVM knob. The pipeline runner invokes sub_29C8000 as a sandwich around each optimization pass:

// Pseudocode for the pipeline runner's verification protocol
// (entry: 0x29C8000, stack: 0x4B8 bytes)
snapshot_debug_metadata(M);
run_optimization_pass(M, "instcombine");
sub_29C8000(M, errs(), dbgCU, hashMap, "instcombine", 11, file, fileLen, jsonOut);

The pass name argument identifies which optimization just ran, so the JSON report can attribute any debug info degradation to the specific pass responsible. The verifier checks the full metadata inventory: subprograms, scopes, variables, types, labels, imported entities, and retained nodes. It produces ERROR diagnostics for dropped subprograms and WARNING diagnostics for dropped debug variable intrinsics.

Activation: -Xcicc -verify-each or -Xcicc -verify-after-all Overhead: One full metadata snapshot + eight hash table constructions + per-function variable scan per optimization pass. Substantial for large modules.

Mode 2: Debugify Synthetic Injection + Verification (debugify-each)

The full Debugify cycle injects synthetic debug metadata before each pass, runs the pass, then verifies the synthetic metadata survived. This mode is more aggressive than Mode 1 because it tests every pass even on code compiled without -g.

// Debugify cycle pseudocode
sub_29C1CB0(M, "llvm.debugify");   // inject synthetic debug info
run_optimization_pass(M, "instcombine");
sub_29C8000(M, errs(), dbgCU, hashMap, "instcombine", ...);  // verify
strip_debugify_metadata(M, "llvm.debugify");  // cleanup

The injector (sub_29C1CB0) creates "llvm.debugify" / "llvm.mir.debugify" named metadata nodes that serve as watermarks. The checker looks for these watermarks to distinguish synthetic from genuine debug info.

Activation: -Xcicc -debugify-each Sub-knobs: debugify-level (locations or location+variables), debugify-quiet, debugify-func-limit, debugify-export

Mode 3: Debug Info Preservation Checking (verify-debuginfo-preserve)

A lighter-weight mode that checks only whether existing debug info survives optimization, without injecting synthetic metadata. This mode is available through the New Pass Manager infrastructure and can export results via verify-di-preserve-export.

Activation: -Xcicc -verify-debuginfo-preserve Sub-knobs: verify-each-debuginfo-preserve, verify-di-preserve-export

Mode Selection Matrix

Pipeline Integration

The verifier operates as an interleaved "check" pass. The New Pass Manager registers it via two wrappers in the pipeline construction code at 0x2270000--0x227FFFF:

The pipeline text parser (sub_2272BE0, 14KB) recognizes these as named module passes:

When debugify-each is active, the pipeline builder (sub_2277440, 60KB -- buildDefaultPipeline() equivalent) wraps every optimization pass in a debugify/check-debugify pair. When verify-each is active, only the check-debugify wrapper is inserted.

Verification Function Signature

The function signature reconstructed from the binary:

bool sub_29C8000(
    Module*       module,       // rdi
    raw_ostream&  output,       // rsi -- diagnostic stream
    NamedMDNode*  dbgCU,        // rdx -- "llvm.dbg.cu" metadata
    DenseMap*     hashMap,      // rcx -- metadata identity table
    const char*   passName,     // r8
    size_t        passNameLen,  // stack+0x00
    const char*   fileName,     // stack+0x08
    size_t        fileNameLen,  // stack+0x10
    raw_ostream*  jsonOutput,   // stack+0x18 -- NULL if no JSON report
    ...
);
// Returns: true = all checks passed, false = any violation detected

Verification Algorithm

The pass proceeds through nine sequential phases within a single function call. The 0x4B8-byte stack frame holds eight separate tracking data structures.

Phase 1: Module-Level Guard (0x29C8000 -- 0x29C807A)

Looks up the "llvm.dbg.cu" named metadata node via sub_BA8DC0 (Module::getNamedMetadata). If absent or empty, prints ": Skipping module without debug info\n" and returns 0. This is the fast path for modules compiled without -g.

Phase 2: Pre-Pass Metadata Snapshot (0x29C8080 -- 0x29C8AE5)

Initializes eight SmallVector/DenseMap structures on the stack and walks the compile unit metadata tree:

For each DICompileUnit operand, the pass walks the subprogram list and retained types, recording every metadata node in hash tables for O(1) identity comparison. The hash function is:

uint64_t hash = ((ptr >> 4) ^ (ptr >> 9)) & (bucket_count - 1);

This is the standard DenseMap pointer hash with LLVM-layer sentinels. See Hash Table and Collection Infrastructure for the complete specification.

Phase 3: DISubprogram Iteration (0x29C82BE -- 0x29C84C8)

Walks the subprogram list attached to each compile unit via linked-list traversal ([node+8] = next pointer). For each subprogram, reads the metadata tag byte at [node-18h]:

The flag byte at [rdx+21h] & 0x20 tests the "definition" bit (only defined, non-declaration subprograms are tracked). Values outside 0x44--0x47 are flagged as invalid scope types.

Phase 4: Hash Table Construction (0x29C8508 -- 0x29C8AC2)

Allocates and populates eight sorted hash tables via sub_C7D670 (aligned_alloc, alignment=8), each holding 16-byte entries [pointer, secondary_key]:

The MDNode operand access pattern used during population:

// MDNode internal layout decoding (0x29C8508+)
byte flags = *(ptr - 0x10);
if (flags & 0x02) {          // distinct metadata
    operands = *(ptr - 0x20);  // operand array is before the node
} else {
    int count = (flags >> 2) & 0x0F;
    operands = ptr - 0x10 - (count * 8);  // inline operands
}

Phase 5: Per-Function Debug Variable Checking (0x29C8B3B -- 0x29C9060)

Iterates every function in the module. For each, looks up its DISubprogram in the hash table and cross-references dbg.value() / dbg.declare() intrinsics against the pre-snapshot. Two diagnostic levels:

ERROR (pass dropped a subprogram entirely):

ERROR: <pass> dropped DISubprogram of <function> from <file>
ERROR: <pass> did not generate DISubprogram for <function> from <file>

WARNING (pass dropped individual variable tracking):

WARNING: <pass> drops dbg.value()/dbg.declare() for <var> from function <func> (file <file>)

The distinction between "dropped" and "did not generate" is significant: "dropped" means metadata existed before the pass and was deleted; "not-generate" means the pass created new IR (e.g., from inlining or outlining) without attaching corresponding debug metadata. This taxonomy is important for GPU compilation because kernel outlining and device function inlining frequently create new IR nodes.

The variable name is resolved by:

1. Getting DISubprogram from the metadata ref
2. Calling sub_AF34D0 (DIScope::getScope()) to walk the scope chain upward
3. Getting the file via operand [10h] of the scope's file ref
4. Calling sub_B91420 (MDString::getString()) to convert MDString to StringRef

Phase 6: Per-Instruction Location Verification (0x29C8D42 -- 0x29C8D85)

Delegated to sub_29C3AB0 (5,592 bytes), which performs detailed checks:

• Every instruction with a DebugLoc has a valid DILocation
• DILocation scope chains resolve to a valid DISubprogram
• No orphaned debug locations reference deleted subprograms
• BB-level consistency: all instructions in a basic block share compatible scopes
• Dropped location tracking: emits "dropped DILocation" diagnostics

The JSON output from this sub-pass uses structured field names: "DILocation", "bb-name", "fn-name", "action" (with values "drop" or "not-generate").

Phase 7: JSON Structured Output (0x29C90BC -- 0x29C94E2)

When a non-null JSON output stream is provided (the jsonOutput parameter), the pass serializes a structured report via sub_2241E40 (YAML/JSON serializer):

{"file":"kernel.cu", "pass":"instcombine", "bugs": [
  {"metadata":"DISubprogram", "name":"_Z6kernelPf", "fn-name":"_Z6kernelPf", "action":"drop"},
  {"metadata":"dbg-var-intrinsic", "name":"idx", "fn-name":"_Z6kernelPf", "action":"not-generate"}
]}

This JSON reporting mechanism is an NVIDIA extension with no upstream LLVM equivalent. It feeds into NVIDIA's internal CI infrastructure to track debug info quality regressions across compiler versions. The "no-name" string serves as fallback when the pass name pointer is NULL.

The serialization calls sub_CB7060 (YAML::IO constructor) and proceeds through sub_C6D380 (object emission), sub_C6C710 (array emission), and sub_C6B0E0 (key writer). After serialization, the stream is flushed via sub_CB7080 and freed via sub_CB5B00. If the file descriptor is valid (fd != -1), it is closed via sub_C837B0 (close(fd)).

Phase 8: Result Reporting and Metadata Reconstruction (0x29C94E2 -- 0x29C9A27)

Prints the summary line ("<pass>: PASS\n" or "<pass>: FAIL\n"), then reconstructs the module's metadata tables from the verified versions -- reallocating subprogram, type, variable, label, and global variable arrays and copying verified metadata back into the compile unit structures.

The result is a 3-way outcome in bit flags (combined at 0x29C9073--0x29C9080 via AND):

• Bit 0: any verification failure (determines PASS/FAIL)
• Bit 1: JSON report was requested and successfully written

The final result is PASS only if all sub-checks passed AND the JSON report (if requested) was successfully written.

Cleanup frees all eight temporary hash tables (each via sub_C7D6A0 -- sized dealloc with alignment 8), linked list nodes via j_j___libc_free_0, and SmallVector inline buffers are detected by pointer comparison (if ptr == stack_addr, skip free).

Phase 9: Return (0x29C9A12 -- 0x29C9A27)

Returns var_420 (bool) in the al register. Standard epilog restores rbx, r12--r15, rbp.

Complete Diagnostic Code Table

Every diagnostic string emitted by the debug verification subsystem, with exact provenance and trigger conditions.

Verification Pass Diagnostics (sub_29C8000)

Per-Instruction Verifier Diagnostics (sub_29C3AB0)

JSON Report Field Schema

Action Value Taxonomy

String Encoding Details

Several diagnostic strings are constructed inline using immediate mov instructions rather than string table references:

These inline immediate constructions avoid string table lookups and are a common LLVM raw_ostream optimization for short fixed strings.

Compile Unit Descriptor Layout

The verification pass reads and reconstructs a per-CU descriptor object (referenced at [rbp+var_440]) with the following layout:

DISubprogram Node Layout

Accessed during Phase 3 scope chain validation:

Debugify Injector (sub_29C1CB0)

The Debugify injector creates synthetic debug metadata to test whether optimization passes preserve debug info correctly. It is the counterpart to the verifier -- the injector sets up the watermarks, and the verifier checks them.

Named metadata markers:

• "llvm.debugify" -- marks the module as containing synthetic debug info (standard Debugify)
• "llvm.mir.debugify" -- marks MIR-level synthetic debug info

Behavior controlled by debugify-level:

• locations -- inject only DILocation on every instruction (cheaper, tests location preservation)
• location+variables -- inject DILocation plus synthetic dbg.value()/dbg.declare() for every SSA value (full coverage, higher overhead)

The injector assigns monotonically increasing line numbers to every instruction and creates one DILocalVariable per SSA value that produces a result. The variable names follow the pattern "dbg_var_N" where N is the SSA value index. After injection, the module has guaranteed 100% debug coverage, making any coverage loss attributable to the subsequent optimization pass.

Verbosity Control

Two global flags provide fine-grained control over verification output:

qword_5008FC8 -- Verbose Diagnostic Output Enable

Boolean flag (byte). Controls the output stream selection:

• When 0: uses sub_CB72A0 (null/discard stream constructor) -- diagnostics silently discarded
• When non-zero: uses sub_CB7330 (stderr stream accessor) -- diagnostics printed to stderr

This flag gates the ERROR and WARNING messages. The JSON structured output is controlled separately by the jsonOutput parameter. Setting qword_5008FC8 = 0 suppresses text diagnostics while still producing JSON output.

qword_5008C88 -- Metadata Depth Threshold

Signed 32-bit integer, read at 0x29C8371. Controls how deep the scope chain walk goes:

• When <= 0: the deep scope chain walk is skipped for non-subprogram metadata. Only top-level DISubprogram validation runs.
• When > 0: full scope chain traversal validates every DILexicalBlock, DILexicalBlockFile, and DINamespace in the hierarchy.

This allows production builds to run lightweight verification (subprogram-only) while development builds run exhaustive scope chain checking.

Debugify-Specific Knobs

GPU Debug Info: What PTX Needs

DWARF for PTX differs fundamentally from DWARF for x86. PTX is a virtual ISA -- there are no physical registers, no real stack, and no fixed instruction encoding. The debug metadata cicc emits serves two consumers: cuda-gdb (which maps PTX locations back to source) and ptxas (which carries debug info forward into SASS/ELF for the hardware debugger).

The .loc Directive

The AsmPrinter (sub_31D55F0) emits DWARF .loc directives before each PTX instruction that has a valid DebugLoc:

.loc 1 42 0          // file 1, line 42, column 0
ld.param.u64 %rd1, [_Z6kernelPf_param_0];
.loc 1 43 5
mul.wide.u32 %rd2, %r1, 4;

The .file directives (sub_31E4280) establish the file table, and sub_31E6100 maintains a file/line-to-MCSymbol mapping for line table construction.

The dwarf-extended-loc knob (enum: Default/Enable/Disable, registered at 0x490000 range) controls whether extended flags appear in .loc directives. When disabled, cicc emits bare .loc file line column without the is_stmt, prologue_end, or discriminator extensions. This is relevant because older ptxas versions do not parse extended .loc flags.

The line-info-inlined-at Extension

The -line-info-inlined-at LLVM knob (registered at ctor_043 / 0x48D7F0, exposed as -no-lineinfo-inlined-at in the cicc CLI, which sets -line-info-inlined-at=0 on the backend) controls whether inlined-at chains are preserved in PTX line info. When enabled (the default), every .loc directive for inlined code carries the full inlining chain so cuda-gdb can reconstruct the call stack at any point in the inlined code. When disabled, only the immediate source location is emitted, losing the inlining context but producing smaller PTX.

The -show-src / nvptx-emit-src Feature

The -show-src CLI flag (stored at flag struct offset +808, routed to the backend as -nvptx-emit-src) enables source line interleaving in PTX output. When active, the AsmPrinter annotates each .loc directive with the corresponding source line as a PTX comment:

// kernel.cu:42    float val = input[idx];
.loc 1 42 0
ld.global.f32 %f1, [%rd2];
// kernel.cu:43    val = val * val;
.loc 1 43 0
mul.f32 %f2, %f1, %f1;

This is purely a readability feature for developers inspecting PTX output. It has no effect on cuda-gdb or debug quality -- the source text is embedded as comments that ptxas ignores.

NvvmDebugVersion

The NVVM container format includes a debug version field (NvvmDebugVersion, packed as {Major:uint16, Minor:uint16} at container offset 0x08--0x09). The current version is Major=3, Minor<=2. The reader (sub_CD41B0) validates that Major equals 3 and warns if Minor exceeds 2. If absent, the default {3, 2} is assumed. This version tracks the debug metadata schema independently of the NVVM IR version, allowing debug format evolution without breaking IR compatibility.

The standalone pipeline (sub_12BFF60) performs a consistency check: if the container declares debug_info_present (bit 4 of flags) AND the debug mode flag is set AND the debug version has not been validated, it returns error code 3 (incompatible).

DbgRecord Format (LLVM 20)

cicc v13.0 uses LLVM 20's DbgRecord format by default (write-experimental-debuginfo = true, registered at ctor_025). This replaces traditional dbg.value()/dbg.declare() intrinsics with non-intrinsic debug records attached directly to instructions. Related knobs:

The verifier handles both formats: it checks for dbg.value()/dbg.declare() intrinsics AND for DbgRecord attachments.

Debug Info Stripping Passes

cicc includes five stripping passes registered in the pipeline parser (at sub_12C6910 and related):

The strip-nonlinetable-debuginfo pass is the key one for the -generate-line-info mode: it strips all debug metadata except .loc / .file directives, producing line-number-only debug info without variable locations, type descriptions, or scope trees. This is what nvcc's --generate-line-info flag triggers -- enough for profiler source correlation but not enough for stepping through code in cuda-gdb.

The core debug info stripping implementation lives at 0xAE0000 (Zone 3 of the type system module), which calls stripDebugInfo() to remove all llvm.dbg.* intrinsics from the module.

Debug Compilation Modes

cicc supports three debug info levels, controlled by CLI flags that route through the flag dispatch table:

When -g is active, cicc emits DICompileUnit with full emission kind, preserves all DISubprogram, DILocalVariable, DIType, and scope metadata through the pipeline, and the backend emits complete DWARF sections. The verifier runs at full depth.

When -generate-line-info is active, the StripNonLineTableDebugInfoPass runs early in the pipeline, leaving only line table metadata. The verifier still runs but only checks DILocation / DISubprogram consistency (variable checks are skipped because the variable metadata was intentionally stripped).

Key routing difference: -g routes to BOTH the linker (-debug-compile) and optimizer (-debug-compile), because libdevice linking needs the debug flag to preserve user debug info during merging. -generate-line-info routes to the optimizer only.

The frontend uses two independent guard mechanisms for debug emission:

• dword_4D046B4 -- global flag checked at statement/parameter level by sub_9433F0 (per-param debug), sub_943430 (per-global debug)
• [ctx+0x170] -- compile unit pointer checked at module finalization level by sub_915400

The NVVM container carries a dedicated DebugInfo enum (3 values: NONE, LINE_INFO, DWARF) at deserialized struct offset +12, separate from the module metadata.

Complete Knob Reference

DWARF Emission Backend

The actual DWARF section emission lives in a separate module at 0x3990000--0x39DF000:

The module-level entry sub_215ACD0 checks *(a1+240)->field_344 to determine if DWARF is enabled, then looks up the "NVPTX DWARF Debug Writer" / "NVPTX Debug Info Emission" pass info. The NVPTX backend does not emit physical register locations (GPUs have no DWARF register numbering scheme that maps to hardware); instead, it emits virtual register references that cuda-gdb resolves through ptxas's SASS-level debug info.

Function Map

LLVM Infrastructure Functions Used

NVIDIA Modifications vs Stock LLVM

The key differences from upstream LLVM's CheckDebugInfoPass:

1. JSON structured output -- Upstream only prints text diagnostics. NVIDIA added a YAML/JSON serializer (sub_2241E40, sub_CB7060) that produces machine-parseable bug reports with "file", "pass", "bugs" fields and per-bug "action" classification ("drop" vs "not-generate").

2. Verbosity control -- Two global flags (qword_5008FC8 for output enable, qword_5008C88 for depth threshold) allow fine-grained control over verification overhead. Upstream has only the debugify-quiet knob.

3. Eight-table metadata tracking -- Upstream CheckDebugInfoPass tracks DISubprograms and debug variable intrinsics. NVIDIA's version maintains eight separate hash tables covering subprograms, scopes, global variables, local variables, types, imported entities, labels, and retained nodes -- a much more comprehensive snapshot.

4. Metadata reconstruction -- After verification, NVIDIA's pass reconstructs the module's metadata tables from the verified versions (Phase 8), which upstream does not do. This means the verifier can also serve as a "repair" pass that normalizes metadata after an optimization pass corrupts it.

5. No kernel-specific handling -- The verifier treats __global__ and __device__ functions identically. CUDA-specific debug info (address space annotations, shared memory debug, warp-level location info) is validated elsewhere, likely during NVPTX backend emission.

6. DbgRecord format support -- cicc v13.0 defaults to the LLVM 20 DbgRecord format (write-experimental-debuginfo = true), so the verifier handles both intrinsic-based and record-based debug info transparently.

Cross-References

• AsmPrinter & PTX Body Emission -- .loc/.file directive emission, per-instruction debug annotation
• PTX Emission -- module-level emission entry, DWARF debug writer lookup
• Debug Info Pipeline -- end-to-end debug info flow from frontend to backend
• CLI Flags -- -g, -generate-line-info, -show-src flag routing
• LLVM Knobs -- debugify-*, verify-each, dwarf-* knobs
• Pipeline & Ordering -- where debug verification fits in the pass pipeline
• Hash Infrastructure -- DenseMap/DenseSet implementation used by tracking tables
• Diagnostics -- broader diagnostic and remark system
• NVVM Container -- NvvmDebugVersion field