CICC v13.0 — Reverse Engineering Reference

CICC is NVIDIA's CUDA C-to-PTX compiler — the binary that transforms CUDA C++ source code (or LLVM bitcode) into PTX assembly for GPU execution. At 60 MB, it is one of the largest single compiler binaries in production use. This wiki documents its internal architecture, recovered from static analysis of the stripped x86-64 ELF binary using IDA Pro 8.x and Hex-Rays decompilation.

Three Subsystems

CICC is not a monolithic compiler. It is composed of three largely independent subsystems, each with its own lineage, coding conventions, and internal data structures:

1. EDG 6.6 C++ Frontend (3.2 MB, 0x5D0000–0x8F0000) — A licensed commercial frontend from Edison Design Group that parses CUDA C++ source code and emits transformed C code. It operates as a source-to-source translator: CUDA kernel launch syntax (<<<>>>) is lowered to CUDA runtime API calls, memory space qualifiers (__shared__, __constant__) are resolved to address space annotations, and C++ templates/constexpr are fully evaluated. The output is not LLVM IR — it is C code that feeds into a second compilation phase. See EDG 6.6 Frontend.

2. NVVM Bridge (~4 MB, 0x8F0000–0x12CFFFF) — The glue layer between EDG and LLVM. It handles CLI parsing, architecture detection (23 SM variants with 3-column flag fan-out), the dual-path compilation dispatch (Path A via LibNVVM API, Path B standalone), the NVVMPassOptions knob system (221 per-pass configuration slots), and the 770-entry builtin resolution table. This layer is entirely NVIDIA-proprietary. See Entry Point & CLI and LLVM Optimizer.

3. LLVM 20.0.0 Backend (~45 MB, 0x12D0000–0x3BFFFFF) — A heavily modified LLVM fork that performs IR optimization and PTX code generation. NVIDIA has added 35 custom passes (MemorySpaceOpt, Rematerialization, BranchDist, LoopIndexSplit, Sinking2, etc.), a proprietary two-phase compilation model with per-function thread parallelism, and extensive modifications to the NVPTX backend for tensor core code generation across 5 GPU architecture generations. See Code Generation and PTX Emission.

Additionally, jemalloc 5.3.x (~400 functions at 0x12FC000) is statically linked, replacing the system allocator for improved memory allocation performance during compilation.

Dual-Path Architecture

A distinctive feature of cicc is its dual-path design — two complete copies of the compilation backend exist within the same binary, selected at runtime:

Runtime selection is controlled by v253 in sub_8F9C90 (the real main function). The default value (2) triggers an environment variable lookup through an obfuscated string comparison to determine which path to take. This design allows a single binary to serve both the nvcc driver toolchain and the LibNVVM runtime compilation API.

Compilation Pipeline

Both paths converge on the same 5-stage pipeline:

CUDA C++ Source (.cu / .ci / .i)
  │
  ├─ EDG 6.6 Frontend (sub_5D2A80)
  │   ├─ lgenfe_main (sub_617BD0): 282-case CLI, 737 #defines
  │   ├─ Parser: recursive-descent + declaration specifier state machine
  │   ├─ Constexpr evaluator: 317KB tree-walking interpreter
  │   └─ Backend: "Generating NVVM IR" → .int.c / .device.c / .stub.c
  │
  └─ NVVM/LLVM Pipeline
      │
      ├─ IRGEN:  EDG IL → LLVM IR translation (cicc's equivalent of Clang CodeGen)
      │            Type translation (fixed-point iteration, address space mapping)
      │            Expression/statement/function codegen (recursive AST walk)
      │            CUDA semantic lowering (threadIdx→intrinsics, printf→vprintf, etc.)
      │            Kernel metadata emission (nvvm.annotations)
      │            Two copies: Path A (0x90xxxx) and Path B (0x126xxxx)
      │
      ├─ LNK:     Module linking + libdevice (455KB embedded bitcode)
      │            Triple validation (must be nvptx64-)
      │            IR version check (nvvmir.version metadata)
      │
      ├─ OPT:     Two-phase compilation (Phase I: whole-module, Phase II: per-function)
      │            ~150 pass insertions via sub_12E54A0
      │            Three language paths: "ptx" / "mid" / default
      │            35 NVIDIA custom passes interleaved with standard LLVM
      │            Optional: concurrent per-function compilation (thread pool + jobserver)
      │
      ├─ OPTIXIR:  OptiX IR generation (optional, --emit-optix-ir)
      │
      └─ LLC:     NVPTX backend code generation
                   SelectionDAG lowering (2.3 MB NVPTXTargetLowering)
                   19 MMA shapes × 11 data types for tensor core codegen
                   9 PTX register classes
                   StructurizeCFG (mandatory for PTX structured control flow)
                   → .ptx output

Subsystem Address Map

Reading This Wiki

The wiki is organized around the compilation pipeline. Every page is written at reimplementation-grade depth for an audience of senior C++ developers with LLVM backend experience.

Section Index

• Pipeline Overview — End-to-end compilation flow diagram with links to every stage.
• Entry Point & CLI — CLI parsing, dual-path dispatch, architecture detection.
• EDG 6.6 Frontend — CUDA C++ to transformed C source-to-source translation.
• NVVM IR Generation — EDG IL tree to LLVM Module: types, expressions, statements, functions.
• LLVM Optimizer — Two-phase compilation, pipeline assembly, NVVMPassOptions.
• Code Generation — SelectionDAG, ISel, register allocation, scheduling.
• PTX Emission — AsmPrinter, directive emission, PTX body output.
• NVIDIA Custom Passes — 35 proprietary passes not in upstream LLVM.
• LLVM Pass Pipeline & Ordering — Complete pass registration, execution order per O-level, tier system.
• NVVM Builtins — 770-entry builtin table: hash structure, ID inventory, category breakdown.
• GPU Targets — SM feature gates, architecture detection, sm_75 through sm_121f.
• Data Structures — IR node layout, pattern database, DAG node, symbol table, NVVM container.
• Infrastructure — Alias analysis, MemorySSA, AsmPrinter, debug verification, NVPTX target.
• LTO & Module Optimization — Cross-TU inlining, devirtualization, GlobalOpt, ThinLTO import.
• Configuration — Three knob systems: ~1,689 cl::opt flags, 222 NVVMPassOptions slots, ~70 codegen knobs.
• Reference — Address spaces, register classes, NVPTX opcodes, GPU execution model.
• Function Map — Address-to-identity lookup for ~350 key functions with confidence levels.
• Binary Layout — Subsystem address map at pass granularity.
• Methodology — How this analysis was performed and how to assess confidence.

Reading Path 1: End-to-End Pipeline Understanding

Goal: understand how CUDA source becomes PTX, what each stage does, and how control flows between subsystems.

Read in this order:

1. Pipeline Overview — The complete flow diagram. Establishes the 10 stages and their address ranges. Read this first to build the mental model that all other pages assume.
2. Entry Point & CLI — How cicc is invoked, the 1,689-flag CLI, dual-path dispatch (Path A LibNVVM vs. Path B standalone), and the sub_8F9C90 real-main function.
3. nvcc-to-cicc Interface — The flag translation layer between nvcc and cicc. The 40+ flag mappings and 3-column architecture fan-out. Necessary context for understanding why certain flags exist.
4. EDG 6.6 Frontend — The commercial C++ frontend. How CUDA syntax is lowered to C, the 737 configuration #defines, and the .int.c / .device.c / .stub.c output split.
5. NVVM IR Generation — The EDG-to-LLVM bridge. Then follow the four sub-pages: Type Translation → Expressions → Statements → Functions.
6. Libdevice Linking — The embedded 455KB bitcode library with 352 __nv_* math functions. Triple validation, version checking.
7. LLVM Optimizer — The two-phase compilation model, the 49.8KB pipeline assembler (sub_12E54A0), pass ordering, and the NVVMPassOptions knob system. This is the longest and densest stage.
8. Pipeline & Pass Ordering — The exact pass execution order at each O-level, the tier system, and the 526 registered passes.
9. Code Generation — SelectionDAG lowering, instruction selection, register allocation, instruction scheduling. Hub page with links to deep dives.
10. PTX Emission — AsmPrinter, directive headers, PTX body output, metadata emission.

Optional extensions after the core path:

• OptiX IR Generation — The alternative output mode for ray tracing workloads.
• Debug Info Pipeline — How -g debug metadata survives the optimizer.
• LTO & Module Optimization — Cross-module optimization when compiling multiple translation units.
• Concurrent Compilation — The Phase II thread pool and GNU Jobserver integration.
• GPU Execution Model — Background on warps, divergence, shared memory, and address spaces if you are new to GPU architecture.

Reading Path 2: Reimplementing a Specific Pass

Goal: reproduce the exact behavior of one NVIDIA custom pass or understand an LLVM pass modification deeply enough to write a compatible replacement.

For an NVIDIA custom pass (e.g., MemorySpaceOpt, Rematerialization, BranchDist):

1. NVIDIA Custom Passes — Overview — Locate the pass in the inventory table. Note its category (module/function/loop/machine), its pipeline position, and its controlling knobs.
2. The pass's dedicated page (e.g., MemorySpaceOpt, Rematerialization, Branch Distribution). Every dedicated page contains the function address, decompiled algorithm, data flow description, controlling knobs, and diagnostic strings.
3. NVVMPassOptions — The 222-slot struct that controls per-pass enable/disable toggles and parametric thresholds. Find which slots your target pass reads.
4. Pipeline & Pass Ordering — Determine exactly where the pass runs in the pipeline. Identify what analyses it depends on (must run before it) and what passes consume its results (run after it).
5. Optimization Levels — Determine at which O-levels the pass is enabled, disabled, or parameterized differently.
6. Function Map — Cross-reference the pass's internal function addresses with the master function map for confidence levels.

For a modified LLVM pass (e.g., InstCombine, GVN, DSE, LICM, LoopVectorize):

1. The pass's dedicated page (e.g., InstCombine, GVN, DSE, LICM). These pages document NVIDIA's modifications relative to upstream LLVM 20.0.0.
2. Alias Analysis & NVVM AA — The custom alias analysis chain. Nearly every optimization pass depends on AA, and NVIDIA's GPU-aware AA behaves differently from upstream (address-space-aware NoAlias for disjoint spaces, __restrict__ propagation).
3. MemorySSA — The memory dependence representation used by DSE, LICM, and other memory-sensitive passes.

For a machine-level pass (e.g., Block Remat, MRPA, Machine Mem2Reg):

1. Machine-Level Passes — The complete machine pass pipeline with per-pass algorithm descriptions.
2. Register Allocation — The greedy RA algorithm with NVIDIA's occupancy-driven spill heuristics.
3. Register Classes — The 9 PTX register classes and their constraints.
4. NVPTX Machine Opcodes — The MachineInstr opcode reference.

Supporting references for any pass reimplementation:

• IR Node Layout — The internal IR data structures that passes operate on.
• Address Spaces — GPU address space semantics that many passes must respect.
• NVPTX Target Infrastructure — TargetMachine, TTI hooks, and target feature queries.
• Diagnostics — The three diagnostic systems (EDG, LLVM remarks, profuse framework) for reproducing pass-level reporting.

Reading Path 3: Debugging Correctness

Goal: diagnose a miscompilation, a crash, or incorrect PTX output by tracing the problem to a specific pass or pipeline stage.

Start with instrumentation and observability:

1. Diagnostics & Optimization Remarks — The three independent diagnostic layers: EDG frontend errors, LLVM optimization remarks (-opt-bisect-limit, -Rpass=, -Rpass-missed=), and NVIDIA's profuse framework (profuseinline, profusegvn). This page tells you how to make cicc talk about what it is doing.
2. Debug Info Verification — The three verification modes (verify-each, debugify-each, and JSON delta reporting). Use verify-each to detect the first pass that corrupts debug metadata.
3. CLI Flags — Locate the flags for dumping IR at specific pipeline points: --print-after-all, --print-before-all, --filter-print-funcs=, --opt-bisect-limit=. Also the --passes= interface for running individual passes in isolation.
4. Optimization Levels — Compare the pass pipeline at different O-levels. If a bug appears at -O2 but not -O1, the diff between their pipelines identifies the suspect passes.

Then isolate the pipeline stage:

5. Pipeline Overview — Determine which stage produces the incorrect output. The pipeline is linear: EDG → IR Generation → Libdevice Linking → Optimizer → Codegen → Emission. The stage boundary where output first goes wrong narrows the search.
6. NVVM IR Verifier — The 230KB three-layer verifier (module + function + intrinsic). It validates triples, address spaces, atomic restrictions, pointer cast rules, and architecture-gated intrinsic availability. A verification failure after a specific pass is a strong signal.
7. Bitcode I/O — If the problem is in bitcode reading/writing (corrupted input, version mismatch), this page documents the reader at sub_9F2A40 and the writer.

Then investigate the suspect pass:

8. NVIDIA Custom Passes or the relevant LLVM pass page — Read the algorithm description for the suspect pass. Look for documented edge cases, known limitations, and diagnostic strings that would appear in verbose output.
9. NVVMPassOptions — Check whether the suspect pass has enable/disable knobs or threshold parameters that could be adjusted to confirm or rule it out.
10. Environment Variables — Some passes are gated by environment variables (including obfuscated ones). Check whether any are influencing behavior.

For correctness issues specific to GPU semantics:

• Address Spaces — Incorrect address space resolution is a common source of silent miscompilation. Global vs. shared vs. local aliasing rules differ from CPU memory models.
• MemorySpaceOpt — This pass resolves generic pointers to specific address spaces. If it infers the wrong space, downstream code will access the wrong memory.
• Alias Analysis — If the alias analysis returns NoAlias for pointers that do alias, DSE/LICM/GVN will misoptimize. The process-restrict propagation is a known source of aggressive alias assumptions.
• StructurizeCFG — PTX requires structured control flow. If structurization produces incorrect flow blocks, the kernel will execute the wrong path.
• Dead Barrier Elimination and Dead Synchronization Elimination — Incorrect elimination of barriers or synchronization can cause race conditions that only manifest under specific warp configurations.

Reading Path 4: Tuning Performance

Goal: understand what cicc does at each optimization level, which passes are the performance-critical ones, and what knobs control their aggressiveness.

Start with the tuning infrastructure:

1. Optimization Levels — The four standard levels (O0--O3) and three fast-compile tiers (Ofcmin/Ofcmid/Ofcmax). This page shows the exact pass pipeline diff between levels, including which passes are added, removed, or reparameterized at each step.
2. NVVMPassOptions — The 222-slot per-pass configuration system. This is the primary tuning mechanism. The page documents every slot's type (boolean/integer/string), its default value, and which pass reads it.
3. CLI Flags — The flag-to-pipeline routing tables. Locate flags that control pass thresholds (--inline-threshold=, --unroll-count=, etc.) and pass enable/disable toggles.
4. LLVM Knobs — The ~1,689 cl::opt flags with their defaults, types, and controlling constructors.
5. Environment Variables — Runtime environment overrides, including the obfuscated variables.

Then study the high-impact optimization passes:

6. LLVM Optimizer — Understand the two-phase model. Phase I (whole-module) determines inlining decisions, inter-procedural memory space propagation, and global optimization. Phase II (per-function, potentially concurrent) does register-pressure-driven rematerialization and instruction scheduling. Tuning decisions in Phase I cascade into Phase II.
7. Inliner Cost Model — Inlining is typically the single highest-impact optimization decision. This page documents the cost model thresholds, the caller/callee size heuristics, and NVIDIA's kernel-specific adjustments.
8. LoopVectorize & VPlan — Loop vectorization for GPU SIMT. The VPlan infrastructure, cost model, and the NVIDIA TTI hooks that influence vectorization width decisions.
9. Loop Unrolling — Unrolling thresholds, the NVIDIA-specific unroll heuristics, and the interaction with register pressure.
10. Rematerialization — NVIDIA's IR-level rematerialization pass (67KB). Trades recomputation for register pressure reduction, which directly affects occupancy on GPU.
11. Register Allocation — The greedy RA with occupancy-driven spill heuristics. Register count directly determines maximum occupancy.
12. Instruction Scheduling — The scheduler subsystems and their interaction with hardware latency models.

For tensor core workloads specifically:

• Tensor / MMA Codegen — 19 MMA shapes across 11 data types. The instruction selection patterns, register allocation constraints, and WGMMA code generation for Hopper and Blackwell.
• Tensor / MMA Builtins — The builtin-to-intrinsic lowering for wmma, mma, and wgmma operations.
• SM 90 — Hopper — Hopper-specific features: TMA, WGMMA, asynchronous barriers, cluster launch.
• SM 100 — Blackwell — Blackwell-specific features: new MMA shapes, FP4/FP6 support, sparsity.

For understanding performance at the target level:

• GPU Targets — The SM feature gate matrix. Which features are enabled at each architecture level, and how architecture detection routes to different codegen paths.
• NVPTX Target Infrastructure — The TTI hooks that passes query for target-specific costs (memory latency, instruction throughput, register file size).
• Concurrent Compilation — If compile time itself is the bottleneck, understand the Phase II thread pool and GNU Jobserver integration to maximize parallelism.

Function Map

Address-to-identity lookup table. Confidence: VERY HIGH = string evidence, HIGH = strong structural evidence, MEDIUM = inferred from context/callgraph.

Top Functions by Size

Pipeline Functions

EDG 6.6 Frontend

Core Orchestration

AST-to-Source Printer Cluster

Parser & Declaration Processing

Type System

IL Tree Infrastructure

Constexpr Evaluator

Preprocessor

Template Engine

Semantic Analysis

CUDA-Specific Frontend

Name Mangling (Itanium ABI)

Diagnostics & Support

Class Layout & Vtable

LLVM cl::opt Registration Infrastructure

Key Constructors (cl::opt registration)

NVIDIA Custom Pass Implementations

MMA / Tensor Core Emission

PTX Emission

Hash Infrastructure

NVVM Builtin Infrastructure

Register Allocation

jemalloc (Statically Linked, v5.3.x)

Optimizer Pipeline Assembly

Functions discovered during wiki writing (W101--W241). These assemble the LLVM optimization pipeline from NVVMPassOptions slots.

Pipeline Builders

Pass Factories (Pipeline Insertion Order)

Each factory creates a pass instance; referenced from sub_12E54A0, sub_12DE330, and sub_12DE8F0.

NVPTX Backend (SelectionDAG & ISel)

NVVM Verifier Subsystem

LTO Subsystem

LLVM IR Utility Functions

Common LLVM IR manipulation functions referenced across many passes.

Machine-Level Infrastructure

Binary Layout

This page is a visual guide to navigating the cicc v13.0 binary in IDA Pro. It covers the ELF structure, section layout, subsystem address ranges, embedded data payloads, and the statically linked jemalloc allocator. If you are opening this binary for the first time, start here to orient yourself before diving into individual subsystems.

ELF Overview

CICC is a statically linked, stripped x86-64 ELF binary. There are no dynamic symbol tables, no .dynsym, no DWARF debug info, and no export table. Every function name was removed at build time. IDA Pro recovers 80,562 functions; Hex-Rays successfully decompiles 80,281 of them (99.65%).

Because the binary is statically linked, libc, libpthread, and libm are all embedded. This inflates the raw function count but also means every call target resolves to a concrete address within the binary itself -- there are no external dependencies at runtime beyond the kernel syscall interface.

Address Space Map

The binary's .text section spans roughly 0x400000 to 0x3C00000. Within that 56 MB range, subsystems occupy contiguous, non-overlapping regions. The map below is the primary orientation tool for IDA Pro navigation.

0x400000 ┌─────────────────────────────────────────┐
         │  CRT startup + libc stubs               │  ~52 KB
0x40D000 ├─────────────────────────────────────────┤
         │  jemalloc stats / vsnprintf              │  ~80 KB
0x420000 ├─────────────────────────────────────────┤
         │  (gap: misc libc, math, string ops)      │  ~64 KB
0x430000 ├─────────────────────────────────────────┤
         │  Global constructors (cl::opt reg)        │  ~1.6 MB
         │  ~1,689 LLVM command-line option objects  │
0x5D0000 ├─────────────────────────────────────────┤
         │  EDG 6.6 C++ Frontend                    │  3.2 MB
         │  Parser, constexpr evaluator, IL walker   │
0x8F0000 ├─────────────────────────────────────────┤
         │  CLI / Real Main / NVVM Bridge            │  520 KB
         │  sub_8F9C90 (real main), dual-path dispatch│
0x960000 ├─────────────────────────────────────────┤
         │  Architecture detection, NVVM options     │  576 KB
0x9F0000 ├─────────────────────────────────────────┤
         │  Bitcode reader (parseFunctionBody)       │  ~1 MB
0xAF0000 ├─────────────────────────────────────────┤
         │  X86 AutoUpgrade (legacy, 457KB fn)       │  ~1 MB
0xBF0000 ├─────────────────────────────────────────┤
         │  LLVM IR Verifier                        │  500 KB
0xC00000 ├─────────────────────────────────────────┤
         │  LLVM Support / ADT library              │  ~3.2 MB
         │  (see detailed sub-map below)             │
0x12D0000├─────────────────────────────────────────┤
         │  PassManager / NVVM bridge                │  4.2 MB
         │  Pipeline assembly (sub_12E54A0)          │
0x12FC000├ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ┤
         │  jemalloc core (~400 functions)           │  ~256 KB
0x1700000├─────────────────────────────────────────┤
         │  Backend / machine passes                 │  8 MB
         │  RegAlloc, Block Remat, Mem2Reg           │
0x1F00000├─────────────────────────────────────────┤
         │  SelectionDAG                            │  2 MB
         │  LegalizeTypes (348KB), LegalizeOp        │
0x2100000├─────────────────────────────────────────┤
         │  NVPTX PTX emission                      │  1 MB
0x2340000├─────────────────────────────────────────┤
         │  New PM / pass registration               │  768 KB
         │  2,816-line registrar at sub_2342890      │
0x2A00000├─────────────────────────────────────────┤
         │  Loop passes                             │  4 MB
         │  LoopVectorize, SLP, Unroll               │
0x3000000├─────────────────────────────────────────┤
         │  NVPTX ISel + lowering                    │  7 MB
         │  343KB intrinsic switch (sub_33B0210)     │
0x3700000├─────────────────────────────────────────┤
         │  Machine-level passes (tail)              │  ~3 MB
         │  BlockPlacement, Outliner, StructurizeCFG │
0x3A00000├─────────────────────────────────────────┤
         │  (trailing code, CRT finalization)        │
         └─────────────────────────────────────────┘

DATA SECTIONS:
0x3EA0080   Embedded libdevice bitcode (Path A)    456 KB
0x420FD80   Embedded libdevice bitcode (Path B)    456 KB
0x4F00000+  Global BSS (cl::opt storage, hash tables, state)

Detailed Subsystem Map at Pass Granularity

The coarse map above partitions the binary into ~18 zones. The following map refines every zone to individual-pass resolution, giving the factory address of each identified pass or subsystem entry point. Addresses prefixed with sub_ are IDA function names. Sizes in parentheses are decompiled C output; actual machine code is typically 2-3x smaller.

Zone 1: CRT, libc, jemalloc stats (0x400000 - 0x42FFFF)

0x400000   _start / CRT entry (ELF entry point)
0x40D5CA   sub_40D5CA   vsnprintf (jemalloc stats formatting)
0x420000   libc math/string helpers (memcpy, memset, strlen, etc.)

No LLVM or NVIDIA code lives here. Pure runtime support.

Zone 2: Global constructors (0x430000 - 0x5CFFFF)

~1,689 cl::opt registration constructors execute before main(). Each registers a command-line option string, description, default value, and storage pointer into the global option registry. The .init_array section holds function pointers to these constructors.

Zone 3: EDG 6.6 C++ Frontend (0x5D0000 - 0x8EFFFF)

The complete Edison Design Group C++ frontend, version 6.6. Contains the lexer, parser, constexpr evaluator, template instantiator, overload resolver, IL walker/copier, diagnostic engine, SARIF output, and CUDA-specific extensions (kernel launch grammar, __shared__/__device__ memory space parsing, atomic builtin stubs).

Zone 4: CLI / Real Main / Dual-Path Entry (0x8F0000 - 0x9EFFFF)

Zone 5: Bitcode Reader / X86 AutoUpgrade / Verifier (0x9F0000 - 0xBFFFFF)

Zone 6: LLVM Support Library (0xC00000 - 0xCAFFFF)

1,653 functions. Pure LLVM infrastructure -- no NVIDIA-specific modifications except a single !Flat address space annotation in the sample profile reader at sub_C29E70.

Zone 7: NVVM Container, SCEV, DWARF, MC Layer (0xCB0000 - 0x10CFFFF)

This 4 MB zone contains LLVM mid-level infrastructure and the NVVM container format.

Zone 8: InstCombine Mega-Region (0x10D0000 - 0x122FFFF)

The single largest contiguous pass in the binary. NVIDIA's modified InstCombine spans 1.4 MB of code with three NVIDIA-custom opcodes (0x254D, 0x2551, 0x255F) for proprietary intrinsic folding.

Zone 9: NVVM Bridge / Builtin System / IR Codegen (0x1230000 - 0x12CFFFF)

This zone is the core NVIDIA bridge between the EDG frontend AST and the LLVM IR optimizer.

Zone 10: Pipeline Builder / Pass Options (0x12D0000 - 0x12FFFFF)

The pipeline assembler constructs the complete LLVM pass pipeline, inserting passes by calling factory functions whose addresses scatter across the entire binary.

Zone 11: IR Infrastructure / PassManager (0x1300000 - 0x16FFFFF)

Dense LLVM infrastructure: IR types, constants, instructions, metadata, use-lists, PassManager execution engine, IR linker, bitcode reader, regex, and DataLayout.

Zone 12: InstCombine (NewPM) + Sanitizers + PGO (0x1700000 - 0x17FFFFF)

946 functions. Dominated by the New Pass Manager version of InstCombine (~600 functions, ~3.5 MB decompiled), with sanitizer instrumentation (MSan, TSan, coverage) and PGO/GCov infrastructure.

Zone 13: GVN + Scalar Passes + NVIDIA Custom IR Passes (0x1800000 - 0x1CFFFFF)

This 5 MB zone contains the bulk of LLVM's scalar optimization passes and all of NVIDIA's custom IR-level passes.

GVN family (0x1900000 - 0x193FFFF):

Standard scalar passes (0x1830000 - 0x1AFFFFF):

Loop unrolling + switch lowering (0x1B00000 - 0x1B7FFFF):

Loop/SLP vectorizer (0x1B80000 - 0x1BFFFFF):

NVVM module validation + configuration (0x1C00000 - 0x1C3FFFF):

NVIDIA custom IR passes (0x1C40000 - 0x1CFFFFF):

This 1 MB block contains the majority of NVIDIA's proprietary IR-level optimization passes. Every pass listed here has no upstream LLVM equivalent.

Zone 14: SelectionDAG ISel / CodeGenPrepare / Backend (0x1D00000 - 0x1EFFFFF)

Zone 15: Backend CodeGen Infrastructure (0x1F00000 - 0x20FFFFF)

Zone 16: NVPTX Target Backend (0x2100000 - 0x21FFFFF)

Zone 17: New PM Pass Registration (0x2340000 - 0x23FFFFF)

Zone 18: IPO / Attributor / OpenMP Optimization (0x2400000 - 0x29FFFFF)

Zone 19: Loop Transforms (0x2A00000 - 0x2CFFFFF)

Note: sub_2C84BA0 is a second copy of the dead synchronization elimination pass located outside the main NVIDIA custom pass zone. This is the 94KB variant analyzed in depth (p2b.6-01), with the four-category fixed-point R/W dataflow algorithm and red-black tree maps.

Zone 20: Codegen Target Options / SelectionDAG Lowering (0x2D00000 - 0x2FFFFFF)

5,217 functions. Contains LLVM TargetMachine option registration and the core SelectionDAG infrastructure used by the NVPTX backend.

Zone 21: NVPTX ISel + SelectionDAG Lowering (0x3000000 - 0x36FFFFF)

7 MB. The NVPTX instruction selection and target-specific DAG lowering.

Zone 22: NVPTX Instruction Selector / Machine Tail (0x3700000 - 0x3BFFFFF)

Pass Factory Address Summary

The pipeline assembler (sub_12E54A0) calls pass factory functions to construct the pipeline. Each factory address below is called directly from the pipeline builder and uniquely identifies a pass in the binary.

Type: M = ModulePass, F = FunctionPass, L = LoopPass.

Embedded Data Payloads

Libdevice Bitcode

Two identical copies of NVIDIA's libdevice are embedded directly in the .rodata section as raw LLVM bitcode. Each copy is approximately 456 KB and contains around 400 math intrinsic implementations (__nv_sinf, __nv_expf, __nv_sqrtf, etc.). The duplication supports the dual-path architecture: Path A (LibNVVM API mode) references one copy at 0x3EA0080; Path B (standalone mode) references the other at 0x420FD80. The bitcode is linked into the user's module during the LNK phase via the bitcode linker at sub_12C06E0.

String Tables

IDA Pro extracts 188,141 strings from the binary. These fall into several categories:

String cross-referencing is the single most productive technique for identifying functions in a stripped binary. The LLVM pass registration pattern is especially reliable: a string like "nvvm-memspace-opt" appears exactly once, in the constructor of that pass, which IDA locates via xref.

NVVM Container Format

The binary includes a proprietary container format for wrapping LLVM bitcode with compilation metadata. The container uses a 24-byte binary header with magic 0x7F4E5C7D, followed by delta-encoded tag/value pairs (only fields that differ from defaults are serialized). There are 144 distinct tag IDs spanning core options (tags 1-39), compression metadata (tag 99), extended target options (tags 101-173), blob data (tags 201-218), and structured hardware descriptors (tags 401-402 for TMA/TCGen05 configurations). Serialization and deserialization are handled by sub_CDD2D0 and sub_CD1D80 respectively.

jemalloc Integration

NVIDIA statically links jemalloc 5.3.x as the process-wide memory allocator. The jemalloc functions cluster around 0x12FC000 (approximately 400 functions). The configuration initialization function sub_12FCDB0 (129 KB, one of the largest functions in the binary) parses 199 configuration strings from the MALLOC_CONF environment variable.

Key jemalloc entry points visible in the binary:

The jemalloc integration is significant for reverse engineering because it means malloc/free calls throughout the binary resolve to jemalloc's arena-based allocator rather than glibc's ptmalloc2. When tracing memory allocation patterns in IDA, look for calls into the 0x12FC000 range.

Global Constructors

The region from 0x430000 to 0x5CFFFF (~1.6 MB) is dominated by global constructors that execute before main(). The primary purpose of these constructors is LLVM cl::opt registration: approximately 1,689 command-line option objects are initialized, each registering a string name, description, default value, and storage location into LLVM's global option registry.

The .init_array section contains function pointers to these constructors. They execute in linker-determined order and populate a global hash table that sub_8F9C90 (the real main) later queries during CLI parsing. In IDA Pro, navigating to any cl::opt constructor reveals the option name string and its associated global variable, which is invaluable for understanding what flag controls what behavior.

Additional global constructors handle:

• LLVM pass registration (RegisterPass<T> and PassInfo objects)
• LLVM target initialization (NVPTX target machine factory)
• jemalloc allocator bootstrapping
• EDG frontend static initialization tables

Dual-Path Code Duplication

A distinctive structural feature of the binary is the presence of two near-complete copies of the NVVM bridge and backend entry points. Path A (LibNVVM API mode) lives around 0x90xxxx; Path B (standalone/nvcc mode) lives around 0x126xxxx. Each path has its own:

In IDA, if you have identified a function in one path, search for a structurally similar function at the corresponding offset in the other path. The code is not byte-identical -- Path B is generally slightly larger due to additional standalone-mode logic -- but the control flow graphs are nearly congruent.

IDA Pro Navigation Tips

When opening cicc in IDA Pro for the first time, the auto-analysis will take several minutes due to the 60 MB size. The following workflow accelerates orientation:

1. Start with strings. Open the Strings window (Shift+F12), filter for known LLVM pass names ("instcombine", "gvn", "nvvm-"). Each xref leads directly to a pass constructor or registration site.

2. Use the address map above. If you are looking at an address in the 0xC00000-0x12CFFFF range, you are in LLVM optimization passes. The 0x3000000-0x36FFFFF range is NVPTX instruction selection. The 0x5D0000-0x8EFFFF range is EDG. Context narrows the search space immediately.

3. Watch for vtable patterns. LLVM passes are C++ classes with virtual methods. IDA's vtable reconstruction reveals inheritance hierarchies. Every FunctionPass, ModulePass, and LoopPass subclass has a vtable with runOnFunction/runOnModule at a consistent slot offset.

4. Anchor on mega-functions. The largest functions are the easiest to locate and serve as landmarks: sub_A939D0 (457 KB, X86 AutoUpgrade), sub_10EE7A0 (396 KB, InstCombine), sub_20019C0 (341 KB, LegalizeTypes). These anchors partition the address space.

5. Follow the pipeline. Entry at sub_8F9C90 calls into EDG at sub_5D2A80, pipeline assembly at sub_12E54A0, and PTX emission starting at 0x2100000. Tracing callgraph edges from these known entry points maps out the entire compilation flow.

6. Mark jemalloc early. Identifying and labeling the jemalloc cluster at 0x12FC000 prevents wasted time reverse-engineering well-known allocator internals. The 199-string malloc_conf_init function is an unmistakable fingerprint.

7. Locate NVIDIA passes via factory addresses. The Pass Factory Address Summary table above maps every pipeline-inserted pass to its constructor address. In IDA, setting a breakpoint at sub_12DE0B0 (AddPass) and logging the second argument reveals the exact pass insertion order at runtime.

Master Address-Range Map

The definitive quick-reference for "what lives at address X?" Every major address range in the cicc v13.0 binary, sorted by start address, consolidated from all subsystem pages in this wiki.

.text Section (0x400000 - 0x3BFFFFF)

.rodata / .data Sections (0x3C00000+)

Usage

Given an IDA address, find the row whose Start <= address < End. The Subsystem column tells you which component of cicc you are looking at. For pass-level detail within a zone, jump to the corresponding Zone section above.

Cross-References

• Pipeline Overview -- compilation flow from entry to PTX emission
• LLVM Pipeline -- 526-pass registration table and tier execution order
• Optimizer -- two-phase model, AddPass mechanism, tier system
• Pass Inventory -- complete pass catalog with dedicated deep-dive pages
• NVVMPassOptions -- 222-slot pass configuration system
• Function Map -- address-to-identity lookup table
• CLI Flags -- flag-to-pipeline routing

Methodology

This page documents how the reverse engineering of cicc v13.0 was performed. It serves as both a transparency record -- so readers can assess the confidence of any claim in this wiki -- and as a practical guide for anyone who wants to reproduce or extend the analysis.

Scope and Scale

CICC is a 60 MB stripped x86-64 ELF binary with no debug symbols, no export table, and no DWARF information. The scale of the analysis:

The 281 functions that Hex-Rays could not decompile are predominantly very small thunks, computed-jump trampolines, or hand-written assembly stubs in the CRT startup and jemalloc fast paths. None are in critical compiler logic.

Toolchain

All analysis was performed with IDA Pro 8.x and the Hex-Rays x86-64 decompiler. No dynamic analysis (debugging, tracing, instrumentation) was used -- the entire effort is static analysis of the binary at rest. Supplementary tools:

No runtime instrumentation, no strace/ltrace, no gdb breakpoints. Every finding derives from static analysis of the binary's code and data sections.

Function Identification Strategies

Identifying functions in a stripped binary of this size requires multiple complementary strategies. They are listed below in order of reliability.

String Cross-References (Highest Confidence)

LLVM is a string-rich codebase. Error messages, pass names, option descriptions, and assertion text are compiled into the binary. A string like "Running pass 'NVVMMemorySpaceOpt'" appears at exactly one address in .rodata, and IDA's xref from that string leads directly to the function that prints it. This is the most reliable identification technique and produces VERY HIGH confidence identifications.

Specific high-value string patterns:

• LLVM pass registration: "instcombine", "gvn", "nvvm-memspace-opt" -- each appears in exactly one RegisterPass constructor or PassInfo initializer.
• cl::opt names: "-nvvm-enable-remat", "-nvvm-branch-dist-threshold" -- each names a global variable and its registration constructor.
• Error messages with context: "parseFunctionBody: ..." (174 unique error strings in the bitcode reader), "visitCallInst: ..." (298 verification messages in the verifier).
• Timer names: "CUDA C++ Front-End", "LibNVVM", "Optimizer" -- appear in timer-creation calls that bracket pipeline stages.
• EDG error templates: "expected a %s", "declaration not allowed here" -- 2,500+ diagnostic strings anchoring the frontend parser.

LLVM Pass Registration Patterns (Very High Confidence)

Every LLVM pass follows a predictable structural pattern. A pass class has a vtable with virtual methods at fixed offsets (runOnFunction at slot N, getAnalysisUsage at slot M). The pass registers itself via a global constructor that stores a PassInfo object containing the pass name string, the pass ID address, and a factory function pointer. By enumerating all .init_array entries that write a PassInfo-shaped structure, all ~437 passes were catalogued systematically.

The New Pass Manager (at sub_2342890, a 2,816-line registrar function) contains a massive string-to-pass-factory dispatch table with ~268 pass name entries. Decompiling this single function yields the name-to-address mapping for every New PM pass in the binary.

Vtable Analysis (High Confidence)

LLVM's class hierarchy is deep and regular. Pass -> FunctionPass -> LoopPass, Pass -> ModulePass, etc. Each level adds virtual methods at predictable vtable slots. By reconstructing vtable layouts (finding pointers to __cxa_pure_virtual for abstract methods, then tracing concrete overrides), the class hierarchy was reconstructed without debug symbols.

For the NVPTX backend specifically, vtable analysis identified NVPTXTargetLowering (2.3 MB of lowering logic), NVPTXInstrInfo, NVPTXRegisterInfo, and NVPTXFrameLowering as distinct classes with their own method tables.

Callgraph Propagation (High Confidence)

Once a function is identified with high confidence, its callees and callers gain contextual identity. If sub_12E54A0 is the pipeline assembly function (confirmed by string refs to pass names it registers), then the functions it calls to create individual passes are the pass factory functions. This propagation is transitive: identifying a factory function identifies its return type's vtable, which identifies the pass's runOnFunction method.

The pipeline orchestrator at sub_12C35D0 (41 KB) is a particularly productive anchor: it calls into the LNK, OPT, OPTIXIR, and LLC stages in sequence, and each stage's entry point was identified by following its callgraph edges.

Size and Structural Fingerprinting (Medium Confidence)

Some functions are identifiable by their size and structural characteristics alone. LLVM's InstCombine::visitCallInst is famously enormous (396 KB in this binary) because it handles every LLVM intrinsic. SelectionDAG::LegalizeTypes (348 KB) contains a switch with 967 case labels. These mega-functions have no structural equivalents and can be identified by size alone with reasonable confidence.

Similarly, the EDG frontend's constexpr evaluator (sub_786210, 317 KB) is identifiable by its 124 case labels corresponding to C++ operator opcodes -- a characteristic that matches the known EDG evaluator design.

Known Library Fingerprinting (Medium Confidence)

jemalloc was identified by its 199 configuration string names ("background_thread", "dirty_decay_ms", "narenas", etc.), which are unique to jemalloc's malloc_conf_init function. Once the allocator library was identified, its ~400 functions were bulk-labeled, removing them from the analysis scope.

The X86 AutoUpgrade function (sub_A939D0, 457 KB) is an LLVM artifact -- leftover x86 intrinsic renaming code that ships in every LLVM-based binary regardless of target. It was identified by its intrinsic name strings ("llvm.x86.sse2.*", "llvm.x86.avx.*") and excluded from NVPTX-specific analysis.

Confidence Levels

Every function identification in this wiki carries one of four confidence levels:

Approximately 60% of identified functions are VERY HIGH or KNOWN confidence. The remaining 40% are HIGH or MEDIUM, concentrated in areas with fewer string anchors (machine-level passes, register allocation internals, EDG IL tree walkers).

Analysis Pipeline and Scripts

The manual IDA Pro work was augmented by a systematic scripted pipeline that processed the exported IDA databases into structured evidence. The pipeline operates in two phases: L0 (foundation) builds indices and classifies all 80,562 functions automatically, and L1 (module analysis) organizes functions into per-module directories with metadata for human review.

All scripts live in cicc/scripts/. The pipeline requires four JSON databases exported from IDA: cicc_functions.json (80,562 function records), cicc_strings.json (188,141 string records), cicc_xrefs.json (cross-reference records), and cicc_callgraph.json (call edge records). These exports are stored in cicc/databases/.

L0 Foundation Pipeline

The L0 pipeline runs as a single sequential batch via scripts/run_foundation_analysis.sh. Each step depends on the output of the previous step.

Step 0: Extract Wiki Knowledge (foundation/00_extract_wiki_knowledge.py)

Scans all existing wiki markdown files for hex addresses (regex \b0x[0-9a-fA-F]{6,}\b) and builds a ground-truth mapping of address-to-module from prior manual analysis. This seed data provides the highest-confidence module assignments (100% confidence) used to bootstrap the automated classifier.

Output: foundation/taxonomy/modules/wiki_known_functions.json, wiki_module_addresses.json.

Step 1: Build Fast Lookup Indices (foundation/01_build_indices.py)

Loads the three IDA JSON databases (functions, strings, xrefs) and builds four pickle-serialized indices for O(1) lookup in subsequent steps:

• addr_to_func.pkl -- address to function metadata (name, size, instruction count, library/thunk flags).
• string_to_xrefs.pkl -- string address to string value and xref list.
• func_to_callers.pkl -- function name to list of caller names.
• func_to_callees.pkl -- function name to list of callee names.

Output: foundation/indices/.

Step 2: Classify Strings (foundation/02_classify_strings.py)

Applies four regex-based pattern sets to all 188,141 strings, classifying each into one or more semantic categories:

• Error messages: strings matching error, failed, invalid, unsupported, expected, etc.
• Optimization passes: strings matching pass, optimize, transform, inline, unroll, gvn, licm, etc.
• Architecture features: strings matching sm_\d+, tensor, warp, FP4, blackwell, hopper, etc.
• Debug messages: strings matching debug, trace, dump, verbose.

Each classified string retains its address and xref list, so the classifier output doubles as a "which functions reference optimization-related strings" index.

Output: foundation/taxonomy/strings/error_messages.json, optimization_passes.json, architecture_features.json, debug_messages.json, extracted_pass_names.json.

Step 3: Build Module Taxonomy (foundation/03_build_module_taxonomy.py)

The core classification engine. Assigns each of the 80,562 functions to one of eight compiler subsystem modules (or unknown) using four strategies applied in decreasing confidence order:

1. Wiki ground truth (100% confidence) -- addresses found in wiki pages in Step 0.
2. String content analysis (80% confidence) -- functions whose string xrefs match module-specific keyword patterns (e.g., a function referencing "tensor", "mma", or "tcgen" strings is classified as tensor_core_codegen).
3. Call proximity propagation (30-60% confidence, 3 iterations) -- unclassified functions are assigned to the module voted by their callers (weighted 2x) and callees. A minimum of 2 votes is required. Each iteration propagates classifications outward from already-classified functions.
4. Code location heuristics (40% confidence) -- address range rules for known code regions (e.g., 0x2F00000-0x3000000 maps to register_allocation).

The eight modules are: optimization_framework, register_allocation, compilation_pipeline, ptx_emission, instruction_selection, error_handling, tensor_core_codegen, architecture_detection.

Output: foundation/taxonomy/modules/function_to_module_map.json, module_list.json.

Step 4: Analyze Call Graph (foundation/04_analyze_callgraph.py)

Computes three structural properties of the call graph:

• Entry points -- functions with zero callers and nonzero callees (top 100 by callee count). These are pipeline entry points, API functions, or global constructors.
• Leaf functions -- functions with zero callees and nonzero callers (top 1,000 by caller count). These are utility functions, allocators, and assertion handlers.
• Hot paths -- functions ranked by caller count (top 1,000). The highest-traffic functions in the binary.

Output: foundation/callgraph/entry_points.json, leaf_functions.json, hot_paths.json.

Step 5: Assign Priorities (foundation/05_assign_priorities.py)

Computes a composite priority score for each function to guide analysis effort allocation. The scoring formula:

• Size component: 1000 points for functions over 10 KB, 700 for 5-10 KB, 400 for 2-5 KB, 200 for 1-2 KB, 100 for 500 B-1 KB.
• Call frequency component: 500 points for 1000+ callers, 300 for 500+, 150 for 100+, 75 for 50+.
• Named function bonus: 200 points if the function has a recovered name (not sub_).
• Critical module bonus: 300 points if the function belongs to a critical module (compilation_pipeline, tensor_core_codegen, architecture_detection, register_allocation, instruction_selection, ptx_emission).

Functions scoring 1000+ are tier CRITICAL, 500+ are HIGH, 200+ are MEDIUM, below 200 are LOW.

Output: foundation/priorities/scoring_report.json, critical.json, high.json, medium.json, low.json.

Step 6: Generate Coverage Tracker (foundation/06_generate_coverage_tracker.py)

Aggregates all prior outputs into a master JSON tracker that records, per module and per function, the analysis status (pending/in-progress/complete), the assigned analyst, and the evidence quality score. This tracker serves as the coordination database for the L1 phase.

Output: foundation/coverage/tracker.json.

L1 Module Analysis Pipeline

The L1 pipeline runs via scripts/run_l1_programmatic.sh and requires L0 completion. It organizes CRITICAL and HIGH priority functions into per-module directories for systematic human review.

Step 1: Create Module Structure (modules/01_create_module_structure.py)

Creates the directory tree modules/{module}/functions/{critical,high}/ for each of the eight modules. MEDIUM and LOW tiers are intentionally excluded from L1 to focus effort on the most important functions.

Step 2: Extract Function Metadata (modules/02_extract_function_metadata.py)

For each CRITICAL and HIGH function, creates a directory modules/{module}/functions/{tier}/{address}/ containing a metadata.json file with: address, name, module, priority score, size, call frequency, scoring reasons, top 50 callers, top 50 callees, and paths to decompiled/disassembly/CFG files if they exist on disk.

Step 3: Generate Module READMEs (modules/03_generate_module_readmes.py)

Generates a skeleton README.md for each module with function counts, analysis progress tracking fields, and section headings for purpose, key functions, integration points, and data structures. These serve as the starting point for human-written module documentation.

Standalone Analysis Scripts

Six additional scripts perform targeted analyses independent of the L0/L1 pipeline:

analyze_nvvm_pipeline.py -- Loads the NVVM call graph (nvvm_callgraph.json, exported from the LibNVVM shared object analysis) and traces the compilation flow from nvvmCompileProgram. Identifies NVVM API entry points, finds LLVM optimization pass function symbols, traces call paths to depth 10, identifies hub functions (nodes with in-degree or out-degree above 10), and extracts the optimization pass ordering reachable from the compile entry point.

deep_pipeline_trace.py -- Performs deep BFS traversal (up to depth 15, width 100 per level) from nvvmCompileProgram through the NVVM call graph. Annotates each function with structural characteristics (LEAF, HUB, FANOUT, FANIN) and groups results by call depth to reveal the pipeline's stage boundaries. Also traces from secondary API entry points (nvvmVerifyProgram, nvvmAddModuleToProgram, nvvmCreateProgram).

extract_pipeline_structure.py -- Parses the 188,141 strings database for disable-*Pass patterns and Disable * description strings to extract the complete list of optimization passes by name. Categorizes passes into groups (Dead Code Elimination, Loop Optimizations, Inlining, Memory, NVVM-Specific, Lowering, etc.) and reconstructs the 13-stage compilation pipeline from NVVM module loading through PTX code generation. Also extracts compilation mode information (fast-compile, split-compile, partial-link).

analyze_performance_hotspots.py -- Loads the full function database (cicc_functions.json) and computes: global hotspot ranking (top 100 most-called functions), hot path chains (BFS from top 50 hotspots through callees, tracking weighted call frequency), size-efficiency analysis (bytes per call for each function), loop depth estimation (regex-based nesting analysis of decompiled C files), bottleneck identification (functions with 500+ callers), and module-level hotspot distribution.

catalog_optimization_framework.py -- Specialized script for the optimization_framework module. Reads per-function metadata from the L1 module directories, builds a critical function registry sorted by size, extracts HIGH-tier statistics (size tier distribution, top 20 most-called), scans decompiled code for optimization-related string patterns (pass references, iteration patterns, technique keywords), and identifies entry points (functions with 2 or fewer callers).

validate_callgraph.py -- Comprehensive validation system that cross-checks the call graph data against module classifications. Performs six verification analyses: cross-module call matrix verification (counting inter-module edges and sampling for spot-checks), entry point validation (confirming claimed entry points have zero callers), reachability analysis (BFS from main to find dead code), module dependency cycle detection (DFS on the module dependency graph), integration hotspot verification (functions called by all 8 modules), and bridge function identification (functions that both call into and are called from 2+ other modules).

Evidence Index Builders

Two versions of the evidence aggregation engine synthesize all data sources into per-function quality scores:

build_evidence_index.py (v1) -- Loads all five databases (functions, callgraph, strings, xrefs, names, comments, module map) into memory. For each of the 80,562 functions, counts eight evidence types (metadata, callers, callees, strings, xrefs, name pattern, size, module consistency) and computes a weighted confidence score (string evidence weighted highest at 20 points, callgraph at 15 each, xrefs at 15, metadata and name at 10 each, module at 10, size at 5). Produces nine output files including quality tier assignments (GOLD >= 80%, SILVER >= 50%, BRONZE < 50%), citation density analysis, cross-reference statistics, and prioritized recommendations for further analysis.

build_evidence_index_v2.py (v2, optimized) -- Memory-efficient reimplementation that avoids loading the full xref list into memory. Instead of building complete xref lookup tables, it streams the xref file line-by-line and counts only. The callgraph is preprocessed into a caller/callee count map rather than a full edge list. Produces the same nine analysis files as v1 with identical quality tier logic. Recommended for systems with less than 32 GB RAM.

Cross-Module Dependency Analysis

07_analyze_cross_module_dependencies.py -- The most complex standalone analysis. Streams the full call graph (using ijson for memory-efficient parsing) four times to compute:

1. Inter-module call matrix -- for each pair of the 8 modules, the number of call edges crossing the boundary.
2. Module dependency depth -- per-module statistics on how many other modules each function depends on, identifying isolated functions and hub functions.
3. Critical bridges -- functions that call into 3 or more other modules (top 100 by bridge count).
4. Integration hotspots -- functions called by 3 or more other modules (top 100 by fan-in).
5. Module dependency graph -- a JSON graph structure with weighted edges suitable for visualization.
6. Integration patterns -- entry point modules (highest out-degree), utility hub modules (highest in-degree), and linear dependency chains.

Data Flow and Directory Structure

The complete analysis data is organized as follows:

cicc/
  databases/                    # IDA exports (input data)
    cicc_functions.json         #   80,562 function records
    cicc_strings.json           #   188,141 string records
    cicc_xrefs.json             #   cross-reference records
    cicc_callgraph.json         #   call edge records
    cicc_names.json             #   recovered names
    cicc_comments.json          #   IDA comments
  foundation/                   # L0 pipeline output
    indices/                    #   pickle indices for fast lookup
    taxonomy/
      modules/                  #   function-to-module map, module list
      strings/                  #   classified string databases
    callgraph/                  #   entry points, leaf functions, hot paths
    priorities/                 #   priority scoring and tier assignments
    coverage/                   #   master progress tracker
    analyses/                   #   evidence index, quality tiers, cross-module data
  modules/                      # L1 pipeline output
    {module}/
      functions/
        critical/{addr}/        #   metadata.json per critical function
        high/{addr}/            #   metadata.json per high function
      analysis/                 #   module-level analysis files
      README.md                 #   module documentation skeleton
  decompiled/                   # Hex-Rays output (per-function C files)
  disasm/                       # IDA disassembly output (per-function ASM files)
  graphs/                       # Control flow graphs (JSON and DOT)
  scripts/                      # All analysis scripts
    foundation/                 #   L0 pipeline scripts (00-07)
    modules/                    #   L1 pipeline scripts (01-03)
    run_foundation_analysis.sh  #   L0 batch runner
    run_l1_programmatic.sh      #   L1 batch runner

Verification Approaches

To verify any specific finding in this wiki:

1. Open IDA at the stated address. Every function identification includes an address. Navigate to it, press F5 to decompile, and check whether the decompiled code matches the described behavior.

2. Check string xrefs. For VERY HIGH and KNOWN identifications, search for the quoted string in IDA's Strings window. The xref should lead to the stated function address or a function that directly calls it.

3. Compare with upstream LLVM. CICC is based on LLVM 20.0.0. The LLVM source tree at the corresponding git tag contains the original implementations of all standard passes. Structural comparison (switch case counts, parameter counts, error message text) between the decompiled code and the LLVM source is the gold standard for verification.

4. Cross-reference the dual paths. Path A and Path B contain near-duplicate code. If a function is identified in Path A, the corresponding Path B function should exhibit the same structure. Agreement between the two paths increases confidence.

5. Trace from known entry points. Start at sub_8F9C90 (real main, KNOWN confidence) and follow the call chain. Every function reachable from main through a chain of identified functions has a verified callgraph path.

6. Run the validation script. Execute scripts/validate_callgraph.py to cross-check the call graph against module classifications. The script produces a CALLGRAPH_VALIDATION_REPORT.json with quantitative metrics: entry point accuracy, cross-module call counts, reachability percentage, bridge function inventory, and module dependency cycles. A healthy analysis should show entry point confidence above 90% and reachability above 80%.

7. Re-run the evidence index. Execute scripts/foundation/build_evidence_index_v2.py to regenerate quality tier assignments. Compare the GOLD/SILVER/BRONZE percentages against the expected distribution (majority SILVER or GOLD for classified functions). Functions that drop to BRONZE after a wiki edit indicate a regression in evidence consistency.

Reproducing the Full Analysis

To reproduce this analysis from scratch:

1. Obtain the binary. Install CUDA Toolkit 13.0. The binary is at <cuda>/nvvm/bin/cicc. SHA-256 and build string cuda_13.0.r13.0/compiler.36424714_0 must match.

2. Run IDA auto-analysis. Open cicc in IDA Pro 8.x with default x86-64 analysis settings. Allow auto-analysis to complete (5-10 minutes for a binary of this size). Accept the detected compiler (GCC).

3. Batch decompile. Run the following IDA Python script to decompile all functions and export per-function C files:

   import idautils, ida_hexrays, idc
   for func_ea in idautils.Functions():
       try:
           cfunc = ida_hexrays.decompile(func_ea)
           name = idc.get_func_name(func_ea)
           addr = f"0x{func_ea:X}"
           with open(f"decompiled/{name}_{addr}.c", "w") as f:
               f.write(str(cfunc))
       except:
           pass
   

4. Export databases. Use IDA Python to export the five JSON databases (functions, strings, xrefs, callgraph, names) to cicc/databases/. The function export should iterate Functions() and record address, name, size, instruction count, is_library, is_thunk, callers, and callees for each. The string export should iterate IDA's string list and record address, value, and xrefs.

5. Run L0 foundation pipeline.

   cd cicc/scripts
   bash run_foundation_analysis.sh
   

   This executes Steps 0-6 in sequence, producing all indices, classifications, and the coverage tracker. Expected runtime: 2-5 minutes on a modern machine.

6. Run L1 module setup.

   bash run_l1_programmatic.sh
   

   This creates the per-module directory structure, extracts metadata for CRITICAL and HIGH functions, and generates module README skeletons. Expected runtime: under 1 minute.

7. Run standalone analyses (optional, for deeper investigation):

   python3 analyze_nvvm_pipeline.py       # NVVM pipeline trace
   python3 deep_pipeline_trace.py          # Deep BFS from nvvmCompileProgram
   python3 extract_pipeline_structure.py   # Pass extraction from strings
   python3 analyze_performance_hotspots.py # Hotspot ranking
   python3 validate_callgraph.py           # Validation report
   

8. Run evidence indexing (optional, for quality assessment):

   cd foundation
   python3 build_evidence_index_v2.py
   

9. Begin manual analysis. With the foundation data in place, start from the CRITICAL priority list and the string anchors described in the Function Identification Strategies section above. The Function Map page is the primary lookup table.

Dependencies

The analysis scripts require only the Python 3.8+ standard library with one exception: 07_analyze_cross_module_dependencies.py uses ijson for streaming JSON parsing of the large callgraph file. Install with pip install ijson. All other scripts use only json, pickle, re, collections, pathlib, statistics, dataclasses, and typing.

Binary Address Sweep Reports

In addition to the automated scripts, the analysis produced 90+ raw binary sweep reports stored in cicc/raw/. Each report covers a contiguous address range (typically 128 KB to 512 KB) and contains per-function identification notes, string evidence citations, structural observations, and confidence assessments. The reports are named by address range (e.g., p1.3-01-sweep-0x8F0000-0x90FFFF.txt covers the compilation pipeline entry region) and organized into 10 sweep phases corresponding to the binary's major sections. A second sweep phase (p2-* and p2a-p2g) provides focused analyses of specific subsystems (EDG frontend, IR generation, optimization passes, SelectionDAG, register allocation, scheduling, configuration).

These raw reports are the primary source material from which the wiki pages were written. They are not cleaned or edited for presentation -- they contain working notes, false starts, and corrections made during the analysis process.

Limitations and Known Gaps

This analysis has several inherent limitations:

• No dynamic validation. All findings are from static analysis. Runtime behavior under specific inputs (unusual SM targets, edge-case CUDA constructs) has not been verified.
• EDG internals are partially opaque. The EDG frontend is a licensed third-party component. Its internal data structures are less well-documented in the LLVM literature, making identification harder. The IL tree format and scope management structures are identified at MEDIUM confidence.
• Inlined functions are invisible. If the compiler inlined a function during the build of cicc itself, that function has no standalone address and cannot be independently identified. Some small LLVM utility functions (SmallVector operations, StringRef comparisons) are likely inlined throughout.
• Proprietary NVIDIA code has no public reference. The 35 custom NVIDIA passes, the NVVM bridge layer, and the NVVMPassOptions system have no upstream source to compare against. These are identified purely from string evidence and structural analysis.
• Version-specific. All findings apply to cicc v13.0 (build cuda_13.0.r13.0/compiler.36424714_0). Addresses, function sizes, and pass counts will differ in other CUDA toolkit versions.
• Module classification accuracy degrades at the boundary. The automated taxonomy assigns ~60% of functions with high confidence (wiki ground truth or strong string evidence). The remaining functions are classified by call proximity propagation or address range heuristics at 30-60% confidence. Functions at module boundaries may be misclassified; the validate_callgraph.py script quantifies this.
• Callgraph completeness depends on IDA's xref analysis. Indirect calls through function pointers (vtable dispatch, callback registrations) are not fully captured by IDA's static analysis. The call graph is therefore a lower bound on the true call relationships. This primarily affects LLVM's pass manager dispatch and the EDG frontend's visitor pattern implementations.

Version Tracking

This page documents the exact version identifiers embedded in the cicc v13.0 binary and the version relationships between its components. Every version listed here was recovered from string constants, constructor initializations, or binary header fields in the stripped ELF binary. This is the single source of truth for version-related questions across the wiki.

Version Summary

LLVM Version: The Dual Identity

CICC has two LLVM version identities. Internally, it is an LLVM 20.0.0 fork -- all modern instruction opcodes, metadata formats, type encodings, and pass infrastructure from LLVM 20 are present. Externally, the bitcode it emits identifies itself as "LLVM7.0.1" in the producer field.

The reason is historical: NVVM IR 2.0 was defined against LLVM 7.0.1. The entire NVVM toolchain ecosystem (libNVVM, nvcc's device pipeline, nvdisasm, third-party NVVM IR consumers) standardized on "LLVM7.0.1" as the format identifier. Changing the producer string would require a coordinated update across the entire CUDA toolkit and all downstream consumers.

Binary evidence:

• ctor_036 at 0x48CC90: reads LLVM_OVERRIDE_PRODUCER environment variable, falls back to "20.0.0" (the true version).
• ctor_154 at 0x4CE640: reads LLVM_OVERRIDE_PRODUCER, falls back to "7.0.1" (the compatibility marker). This is the constructor that runs for the bitcode writer path.
• sub_E7A190: contains the string "llvm-mc (based on LLVM 20.0.0)".
• sub_1538EC0 (writeModule): emits "LLVM" + "7.0.1" = "LLVM7.0.1" as the IDENTIFICATION_BLOCK producer.

Both constructors accept the LLVM_OVERRIDE_PRODUCER environment variable to override the default. Setting it changes the embedded producer string in output bitcode.

See Bitcode Reader/Writer for the full dual-producer mechanism.

EDG 6.6 Frontend

The EDG (Edison Design Group) frontend is a licensed commercial C/C++ frontend. Version 6.6 occupies 3.2 MB of code at 0x5D0000--0x8F0000. The version string is embedded literally as "Based on Edison Design Group C/C++ Front End, version 6.6" and is accessible via the --version flag.

EDG 6.6 in cicc is configured to emulate GCC 8.1 (DEFAULT_GNU_VERSION = 80100) and Clang 9.1 (DEFAULT_CLANG_VERSION = 90100). It supports C++23 as the newest C++ standard and C23 as the newest C standard, with C++17 as the default mode.

See EDG 6.6 Frontend for the full frontend documentation.

NVVM IR Version

The NVVM IR version is a metadata tuple (major, minor) embedded in every NVVM bitcode module via the !nvvmir.version named metadata node. CICC v13.0 has two distinct version contexts:

User code: the IR generation phase (sub_9151E0) emits !nvvmir.version with the current version tuple. The version checker at sub_157E370 enforces major == 3 and minor <= 2, making 3.2 the current maximum accepted version. Modules with major != 3 or minor > 2 are rejected with "Broken module found, compilation aborted!".

Libdevice: the embedded libdevice.10.bc carries !nvvmir.version = !{i32 2, i32 0}. The version (2, 0) is hard-coded in the version checker (sub_12BDA30) as an always-compatible sentinel -- it passes the check regardless of the current NVVM IR version. This ensures the embedded math library is compatible with any user module.

Container format: the NVVM container binary header stores version fields separately at offsets 0x06--0x07 (nvvm_ir_major, nvvm_ir_minor). These track the container-level IR spec version and may differ from the bitcode-level metadata tuple.

Bypass: setting NVVM_IR_VER_CHK=0 in the environment disables version validation entirely, allowing any version tuple to pass.

See Bitcode Reader/Writer for the version gate implementation and NVVM Container for the container-level version fields.

Embedded Libdevice

The embedded libdevice is libdevice.10.bc, a 455,876-byte LLVM bitcode library containing 352 GPU-optimized math functions. Two identical copies are statically embedded in the binary:

Key properties:

• Target triple: nvptx64-nvidia-gpulibs
• Function count: 352 (all alwaysinline nounwind)
• NVVM IR version: (2, 0) -- always-compatible sentinel
• Producer: "clang version 3.8.0 (tags/RELEASE_380/final)" -- the Clang version that originally compiled libdevice (not indicative of cicc's own compiler version)
• NVVMReflect calls: uses __nvvm_reflect("__CUDA_FTZ"), __nvvm_reflect("__CUDA_ARCH"), and __nvvm_reflect("__CUDA_PREC_SQRT") for runtime specialization

The libdevice.10.bc naming convention carries forward from the CUDA 5.0 era. The 10 in the filename originally indicated "compute capability 1.0 and above" (i.e., universal), not a version number.

See Libdevice Linking for the linking algorithm, version validation, and NVVMReflect interaction.

NVVM Container Format Version

The NVVM container binary envelope uses its own versioning scheme, independent of the NVVM IR version:

The container's LLVM version encoding stores the combined value 20 * 100 + 0 = 2000, confirming the internal LLVM 20.0.0 base.

See NVVM Container for the full binary format specification.

Version Cross-Reference Matrix

How versions flow through the pipeline:

                    EDG 6.6 Frontend
                         |
                         v
                    NVVM IR Generation
                    (emits nvvmir.version = {3, 2})
                         |
                    +----+----+
                    |         |
              libdevice    user IR
           (version 2,0) (version 3,2)
                    |         |
                    +----+----+
                         |
                    NVVM IR Version Check
                    (gate: major==3, minor<=2)
                    (sentinel: 2,0 always passes)
                         |
                    LLVM 20.0.0 Optimizer
                         |
                    Bitcode Writer
                    (producer: "LLVM7.0.1")
                         |
                    NVVM Container Serializer
                    (container version 1.x, LLVM encoded as 2000)
                         |
                         v
                    .ptx / .optixir output

Future Updates

This wiki documents cicc v13.0 from CUDA 12.8. When a new CUDA toolkit release ships a newer cicc binary, the following version fields are the most likely to change:

• LLVM base version: NVIDIA periodically rebases on newer LLVM releases. A jump from 20.0.0 to a later version would change the internal string, the container LLVM encoding, and potentially add new passes, opcodes, and metadata formats.
• EDG version: EDG releases track independently of LLVM. A bump from 6.6 to a later version would affect C++ standard support, keyword handling, and the frontend error catalog.
• NVVM IR version minor: the minor field (currently 2 in the major == 3 series) may increment to accommodate new metadata kinds or intrinsic conventions without breaking the major version.
• PTX ISA version: new SM targets require new PTX versions. sm_100 Blackwell already uses a higher PTX version than sm_90 Hopper.
• SM target range: new GPU architectures add new SM numbers. The sm_75--sm_121 range in v13.0 will expand in future releases.

The bitcode producer string ("LLVM7.0.1") is unlikely to change in the near term -- doing so would break backward compatibility with the entire NVVM IR ecosystem. The libdevice version sentinel (2, 0) is similarly stable because the version checker special-cases it.

To update this wiki for a new cicc version:

1. Extract the build string (search for cuda_XX.Y.rXX.Y/compiler.).
2. Check ctor_036 for the LLVM version fallback string.
3. Check the EDG version string at sub_617BD0.
4. Check the NVVM IR version gate constants at the version checker function.
5. Measure the embedded libdevice size and function count.
6. Verify the NVVM container header version fields.

Cross-References

• Bitcode Reader/Writer -- producer string mechanism, version gate implementation
• NVVM Container -- container binary format with version header fields
• Libdevice Linking -- embedded math library, version sentinel
• EDG 6.6 Frontend -- frontend version, GCC/Clang emulation modes
• Binary Layout -- build ID string, ELF properties
• Entry Point & CLI -- dual-path dispatch, version string arguments
• Environment Variables -- LLVM_OVERRIDE_PRODUCER, NVVM_IR_VER_CHK
• Debug Info Verification -- debug version field (3.2)

Compilation Pipeline Overview

This page maps the complete end-to-end flow of a CUDA compilation through cicc v13.0, from the initial CLI invocation to the final PTX text output. Each stage is a self-contained subsystem with its own address range, data structures, and failure modes. The links below lead to dedicated pages with reimplementation-grade detail for every stage.

Pipeline Diagram

nvcc
  |
  v
+===========================================================+
| cicc (60 MB, 80,562 functions)                            |
|                                                           |
|  1. CLI Parsing & Dispatch -----> [entry.md]              |
|     |  argv/envp, flag translation, arch detection        |
|     |  dual-path select: Path A (LibNVVM) / Path B        |
|     v                                                     |
|  2. nvcc-to-cicc Interface -----> [nvcc-interface.md]     |
|     |  flag tree (40+ mappings), 3-column arch fan-out    |
|     |  mode cookies: 0xABBA=CUDA, 0xDEED=OpenCL          |
|     v                                                     |
|  3. EDG 6.6 Frontend -----------> [edg.md]                |
|     |  CUDA C++ --> transformed C (.int.c/.device.c)      |
|     |  737 config #defines, GCC 8.1 / Clang 9.1 emu      |
|     v                                                     |
|  4. NVVM IR Generation ---------> [ir-generation.md]      |
|     |  EDG IL tree --> LLVM Module (NVVM IR)              |
|     |  address spaces, kernel metadata, builtins          |
|     v                                                     |
|  5. Libdevice Linking ----------> [../infra/libdevice-linking.md]
|     |  embedded 455KB bitcode, 352 __nv_* math fns        |
|     |  target triple validation, NVVM version check       |
|     v                                                     |
|  6. LLVM Optimizer -------------> [optimizer.md]          |
|     |  two-phase model (analysis -> codegen-oriented)     |
|     |  49.8KB pipeline assembler, ~150 pass insertions    |
|     |  concurrent per-function Phase II                   |
|     v                                                     |
|  7. LTO Pipeline ---------------> [../lto/index.md]       |
|     |  cross-TU inlining, devirt, GlobalOpt               |
|     |  closed-world GPU model: no dlopen, no .so          |
|     v                                                     |
|  8. Code Generation ------------> [codegen.md]            |
|     |  SelectionDAG, ISel, RegAlloc, MachineIR passes     |
|     |  37 MB of code, largest subsystem                   |
|     v                                                     |
|  9. PTX Emission ---------------> [emission.md]           |
|     |  .entry/.func headers, register decls, .loc/.file   |
|     |  AsmPrinter, GenericToNVVM addrspace rewrite        |
|     v                                                     |
|  OUTPUT: .ptx file (or NVVM bitcode, or OptiX IR)        |
+===========================================================+

Side paths:
  * OptiX IR (--emit-optix-ir) ----> [optix-ir.md]
  * Debug info (all stages) -------> [debug-info-pipeline.md]

Stage Descriptions

1. Entry Point & CLI Parsing

The real main (sub_8F9C90, 10KB) parses argv, detects wizard mode via NVVMCCWIZ=553282, selects the target architecture (default sm_75), and dispatches into one of two compilation paths. Path A serves the LibNVVM API; Path B serves standalone nvcc invocations. Both paths are functionally identical but duplicated in the binary at different address ranges. See Entry Point & CLI.

2. nvcc-to-cicc Interface

The flag translation layer (sub_8FE280) rewrites nvcc-facing flags into cicc-facing flags through a std::map red-black tree, then a second stage (sub_95EB40) fans each flag out into three columns targeting EDG, OPT, and LLC separately. Mode cookies (0xABBA for CUDA, 0xDEED for OpenCL) select language-specific behavior. See nvcc-to-cicc Interface.

3. EDG 6.6 Frontend

A licensed commercial frontend (3.2 MB, 0x5D0000--0x8F0000) parses CUDA C++ source and emits transformed C code into .int.c, .device.c, and .stub.c files. CUDA syntax (<<<>>>, __shared__, __device__) is fully resolved in this stage. The output is C source, not LLVM IR. See EDG 6.6 Frontend.

4. NVVM IR Generation

Translates the EDG intermediate language (IL) tree into an LLVM Module with proper NVPTX address space annotations, nvvm.annotations kernel metadata, and lowered builtins. This is cicc's equivalent of Clang's lib/CodeGen, but operates on EDG's proprietary IL node format. See NVVM IR Generation and its sub-pages for expressions, statements, functions, and types.

5. Libdevice Linking

A 455,876-byte LLVM bitcode library containing 352 GPU-optimized math functions (__nv_sinf, __nv_expf, etc.) is embedded directly in the cicc binary. The linker validates the nvptx64- target triple, checks NVVM IR version metadata, and merges the library into the compilation module. No filesystem access is required. See Libdevice Linking.

6. LLVM Optimizer

A proprietary two-phase pipeline (sub_12E54A0, 49.8KB) runs ~150 passes: Phase I performs module-wide analysis, Phase II performs codegen-oriented transforms with optional per-function parallelism using a jobserver or thread pool. All behavior is controlled by the 222-slot NVVMPassOptions system. See LLVM Optimizer and Pipeline & Ordering.

7. LTO Pipeline

Exploits the GPU's closed-world compilation model (no dlopen, no shared libraries, no symbol interposition) for aggressive cross-TU inlining, whole-program devirtualization, and global variable promotion. Activated in separate compilation mode (nvcc -dc), but GlobalOpt and the inliner run even in single-TU mode. See LTO & Module Optimization.

8. Code Generation

The largest subsystem (37 MB, 0x1700000--0x35EFFFF) lowers optimized LLVM IR to NVPTX MachineInstr through SelectionDAG construction, type legalization, instruction selection via a three-level pattern match engine (900KB), pressure-driven greedy register allocation, and ~30 machine-level passes including tensor core codegen for HMMA/IMMA/WGMMA/tcgen05. See Code Generation.

9. PTX Emission

The AsmPrinter (sub_31EC4F0, 72KB) walks the final MachineFunction and emits PTX text: .entry/.func headers with kernel attributes, register declarations for 9 register classes, .loc/.file debug directives, and instruction mnemonics. A GenericToNVVM pass rewrites any remaining generic address space references before emission. See PTX Emission.

Side Paths

OptiX IR -- When --emit-optix-ir is passed, the pipeline replaces LLC with an OPTIXIR stage that serializes the optimized LLVM module for the OptiX ray tracing runtime's continuation-based execution model. See OptiX IR Generation.

Debug Info -- Debug metadata flows through all stages: generated in IR-gen, preserved or stripped in the optimizer (5 stripping passes), verified after each pass, and emitted as .loc/.file PTX directives. See Debug Info Pipeline.

Internal Pipeline Encoding

Internally, cicc represents the active pipeline stages as a bitmask:

The standard CUDA compilation bitmask is LNK | OPT | LLC = 0x07. OptiX mode uses 0x43.

Cross-References

• Binary Layout -- address ranges for every subsystem
• Function Map -- master index of recovered function addresses
• CLI Flags -- complete flag catalog
• Optimization Levels -- what changes at -O0/-O1/-O2/-O3
• NVIDIA Custom Passes -- 35 proprietary passes inserted into the LLVM pipeline
• NVPTX Target Infrastructure -- TargetMachine, TTI, SubtargetFeatures

Entry Point & CLI

The cicc binary has a surprisingly complex entry point. Rather than a straightforward main → compile → exit flow, it implements a dual-path architecture where the same binary can operate as either a LibNVVM-based compiler (Path A) or a standalone compiler (Path B), selected at runtime through environment variables and obfuscated string comparisons. This design allows NVIDIA to ship a single binary that serves both the nvcc toolchain and the LibNVVM API.

The entry point region (0x8F0000–0x96FFFF, ~520 KB) handles CLI parsing, architecture detection with a 3-column flag fan-out system, and dispatch into one of several compilation pipelines. A hidden "wizard mode" gated behind an environment variable with a magic number enables developer diagnostics that are otherwise completely inaccessible.

Architecture

main (0x4396A0, 16B thunk)
  │
  └─ sub_8F9C90 (10KB, REAL MAIN)
       │
       ├─ getenv("NVVMCCWIZ") == 553282 → wizard mode
       ├─ sub_16C5290: extract program name from argv[0]
       │
       ├─ ARGUMENT LOOP (v15 = 1..argc)
       │    ├─ -o <file>              → v257 (output)
       │    ├─ -nvvmir-library <path> → v256 (libdevice)
       │    ├─ -lgenfe/-libnvvm/-lnk/-opt/-llc → v263 (mode)
       │    ├─ -arch/-mcpu/--nv_arch  → v242 (SM number)
       │    ├─ --emit-optix-ir        → v243=1, v258=1
       │    ├─ -nvc                   → v258=1
       │    ├─ -irversion             → print IR version, exit
       │    ├─ .bc/.ci/.i/.ii/.cup/.optixir → s (input file)
       │    └─ obfuscated option      → v253 (0 or 1)
       │
       ├─ v253 RESOLUTION (if still == 2)
       │    └─ getenv(obfuscated) → compare → set v253 = 0 or 1
       │
       ├─ DISPATCH (v263 × v253)
       │    ├─ v263==0, v253==1 → sub_902D10  (simple Path A)
       │    ├─ v263==0, v253==0 → sub_1262860 (simple Path B)
       │    ├─ v263==1          → sub_905E50 / sub_12658E0 (lgenfe)
       │    ├─ v263≥2, v253==1  → sub_905EE0  (multi-stage Path A)
       │    └─ v263≥2, v253==0  → sub_1265970 (multi-stage Path B)
       │
       └─ CLEANUP: free all vectors, strings, argv copy

Real Main — sub_8F9C90

The exported main() at 0x4396A0 is a 16-byte thunk that immediately tail-calls sub_8F9C90 — the actual entry point. This function is a monolithic CLI parser and dispatcher: it copies argv into a local buffer, checks for wizard mode, iterates over all arguments accumulating state in ~12 local variables, resolves the compilation path, and finally dispatches to the appropriate pipeline function. The entire function is a single 10KB basic-block-heavy control flow graph with ~80 branch targets.

Argument Handling and Argv Copy

The function begins with a defensive copy of argv into a local buffer. When 8 * argc fits within 0x800 bytes (argc ≤ 256), the copy lives in v284[2096] on the stack. For larger argument lists -- which can occur during complex nvcc invocations with many pass-through flags -- it allocates heap memory via sub_16CD150. This copy is necessary because the argument loop modifies pointers (advancing i to skip flag values), and the caller's argv must not be disturbed.

if (8 * argc > 0x800)
    v284 = sub_16CD150(8 * argc);   // heap alloc for large argc
// else use stack buffer v284[2096]
memcpy(v284, argv, 8 * argc);       // copy all pointers

After copying, sub_16C5290 extracts the base program name from argv[0] -- stripping directory prefixes -- and stores it in dest. This name appears in error messages and verbose output throughout the pipeline.

Key Local Variables

The function's behavior is controlled by two critical dispatch variables: v253 (which compilation backend to use) and v263 (which phase of the pipeline to invoke). These are accumulated during the argument loop and combined after parsing to select one of ~10 possible code paths. The interaction between them creates a matrix of behaviors that covers everything from simple single-file compilation to multi-stage LibNVVM pipeline processing.

Wizard Mode

v10 = getenv("NVVMCCWIZ");                    // 0x8F9D36
if (v10 && strtol(v10, NULL, 10) == 553282)   // 0x8F9D92
    byte_4F6D280 = 1;

Global byte_4F6D280 gates the effectiveness of -v, -keep, -dryrun. Without wizard mode, these flags are silently ignored — v259 and v262 stay 0. This is a deliberate anti-reverse-engineering measure: even if someone discovers the -v flag, it does nothing without the magic environment variable. The magic number 553282 (0x87142) appears to be arbitrary.

Invocation Modes (v263)

The v263 variable determines which stage of the compilation pipeline cicc enters. When nvcc invokes cicc directly, v263 stays at 0 (default). But cicc can also be invoked in sub-pipeline mode — for example, -lnk runs only the linking phase, -opt runs only the optimizer, and -llc runs only code generation. This is how the multi-stage pipeline works: the outer driver calls cicc multiple times with different -lXXX flags, or a single invocation with -libnvvm runs all stages internally.

Each mode has its own format for the -discard-value-names flag, which tells the LLVM backend whether to strip IR value names (reducing memory usage). The different formats exist because each sub-pipeline stage has its own option namespace:

Input File Extensions

Input files are identified by extension during the argument loop. The last matching file wins (s is overwritten each time). Unrecognized arguments are added to the v266 pass-through vector and forwarded to sub-pipelines. The .cup extension has a special restriction — it's only accepted when the preceding argument is --orig_src_path_name or --orig_src_file_name, which are metadata flags inserted by nvcc to track the original source file.

Obfuscated Strings

At 0x8F98A0, sub_8F98A0 decrypts strings using an XOR + ROT13-like cipher:

v40 = v37 ^ (-109 * ((offset + 97) ^ 0xC5));
// then ROT13 on alphabetic characters

This hides an environment variable name and option prefix from static analysis. The decrypted strings control the v253 (Path A vs Path B) resolution when no explicit mode is specified.

Error Messages

The v253 Dispatch Variable

The v253 variable is the single most important dispatch control in the entire entry point. It determines whether the compilation uses Path A (the EDG/PTX-producing pipeline) or Path B (the standalone LLVM-based pipeline). Understanding its resolution logic is essential to reproducing cicc's behavior.

Initialization and Explicit Setting

v253 begins at 2 (unresolved default). During the argument loop, obfuscated string matching can set it directly:

Environment Variable Resolution

When v253 remains at 2 after argument parsing (the common case), cicc resolves it through the obfuscated environment variable NV_NVVM_VERSION (decrypted from byte_3C23A9F). The resolution has two sub-cases depending on the target architecture:

if (v253 == 2) {
    env = getenv(decrypt(byte_3C23A9F));   // NV_NVVM_VERSION
    if (env matches decrypt(byte_3C23A82))       // "nvvm-latest"
        v253 = 1;  // Path A
    else if (env matches decrypt(byte_3C23A7B))  // "nvvm70"
        v253 = 0;  // Path B
    else if (v242 > 99 && !v258)                 // SM >= 100, not -nvc
        v253 = 0;  // Path B (new architectures default to standalone)
    else
        v253 = 1;  // Path A (legacy default)
}

The architectural threshold at SM 100 (Blackwell) is notable: for SM < 100, the default is Path A (the EDG frontend path). For SM >= 100, unless the -nvc flag is present, the default switches to Path B. This suggests NVIDIA is migrating newer architectures toward the standalone LLVM pipeline, possibly as a precursor to eventually deprecating the EDG-based path.

Version Strings Injected per Path

After v253 is resolved and for multi-stage modes (v263 >= 3), the entry point injects a version string into the pass-through options:

This version string propagates through the entire pipeline, controlling bitcode compatibility, intrinsic name resolution, and metadata format expectations.

Post-Parse Dispatch Logic

After the argument loop terminates, the dispatch logic combines v253 and v263 to select the target function. The combined keep-and-verbose flag v260 = v262 & v259 is also computed -- both wizard-mode flags must be active for intermediate file retention and verbose logging to function simultaneously.

Simple Dispatch (v263 == 0)

When cicc is invoked without any -lXXX mode flag (the standard nvcc invocation path):

if (v253 == 1)
    v8 = sub_902D10(dest, 0, &v266, s, v257, v256, v260, v262, v261);
    // Path A: CLI → lgenfe → LibNVVM pipeline
else
    v8 = sub_1262860(dest, 0, &v266, s, v257, v256, v260, v262, v261);
    // Path B: CLI → standalone LLVM pipeline

Both functions receive identical parameter signatures: program name, zero (unused), pass-through options, input file, output file, libdevice path, verbose+keep, keep, and dryrun. The return value becomes the process exit code.

lgenfe Dispatch (v263 == 1)

The -lgenfe mode builds a full argv-style array with the program name as the first entry, followed by all v266 pass-through options. This argv is then passed to one of two function pairs:

The init functions create the LLVM context and register the 42+ metadata kinds used throughout the pipeline (dbg, tbaa, prof, noalias, etc.). These must be registered before any IR construction begins.

Multi-Stage Dispatch (v263 >= 2)

For -libnvvm, -lnk, -opt, and -llc modes, the dispatch constructs a CompilationState structure with input/output strings, extra arguments, and the v278 mode byte, then calls:

For -libnvvm (v263 == 2), the extra args are taken directly from v266 without prepending the program name. For -lnk/-opt/-llc (v263 >= 3), the appropriate version string (nvvm-latest or nvvm70) is appended to the pass-through options before dispatch.

Cleanup

After the pipeline function returns, sub_8F9C90 performs deterministic cleanup in reverse allocation order: the v281 extra-argument char** array and each entry, the v275 output string, the s2 input string, each element of the v266 pass-through vector, the vector's backing buffer, the dest program name, and the v282 argv copy buffer (if heap-allocated). The return value v8 is 0 on success, 1 on argument errors, or the pipeline function's return code (stored in v264).

Path A — EDG → LibNVVM Pipeline

Path A is the full CUDA C++ compilation path. It starts with the EDG 6.6 C++ frontend parsing CUDA source code into an IL tree, then converts that IL into LLVM IR via the lgenfe (LLVM Generation Front End) stage, and finally runs the LibNVVM pipeline to optimize and lower the IR to PTX. This is the path taken when cicc is invoked by nvcc for .cu file compilation, and it represents the standard CUDA compilation flow that most users encounter.

Path A Orchestrator — sub_902D10

The orchestrator is a 9 KB function that sequences the three major stages of Path A compilation. It acts as the conductor between the CLI processing layer, the EDG frontend, and the LibNVVM optimizer/codegen.