cudafe++ v13.0 -- Reverse Engineering Reference

cudafe++ is NVIDIA's CUDA frontend compiler -- the first stage of the CUDA compilation pipeline. It is built on the Edison Design Group (EDG) C++ Front End v6.6, a commercial compiler frontend licensed by compiler vendors worldwide. NVIDIA ships cudafe++ as a statically-linked, stripped ELF binary inside every CUDA Toolkit installation. This binary accepts .cu source files, parses them as C++ with CUDA extensions, separates device code from host code, and produces two outputs: an EDG Intermediate Language (IL) stream consumed by cicc (the NVIDIA PTX code generator), and a transformed .int.c host file consumed by the system C++ compiler (gcc, clang, or cl.exe).

This wiki documents the complete internals of the cudafe++ binary from CUDA Toolkit 13.0, reverse-engineered through static analysis (IDA Pro + Hex-Rays decompilation) of all 6,483 functions. The goal is reimplementation-grade documentation: every page should give a senior compiler engineer enough information to build equivalent functionality from scratch.

Binary Identity

Segment Layout

Role in the CUDA Toolchain

  input.cu
     |
     v
 cudafe++  ──────── THIS BINARY ────────
     |                                   |
     v                                   v
  device.gpu  (EDG IL)             input.int.c  (transformed host C++)
     |                                   |
     v                                   v
   cicc                            gcc / clang / cl.exe
     |                                   |
     v                                   v
  device.ptx                       host.o
     |                                   |
     v                                   v
   ptxas                              ld
     |                                   |
     v                                   v
  device.cubin ──────────────────> final executable

cudafe++ is a source-to-source compiler. It never generates machine code directly. Its job is to take a single .cu translation unit, understand which code is device (__device__, __global__) and which is host, then:

1. For the device track: Emit EDG IL -- a typed, scope-linked intermediate representation containing every declaration, type, expression, and statement. This IL is consumed by cicc, which lowers it through LLVM to PTX assembly.

2. For the host track: Emit a .int.c file -- valid C++ source where device function bodies are suppressed inside #if 0/#endif, __global__ kernels are replaced by __wrapper__device_stub_<name>() forwarding functions, and CUDA runtime registration boilerplate is appended.

The binary runs as a single-threaded, single-pass-per-stage pipeline with 8 stages: pre-init, CLI parsing (276 flags), one-time init (38 subsystem initializers), TU state reset, frontend parse (EDG parser + CUDA extensions), 5-pass IL finalization, backend .int.c emission, and exit. See Pipeline Overview for the full stage diagram.

Source Attribution

The binary embeds __FILE__ strings from the EDG build system, revealing the original source file structure. From these strings plus address-range analysis of decompiled code, 52 .c source files and 13 .h header files have been identified:

The NVIDIA-specific source files are:

• nv_transforms.c (~34 functions, ~14 KB of .text): The heart of CUDA support. Implements device/host-device lambda wrapper template generation (__nv_dl_wrapper_t, __nv_hdl_wrapper_t, __nv_hdl_create_wrapper_t), CUDA attribute validation (__launch_bounds__, __cluster_dims__, __block_size__, __maxnreg__), host reference array emission (.nvHRKI/.nvHRDE/.nvHRCE ELF sections), lambda preamble injection (sub_6BCC20), and array capture helper generation.

• nv_transforms.h: Header with NVIDIA-specific declarations, type trait template names, and bitmask table definitions.

• 3 modified EDG files: cmd_line.c (CUDA CLI flags spliced into EDG's flag table), fe_init.c (CUDA-specific initialization at stage 3), and cp_gen_be.c (device stub generation, lambda wrapper emission, registration table output in the backend).

Key Discoveries

Execution Space Bitfield

Every entity node in the EDG IL carries CUDA execution-space information at byte offset +182 (relative to the entity node base). The bitfield encoding:

This bitfield is checked throughout the pipeline -- in cross-space call validation, device/host code separation, the keep-in-IL predicate, and backend stub generation.

Lambda Wrapper Template Injection

CUDA extended lambdas (__device__ and __host__ __device__ lambdas) cannot be passed directly across the host/device boundary. cudafe++ solves this by injecting a library of template wrapper structs into the compilation at backend time. The master emitter sub_6BCC20 (nv_emit_lambda_preamble) generates all __nv_* templates in a single function call, driven by two 1024-bit bitmasks that record which capture counts were actually needed during parsing:

• unk_1286980: Device lambda capture counts (bit N = need __nv_dl_wrapper_t for N captures)
• unk_1286900: Host-device lambda capture counts (need __nv_hdl_wrapper_t for N captures)

Only the required specializations are emitted, keeping the generated code minimal.

CUDA Error Catalog

The binary contains 3,795 diagnostic messages in the EDG error table. Of these, 338 are CUDA-specific (error numbers in the 20000+ range and the 3500-3800 range). These cover:

• Execution space violations (calling __device__ from __host__ and vice versa)
• __global__ function constraints (no return value, no variadic args, no virtual)
• Lambda restrictions (35+ distinct error categories for extended lambda misuse)
• Attribute conflicts (__launch_bounds__ + __maxnreg__ mutual exclusion)
• RDC mode restrictions (user-defined copy constructors in kernel arguments)
• Architecture feature gates (feature X requires SM_YY or higher)

IL Entry Kind System

The EDG IL uses 85 defined entry kinds (0-84), each representing a distinct node type in the typed, scope-linked IL graph. Key node types include: routine (288 bytes, functions/methods), variable (232 bytes), type (176 bytes, 22 sub-kinds), expr_node (72 bytes, 36 sub-kinds), statement (80 bytes, 26 sub-kinds), and scope (288 bytes, 9 sub-kinds). All nodes live in a region-based arena allocator with 64 KB blocks. See IL Overview for the complete entry kind table.

CLI Flag Inventory

cudafe++ accepts 276 command-line flags parsed in sub_459630 (cmd_line.c). These control:

• Language mode and C++ standard version (__cplusplus value)
• Host compiler identity (MSVC, GCC, Clang) and version
• CUDA-specific modes: extended lambdas, RDC, JIT, architecture target
• Diagnostic suppression and promotion
• Include paths and macro definitions
• Output format and timing

Flags are passed from nvcc via the -Xcudafe forwarding mechanism. Many flags are undocumented EDG internals.

Wiki Structure

This wiki is organized into 10 sections covering the binary from top-level pipeline down to individual data structures.

Overview

• Function Map -- address-to-identity table for all 2,208 mapped functions
• Binary Layout -- segment map, memory regions, address space organization
• Methodology -- RE tools, approach, confidence scoring

Compilation Pipeline

The 8-stage pipeline from main() at 0x408950 through exit. Covers initialization, CLI parsing, EDG frontend invocation, 5-pass IL finalization, backend .int.c emission, and exit code mapping.

CUDA Execution Model

How cudafe++ handles __device__, __host__, and __global__ execution spaces. Device/host code separation, cross-space call validation, kernel stub generation, RDC (relocatable device code) mode, JIT mode, and SM architecture feature gating.

CUDA Attributes

The internal attribute system: __global__ function constraints, __launch_bounds__ / __cluster_dims__ / __block_size__ / __maxnreg__ validation, __grid_constant__ parameter handling, __managed__ variable support, and minor attributes (__nv_pure__, __nv_register_params__).

Lambda Transformations

Extended lambda support architecture: device lambda wrapper (__nv_dl_wrapper_t), host-device lambda wrapper (__nv_hdl_wrapper_t / __nv_hdl_create_wrapper_t), capture handling (field types, array wrappers for up to 8D), preamble injection (sub_6BCC20), and the 35+ lambda restriction error categories.

EDG Intermediate Language

The 85-entry-kind IL format: node allocation (region-based arena), tree walking (5 callback traversal), device code selection (keep-in-IL predicate), display (debug dump), and comparison/copy operations.

Host Output Generation

The .int.c file format, CUDA runtime boilerplate (__nv_managed_rt initialization, crt/host_runtime.h inclusion), host reference arrays (.nvHRKI/.nvHRDE/.nvHRCE ELF sections for device symbol registration), and CRC32-based module ID generation.

EDG Frontend Internals

The stock EDG 6.6 subsystems: lexer/tokenizer (357 token kinds), expression parser, declaration parser, overload resolution, template engine (instantiation worklist), CUDA-specific template restrictions, constexpr interpreter, Itanium ABI name mangling with CUDA extensions, and the type system (176-byte type node, 22 type kinds).

Error & Diagnostic System

The 3,795-entry diagnostic table, CUDA-specific error catalog (338 entries), format specifier system (%t/%s/%n/%sq/%p/%d), and SARIF output / pragma control.

Data Structures

Byte-level layouts for the core IL node types: entity node (execution/memory space at +182), scope entry (784 bytes), translation unit descriptor (424 bytes), type node (176 bytes, 22 kinds), and template instance record (128 bytes).

Configuration

CLI flag inventory (276 flags by category), EDG build configuration (compile-time constants baked into the binary), architecture detection (--nv_arch and SM version mapping), and experimental feature flags.

Reference

EDG source file map (52 .c + 13 .h), global variable index, token kind table (357 types), full error message catalog, and virtual override mismatch matrix.

Navigating This Wiki

If you want to understand the compilation pipeline: Start with Pipeline Overview, then follow the stage-by-stage links.

If you want to understand CUDA-specific behavior: Start with the CUDA Execution Model section. The execution spaces page explains the fundamental bitfield encoding that everything else depends on.

If you want to understand lambda transformations: Start with the Lambda Transformations overview. Lambda support is the most complex NVIDIA addition and involves template injection, capture-count bitmasks, and 5 distinct wrapper template families.

If you want to understand the IL format: Start with IL Overview for the 85 entry kinds, then Keep-in-IL for how device code is selected.

If you want to look up a specific function: The Function Map provides address-to-identity mappings for all 2,208 identified functions. The EDG Source File Map shows which source file each address range belongs to.

Data Sources

This wiki is derived from:

• 6,202 Hex-Rays decompiled C pseudocode files -- one per function with recognizable control flow
• 6,342 x86-64 disassembly files -- full instruction-level coverage
• 9.5 MB strings database with cross-references to every function that uses each string
• 161 MB cross-reference database -- complete caller/callee and data-reference mappings
• 7.7 MB call graph in JSON and DOT format
• 6,483 control flow graphs with basic block boundaries
• 247 MB IDA Pro database (.i64)

All analysis was performed on the binary shipped with CUDA Toolkit 13.0, obtained from NVIDIA's public distribution channels.

Function Map

Every function in the cudafe++ binary that triggers an EDG assertion encodes three pieces of data in the assertion string: the source file path, the line number, and the enclosing function name. These strings survive in .rodata and cross-reference back to the compiled functions, providing a ground-truth mapping from binary address to EDG source file. This page catalogs that mapping for all 52 .c source files and 13 .h header files identified in the CUDA 13.0 build of cudafe++ (EDG 6.6).

The mapping was produced by extracting all string literals matching /dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/drivers/compiler/edg/EDG_6.6/src/*.c and *.h from the binary's .rodata section, then tracing their cross-references to determine which functions load each path. A function that references attribute.c in an assertion string was compiled from attribute.c. Functions that reference no source path at all (the "unmapped" pool) are either too small to contain assertions, are inlined from headers, or belong to the statically-linked C++ runtime.

Coverage Summary

The 2,906 unmapped functions in the EDG region include inlined header expansions (e.g., util.h vector/hash helpers, types.h type queries), small leaf functions below the assertion threshold, switch-table dispatch fragments, and functions from translation units compiled without assertions enabled (notably il_to_str.c display routines and parts of floating.c).

Binary Layout

The EDG .text region (0x403300--0x7E0000) has a three-part structure:

1. Assert stub region (0x403300--0x408B40): 235 small __noreturn functions, one per assertion site. Each encodes a source file path, line number, and function name, then calls sub_4F2930 (the internal error handler). These stubs are sorted by source file name -- the linker grouped them from all 52 .c files into one contiguous block. 200 stubs map to .c files; the remaining 35 are from .h files inlined into .c compilation units.

2. Constructor region (0x408B40--0x409350): 15 C++ static constructor functions (ctor_001 through ctor_015) that initialize global tables at program startup.

3. Main body region (0x409350--0x7DFFF0): The bulk of the compiler. Source files are laid out roughly in alphabetical order by filename, a consequence of the linker processing object files in directory-listing order. The alphabetical ordering holds across the entire range: attribute.c starts at 0x409350, class_decl.c at 0x419280, progressing through to types.c at 0x7A4940, modules.c at 0x7C0C60, and floating.c at 0x7D0EB0.

Source File Address Table

The table below lists all 52 .c source files sorted by their main body start address. "Total Funcs" counts all functions referencing the file (stubs + main body). "Stubs" counts assert stubs in 0x403300--0x408B40. "Main Funcs" counts functions in the main body region.

Totals: 5,338 cross-references across 52 .c files, resolving to 2,129 unique functions. With .h file references added, 2,209 unique functions are mapped.

Largest Source Files by Function Count

Header File Cross-References

Thirteen .h header files appear in assertion strings. These are headers that contain non-trivial inline functions or macros that expand to assertion-bearing code. When a function compiled from decls.c triggers an assertion whose __FILE__ is types.h, that assertion was inlined from types.h into the decls.c compilation unit.

Notable Header Patterns

util.h is the most widely-included header, with 124 cross-references (114 in main body) spanning nearly the entire EDG .text region from 0x430E10 to 0x7C2B10. It provides generic container templates (dynamic arrays, hash tables, sorted sets) used by every major subsystem. The EDG linker inlined these templates into each compilation unit, creating many small util.h-attributed functions scattered across the binary.

float_type.h is concentrated in a single 52 KB block at 0x7D1C90--0x7DEB90, immediately after floating.c. It contains 63 template instantiations for IEEE 754 floating-point type operations (comparison, conversion, arithmetic) for each target floating-point width. These templates were instantiated in the floating.c compilation unit.

walk_entry.h contributes 51 functions in the tight range 0x604170--0x618660, all within the il_walk.c region. These are the per-entry-kind callback dispatch functions generated by preprocessor macros in the IL walker header.

nv_transforms.h is NVIDIA-specific. Its 3 cross-references appear in class_decl.c (sub_432280 at 0x432280), cp_gen_be.c (sub_47ECC0 at 0x47ECC0), and src_seq.c (sub_719D20 at 0x719D20). These are the integration points where NVIDIA's CUDA transform hooks are called from standard EDG code paths -- class definition processing, backend code generation, and source sequence ordering.

NVIDIA-Specific Files

nv_transforms.c

The only NVIDIA-authored .c file in the EDG source tree. Despite having only 1 mapped function via __FILE__ (sub_6BE300 at 0x6BE300), the sweep analysis of the 0x6BAE70--0x6BE4A0 range identified approximately 40 functions compiled from this file. The discrepancy exists because nv_transforms.c uses NVIDIA's own assertion macros (not EDG's standard internal_error path), so most functions do not reference the EDG-style __FILE__ string.

Functions confirmed in the nv_transforms.c region:

Key infrastructure in this file:

• __nv_dl_wrapper_t<> / __nv_hdl_wrapper_t<> struct template generation
• Host reference array emission (.nvHRKE, .nvHRKI, .nvHRDE, .nvHRDI, .nvHRCE, .nvHRCI)
• Capture count bitmask tables: unk_1286980 (device) and unk_1286900 (host-device), 128 bytes each
• Lambda-to-closure entity mapping via hash table at qword_12868F0

nv_transforms.h

NVIDIA's hook header, #include-d from three EDG source files. It declares the functions that bridge standard EDG processing to NVIDIA's CUDA transform layer. The three inclusion sites represent the three points where EDG's standard C++ frontend cedes control to NVIDIA-specific logic:

1. class_decl.c (sub_432280 at 0x432280): Called during class definition processing to apply CUDA execution-space attributes to closure types and validate lambda capture constraints.

2. cp_gen_be.c (sub_47ECC0 at 0x47ECC0): Called during backend code generation to emit CUDA-specific output constructs (device stubs, host reference arrays, registration calls).

3. src_seq.c (sub_719D20 at 0x719D20): Called during source sequence processing to inject NVIDIA preamble declarations and wrapper type definitions into the correct position in the declaration order.

Unmapped Regions (Gap Analysis)

Several address ranges within the EDG .text region contain functions that could not be mapped to any source file via __FILE__ strings. The major gaps and their probable contents:

The largest unmapped gap within EDG code is the IL display region at 0x5E8300--0x5F7FD0 (87 KB). These functions were compiled from il_to_str.c but contain no assertions because the display/dump subsystem was built without assertion macros -- it is purely diagnostic code that formats IL trees to stdout.

The float_type.h block at 0x7D1C90--0x7DEB90 (52 KB) is technically mapped via .h cross-references but has no .c file attribution because the template instantiations carry only the header's __FILE__ path.

Alphabetical Ordering Observation

The files are laid out in the binary in rough alphabetical order, consistent with a build system that compiles object files in directory-listing order and a linker that processes them sequentially:

0x409350  attribute.c      (a)
0x419280  class_decl.c     (c)
0x44B250  cmd_line.c       (c)
0x461C20  const_ints.c     (c)
0x466F90  cp_gen_be.c      (c)
0x48A1B0  debug.c          (d)
0x48B3F0  decl_inits.c     (d)
0x4A1BF0  decl_spec.c      (d)
0x4B3970  declarator.c     (d)
0x4C0910  decls.c          (d)
0x4E9E70  disambig.c       (d)
0x4EDCD0  error.c          (e)
0x4F9870  expr.c           (e)
0x558720  exprutil.c       (e)
0x584CA0  extasm.c         (e)
0x585B10  fe_init.c        (f)
0x588D40  fe_wrapup.c      (f)
0x589550  float_pt.c       (f)
0x594B30  folding.c        (f)
0x5A51B0  func_def.c       (f)
0x5AD540  host_envir.c     (h)
0x5B28F0  il.c             (i)
0x5E0600  il_alloc.c       (i)
0x5F7FD0  il_to_str.c      (i)
0x603FE0  il_walk.c        (i)
0x620CE0  interpret.c      (i)
0x65EA50  layout.c         (l)
0x666720  lexical.c        (l)
0x68ACC0  literals.c       (l)
0x68FAB0  lookup.c         (l)
0x69C980  lower_name.c     (l)
0x6AB6E0  macro.c          (m)
0x6B6DD0  mem_manage.c     (m)
0x6BAE70  nv_transforms.c  (n)  [region start; mapped func at 0x6BE300]
0x6BE4A0  overload.c       (o)
0x6F2790  pch.c            (p)
0x6F61B0  pragma.c         (p)
0x6F9B00  preproc.c        (p)
0x6FE160  scope_stk.c      (s)
0x710F10  src_seq.c        (s)
0x719300  statements.c     (s)
0x726F20  symbol_ref.c     (s)
0x72D950  symbol_tbl.c     (s)
0x74C690  sys_predef.c     (s)
0x7525F0  target.c         (t)
0x7530C0  templates.c      (t)
0x796BA0  trans_copy.c     (t)
0x796E60  trans_corresp.c  (t)
0x7A3BB0  trans_unit.c     (t)
0x7A4940  types.c          (t)
0x7C0C60  modules.c        (m)  [breaks alphabetical order]
0x7D0EB0  floating.c       (f)  [breaks alphabetical order]

Two files break the alphabetical pattern: modules.c at 0x7C0C60 and floating.c at 0x7D0EB0. Both appear after types.c instead of in their expected positions (between mem_manage.c and nv_transforms.c for modules.c, between float_pt.c and folding.c for floating.c). This suggests these two files are compiled as separate translation units outside the main EDG source directory, or are added to the link line after the alphabetically-sorted EDG objects.

Data Source

All mappings were extracted from the binary's .rodata string table. The extraction command:

jq '[.[] | select(.value | test("/dvs/p4/.*\\.c$")) |
  {file: (.value | split("/") | last),
   xrefs: [.xrefs[].func] | length}
] | sort_by(.file)' cudafe++_strings.json

The full build path for every source file is:

/dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/drivers/compiler/edg/EDG_6.6/src/<filename>

Address ranges were verified against the 20 sweep reports (P1.01 through P1.20) produced during the binary analysis phase.

Binary Layout

cudafe++ ships as a single statically-linked, stripped ELF 64-bit x86-64 executable. Static linking pulls in the entirety of libstdc++ (locale facets, iostream, exception handling), Berkeley SoftFloat 3e (half/quad-precision arithmetic), and glibc CRT startup code. The resulting 8.5 MB binary has no external shared library dependencies -- it runs identically on any Linux x86-64 host regardless of installed C++ runtime version.

This page documents the complete segment and section layout, the internal organization of each major section, and the key data structures located within each region. All addresses are virtual addresses from the ELF load image.

ELF Header

Complete Section Table

Total virtual address space consumed: 0x12D73A8 - 0x400000 = 18.9 MB.

.text -- Executable Code (4.15 MB)

The .text section contains all 6,483 functions in the binary. It divides into four distinct regions, laid out contiguously by the linker:

0x403300                                                          0x829722
|-- assert stubs --|-- ctors --|---- EDG main body ----|-- C++ runtime ----|
0x403300    0x408B40  0x409350                  0x7DF400          0x829722
   34 KB      8 KB              3.61 MB                  304 KB

Assert Stub Region (0x403300 -- 0x408B40, 34 KB)

Contains 235 small __noreturn functions, each encoding a single assertion site. Every stub loads three string constants -- source file path, line number, and function name -- then calls sub_4F2930 (the internal_error handler in error.c). These stubs are called from the bodies of larger functions when an impossible condition is detected.

The linker groups all stubs from all 52 .c source files into this contiguous block, sorted approximately by source file name. Of the 235 stubs:

• 200 map to .c source files (e.g., attribute.c:10897 at 0x403300, cp_gen_be.c:22342 at 0x4036F6)
• 35 map to .h header files inlined into .c compilation units (e.g., types.h at 0x40345C)

Each stub is exactly 29 bytes: a lea for the file path, a mov for the line number, a lea for the function name, then a call to sub_4F2930.

Constructor Region (0x408B40 -- 0x409350, 8 KB)

Contains 9 C++ global constructor functions (ctor_001 through ctor_009) registered in the .ctors table. These run before main() via __libc_start_main's init callback at 0x829640. The constructors, in execution order:

Constructors 4--9 belong to statically-linked libstdc++. Only constructors 1--3 initialize EDG/NVIDIA state.

EDG Main Body (0x409350 -- 0x7DF400, 3.61 MB)

The core of the compiler. Contains 5,115 functions compiled from 52 EDG .c source files plus 3 NVIDIA-specific source files. Functions are laid out in approximate alphabetical order by source file name -- the linker processed object files in directory-listing order:

0x409350   attribute.c     (170 functions)
0x419280   class_decl.c    (264 functions)
0x44B250   cmd_line.c      (43 functions)
0x461C20   const_ints.c    (3 functions)
0x466F90   cp_gen_be.c     (201 functions)
  ...
0x6BE300   nv_transforms.c (1 mapped function, NVIDIA)
0x6BE4A0   overload.c      (281 functions)
  ...
0x7A4940   types.c
0x7C0C60   modules.c
0x7D0EB0   floating.c
  ~0x7DF400  end of EDG code

The 52 source files break down by subsystem:

See Function Map for the complete address-to-source-file table.

C++ Runtime Region (0x7DF400 -- 0x829722, 304 KB)

Statically-linked library code with no EDG source attribution. Contains approximately 900 functions from three libraries:

Berkeley SoftFloat 3e (0x7E0D30 -- 0x7E4150, ~80 functions). IEEE 754 arithmetic for half-precision (float16), extended precision (float80), and quad-precision (float128). Operations: add, sub, mul, div, sqrt, comparisons, int/float conversions. Global state at 12D4820 (exception flags) and 12D4821 (rounding mode). Used by the EDG floating.c subsystem for constant folding of non-native float types.

libstdc++ / libsupc++ (0x7E42E0 -- 0x829600, ~800 functions). The C++ runtime:

• operator new/operator delete with new-handler retry loop (0x7E42E0)
• Exception handling: __cxa_throw (0x823050), __cxa_begin_catch (0x822EB0), __cxa_allocate_exception (0x7E4750), std::terminate (0x8231A0)
• Emergency exception pool: 72,704-byte fallback allocator for OOM during exception handling (0x7E45C0)
• iostream initialization: ios_base::Init constructor/destructor (0x7E5650/0x7E5F20) setting up cout/cin/cerr + wide variants
• Full locale system: 600+ functions implementing ctype, num_get, num_put, numpunct, collate, time_get/put, money_get/put, moneypunct, messages, and codecvt facets for both char and wchar_t

CUDA-aware name demangler (at 0x7CABB0, technically in the EDG tail region). NVIDIA's custom Itanium ABI demangler with extensions for CUDA lambda wrapper templates. Recognizes mangled prefixes: "Unvdl" for __nv_dl_wrapper_t<>, "Unvdtl" for __nv_dl_wrapper_t<> with trailing return, and "Unvhdl" for __nv_hdl_wrapper_t<>.

CRT startup (0x40918C and 0x829640 -- 0x829722). _start at 0x40918C calls __libc_start_main(main@0x408950, init@0x829640, fini@0x8296D0). The .fini_array processor at 0x8296E0 iterates backwards through function pointers at off_D428A0.

.rodata -- Read-Only Data (2.48 MB)

The .rodata section at 0x829740 -- 0xAA3FA3 holds all constant data: string literals, jump tables, error message templates, IL metadata tables, and format strings. Major structures:

Error Message Table (off_88FAA0)

The EDG diagnostic system's message template table. An array of 3,795 const char* pointers, indexed by error code 0--3794:

off_88FAA0[0]    = ""                           // error 0: unused
off_88FAA0[1]    = "last line of file ends ..."  // error 1
  ...
off_88FAA0[3794] = "..."                         // error 3794

Each pointer references a NUL-terminated format string elsewhere in .rodata containing % fill-in specifiers (%t = type, %s = string, %n = name, %sq = quoted string, %p = position, %d = decimal). Error codes above 3456 are CUDA-specific (338 entries covering execution space violations, lambda restrictions, architecture feature gates). See Diagnostic Overview.

IL Entry Kind Name Table (off_E6DD80)

Maps the 85 entry_kind enum values (0--84) to human-readable strings. Used by the IL display subsystem (il_to_str.c) for debug output:

off_E6DD80[0]  = "scope"
off_E6DD80[6]  = "type"
off_E6DD80[11] = "routine"
off_E6DD80[23] = "variable"
  ...
off_E6DD80[84] = "last"       // sentinel

The il_one_time_init function (sub_5CF7F0) validates at startup that this table ends with the "last" sentinel, catching version mismatches between the table and the enum.

EDG Source File Path Strings

Approximately 65 string literals of the form /dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/drivers/compiler/edg/EDG_6.6/src/<file>.<ext>. These are __FILE__ expansions embedded in assertion macros. Each is referenced by the corresponding assert stub in the 0x403300 region.

Jump Tables

Switch-statement jump tables for the major dispatch functions. The largest are:

• Expression parser dispatch (~120 case targets)
• Declaration specifier dispatch (~80 case targets)
• IL walker entry-kind dispatch (~85 case targets)
• Backend code generation dispatch (~90 case targets)

Format Strings

Printf-style format strings for the .int.c backend emitter. These include CUDA runtime boilerplate templates ("#include \"crt/host_runtime.h\"", "static __nv_managed_rt ...", "void __device_stub__...") and IL display format strings.

.data -- Initialized Globals (1.22 MB)

The .data section at 0xD46480 -- 0xE7EFF0 holds all initialized global variables. Major structures, ordered by address:

Attribute Descriptor Table (off_D46820)

The master attribute dispatch table, starting at 0xD46820 and extending to approximately 0xD47A60. Each entry is 32 bytes and describes one EDG/CUDA attribute kind: kind code (1 byte), flags (1 byte), name string pointer, validation function pointer, and application function pointer. See Attribute System Overview.

Diagnostic Fill-in Tables (off_D481E0)

Named-label fill-in descriptors for the diagnostic system. Maps fill-in label strings to format specifier dispatch codes. Located at 0xD481E0.

Keyword Tables

The EDG keyword registration system stores keyword-to-token-ID mappings. Initialized during fe_translation_unit_init (sub_5863A0) with 200+ C/C++ keywords (from auto through co_yield), 60+ type trait intrinsics (__is_class, __has_trivial_copy, etc.), and CUDA extension keywords (__device__, __global__, __shared__, __constant__, __managed__, __launch_bounds__, __grid_constant__).

Error Severity Override Table

Maps error codes to their overridden severity levels. Populated by --diag_suppress, --diag_warning, --diag_error CLI flags.

libstdc++ Vtables (0xD428C0 -- 0xD45E00, in .data.rel.ro)

The .data.rel.ro section holds vtables for all statically-linked C++ classes. Key vtables:

Exception Handler Pointers (0xE7EExx)

Located at the tail of .data:

EDG Diagnostic List Head (0xE7FE40)

A 40-byte doubly-linked list structure at 0xE7FE40..0xE7FE68. Initialized by ctor_001 as an empty self-referencing sentinel (both forward and backward pointers point to the list head). Used to chain diagnostic records during compilation.

.bss -- Zero-Initialized Globals (4.34 MB)

The .bss section at 0xE7F000 -- 0x12D6F20 is the largest section by virtual size. It contains all zero-initialized global state for both the EDG compiler and the statically-linked runtime. The .bss occupies no space in the ELF file on disk -- it is allocated and zeroed by the OS loader.

The 4.34 MB .bss divides into three logical regions:

EDG Compiler State (0xE7F000 -- 0x1290000, ~4.1 MB)

The bulk of .bss holds the EDG frontend's global state. Major structures:

Scope stack and symbol tables (~1.5 MB). The EDG scope stack (scope_stk.c) maintains nested scope contexts during parsing. Each scope entry is 784 bytes. The scope stack globals, various hash tables for name lookup, and the associated symbol table arrays consume the largest contiguous blocks.

IL region tracking (~800 KB). Region indices, region-to-scope mappings (qword_126EB90), region memory tables (qword_126EC88), and IL entry list heads. The region counter at dword_126EC80 tracks active regions. Each function definition creates a new region.

Translation unit state (~400 KB). The TU descriptor itself is dynamically allocated (424 bytes), but the per-TU global variables -- source file table, include stack, macro state, conditional compilation depth -- live in .bss. sub_7A4860 (reset_tu_state) zeroes these between compilations.

Parser state (~600 KB). Token lookahead buffers, declaration nesting depth, template argument stacks, expression evaluation context. The lexer maintains character classification tables and identifier hash buckets.

Error and diagnostic state (~200 KB). Error count (qword_126ED90), warning count (qword_126ED98), error limit (qword_126ED60), diagnostic suppression bitmaps, and the stream state table at 126ED80..126EDE0 (13 qwords including the stderr FILE* at qword_126EDF0).

Configuration flags (~100 KB). The 0x106xxxx region contains hundreds of dword flags set by CLI parsing and used throughout compilation. Examples:

Lambda capture bitmasks (~256 bytes). Two 1024-bit bitmasks recording which lambda capture counts were used during parsing:

Bit N set means a lambda with N captures was encountered, triggering emission of the corresponding __nv_dl_wrapper_t or __nv_hdl_wrapper_t specialization in the backend.

IL walker callbacks (5 function pointers at qword_126FB68..126FB88). The five IL tree-walk callback slots: entry filter, entry replace, pre-walk check, string callback, and entry callback. Swapped in and out by different IL traversal passes.

libstdc++ Runtime State (0x1290000 -- 0x12D6F20, ~280 KB)

SoftFloat globals (16 bytes). Exception flags at byte_12D4820, rounding mode at byte_12D4821.

Emergency exception pool (24 bytes of metadata). Free-list head (qword_12D4868), base address (qword_12D4870), capacity (qword_12D4878 = 72,704 bytes). The pool itself is heap-allocated at startup by ctor_004.

Locale system (~2 KB). The "C" locale singleton (unk_12D5E60), global locale impl pointer (qword_12D5E70), classic locale impl pointer (qword_12D5E78), character classification tables (12D5BE0..12D5D50), locale ID counter (dword_12D5E58), and pthread_once control variables.

iostream objects (~2 KB). The six standard stream objects and their backing file buffers:

Each stream object is backed by a basic_filebuf at a known offset (e.g., cout's filebuf at 0x12D67E0).

Demangler caches (40 bytes). Template argument cache at qword_12C7B40/12C7B48/12C7B50 (capacity/count/buffer pointer, grows by 500 entries via realloc). Block-scope suppress flag at dword_12C6A24.

EDG internal lists (7 x 48 bytes). Seven doubly-linked list structures at 12868C0..1286780 initialized by ctor_003. Serve as symbol/scope/type caches with destructor sub_6BD820.

Thread-Local Storage (0x12D6F20 -- 0x12D6F38, 24 bytes)

The .tls section holds exactly 24 bytes of thread-local data. This is the __cxa_eh_globals structure (accessed via __readfsqword(0) - 16):

struct __cxa_eh_globals {
    void     *caught_exception_stack;   // +0x00: linked list of caught exceptions
    uint32_t  uncaughtExceptions;       // +0x08: count of in-flight exceptions
};

Despite cudafe++ being single-threaded, the TLS infrastructure exists because libstdc++ exception handling unconditionally uses TLS offsets compiled into the static library.

.ctors / .dtors -- Constructor/Destructor Tables

The .ctors section at 0xD42858 is 88 bytes: a -1 sentinel (8 bytes), 9 constructor function pointers (72 bytes), and a 0 terminator (8 bytes). The 9 constructors are ctor_001 through ctor_009 documented above.

The .dtors section at 0xD428B0 is 16 bytes: a -1 sentinel and a 0 terminator. No destructors are registered -- all cleanup is done via __cxa_atexit handlers registered during construction.

.eh_frame / .gcc_except_table -- Exception Handling

The .eh_frame section (582 KB) contains DWARF Call Frame Information (CFI) records for stack unwinding during C++ exception propagation. The .gcc_except_table section (13.2 KB) contains GCC Language-Specific Data Area (LSDA) records that map program counters to catch handlers and cleanup functions.

The .eh_frame_hdr section (48.9 KB) is a binary search index into .eh_frame, enabling O(log n) lookup of unwind information by instruction pointer during exception throw.

These sections exist because libstdc++ exception handling requires them. cudafe++ itself rarely throws exceptions -- the EDG frontend uses longjmp-based error recovery. However, the statically-linked libstdc++ code (particularly operator new and locale initialization) uses C++ exceptions internally.

.plt / .got.plt -- PLT Stubs

The .plt section (2.2 KB, 141 entries) and .got.plt (1.1 KB) implement lazy binding for the 141 libc functions that cudafe++ imports despite static linking. These are glibc internal symbols resolved at load time. The PLT stubs are the standard x86-64 two-instruction pattern: indirect jump through GOT, then fallback to the dynamic linker (which never executes since the binary is statically linked -- the GOT is pre-resolved by the static linker).

Static Libraries Linked

The binary statically links four library components:

No shared libraries are loaded at runtime. The binary is fully self-contained.

Virtual Address Space Map

0x400000 +-----------------------+
         | ELF headers           |  10.5 KB
0x402A18 | .init                 |  24 B
0x402A30 | .plt                  |  2.2 KB
0x403300 | .text                 |  4.15 MB
         |   assert stubs        |    34 KB    (0x403300 - 0x408B40)
         |   constructors        |    8 KB     (0x408B40 - 0x409350)
         |   EDG main body       |    3.61 MB  (0x409350 - 0x7DF400)
         |   C++ runtime         |    304 KB   (0x7DF400 - 0x829722)
0x829722 | padding               |
0x829724 | .fini                 |  14 B
0x829740 | .rodata               |  2.48 MB
         |   error table         |    30 KB    (off_88FAA0)
         |   string literals     |    ~2 MB
         |   IL kind names       |    <1 KB    (off_E6DD80)
         |   jump tables         |    ~400 KB
0xAA3FA3 | .eh_frame_hdr         |  48.9 KB
         |       [gap]           |
0xCB1210 | .eh_frame             |  568 KB
0xD3F398 | .gcc_except_table     |  13.2 KB
0xD42858 | .ctors                |  88 B
0xD428B0 | .dtors                |  16 B
0xD428C0 | .data.rel.ro          |  13.3 KB   (vtables)
0xD45E00 |       [padding/GOT]   |
0xD46480 | .data                 |  1.22 MB
         |   attribute table     |    ~5 KB    (off_D46820)
         |   keyword tables      |    variable
         |   handler pointers    |    (at 0xE7EExx)
         |   diagnostic list     |    (at 0xE7FE40)
0xE7EFF0 |       [padding]       |
0xE7F000 | .bss                  |  4.34 MB
         |   EDG compiler state  |    ~4.1 MB  (0xE7F000 - 0x1290000)
         |   libstdc++ state     |    ~280 KB  (0x1290000 - 0x12D6F20)
0x12D6F20| .tls                  |  24 B
0x12D6F38| extern                |  1.1 KB
0x12D73A8+-----------------------+

Key Observations

The .bss dominates. At 4.34 MB, the .bss is the largest section -- larger than .text. This reflects the EDG frontend's design: hundreds of global variables hold parser state, scope stacks, symbol tables, and IL region metadata. A reimplementation should strongly consider replacing these globals with a context struct passed through the call chain.

Static linking adds 304 KB of dead-weight code. The C++ runtime region (0x7DF400 -- 0x829722) contains 900 functions, the majority of which (600+ locale facet methods) are never called by cudafe++. The locale system is pulled in transitively through iostream initialization. A reimplementation that avoids std::cout/std::cerr could eliminate this entirely.

The EDG code is tightly packed. The 3.61 MB EDG main body has almost no inter-function padding. Functions from the same source file are contiguous, and the alphabetical ordering by filename is consistent across the entire range. This makes address-to-source-file attribution reliable.

The binary is position-dependent. No PIE (Position-Independent Executable) flag is set. All code references use absolute addressing. The .got is minimal (56 bytes / 7 entries) -- almost all data references are direct.

Methodology

This page documents the reverse engineering methodology used to produce every page in this wiki. The goal is full transparency: a reader should be able to reproduce any finding by following the same techniques against the same binary. Every claim in the wiki traces back to one of four evidence categories (CONFIRMED, HIGH, MEDIUM, LOW), and this page defines exactly what each level means, what tools produced the raw data, and how that data was refined into the structured documentation that follows.

Toolchain

The binary was loaded into IDA Pro 9.0 with default x86-64 analysis settings. IDA's auto-analysis resolved all code/data boundaries, generated function boundaries for 6,483 functions, and identified 52,489 string literals. The Hex-Rays decompiler was invoked on all 6,483 functions; the IDAPython extraction log reports 6,343 successful decompilations (the remaining 140 failures are exception personality routines, SoftFloat leaf functions, and tiny thunks where Hex-Rays cannot reconstruct a valid C AST). However, due to function-name collisions in the output filenames (multiple sub_XXXXXX entries mapping to the same sanitized name after / replacement), the actual decompiled output directory contains 6,202 unique .c files -- the number used throughout this wiki.

Extraction Script

All raw data was exported from the IDA database in a single automated pass using analyze_cudafe++.py, an IDAPython script that runs inside IDA's scripting environment. The script produces 12 output artifacts:

Script Architecture

The script is structured as a main() function that calls idaapi.auto_wait() to block until IDA's auto-analysis completes, then executes 12 extraction passes in a fixed order. Output is written to four directories: the root output directory (JSON databases), graphs/ (per-function CFGs), disasm/ (per-function disassembly), and decompiled/ (per-function pseudocode). Directories are created if they do not exist.

The 12 passes, in execution order:

1. export_all_strings() -- Enumerates idautils.Strings(), then for each string walks XrefsTo(string_ea) to record every function that references it. Each string entry captures the address, string value, string type code, and a list of xref records ({from_addr, func_name, xref_type}). This is the foundation for source attribution (see below).

2. export_all_functions() -- For each function in idautils.Functions(), records start/end address, size, instruction count (via idc.is_code() on each head), library flag (FUNC_LIB), thunk flag (FUNC_THUNK), and builds caller/callee lists. Callers are found via XrefsTo(func_start); callees via XrefsFrom(head) filtered to call-type xrefs (fl_CN = type 17, fl_CF = type 19).

3. export_imports() -- Enumerates all imported modules via idaapi.get_import_module_qty() and idaapi.enum_import_names(). Records module name, symbol name, address, and ordinal for each of the 142 glibc imports.

4. export_segments() -- Iterates idautils.Segments() to record each ELF section's name, start/end address, size, type code, and permission bits.

5. export_xrefs() -- Full enumeration of all cross-references from every instruction head in every function. For each xref, records source address, source function, target address, target function (if any), and xref type code. Produces the 1,243,258-record xref table. The six xref type codes in the output:

   Type	Code	Count	Meaning
   dr_O	1	29,631	Data offset reference
   dr_W	2	11,488	Data write reference
   dr_R	3	42,364	Data read reference
   fl_CN	17	67,756	Code near call
   fl_CF	19	189,364	Code far/ordinary flow
   fl_JN	21	902,655	Code near jump (including fall-through)

6. export_comments() -- Walks every instruction head in the database via idautils.Heads(), extracting both regular comments (idc.get_cmt(ea, 0)) and repeatable comments (idc.get_cmt(ea, 1)).

7. export_names() -- Iterates idautils.Names() to export all named locations (function names, data labels, IDA auto-generated names).

8. extract_rodata() -- Reads the raw bytes of the .rodata segment via ida_bytes.get_bytes() and writes them to a binary file. Used for offline string scanning and jump table analysis.

9. export_callgraph() -- Builds the 67,756-edge call graph by iterating every function and scanning its instruction heads for outgoing call xrefs (fl_CN, fl_CF). Output in both JSON (array of {from, from_addr, to, to_addr} edge records) and Graphviz DOT format (67,759 lines).

10. export_complete_disassembly() -- Per-function disassembly files. For each function, iterates all instruction heads within the function's address range, generating hex byte dumps alongside disassembly text via idc.generate_disasm_line(). Each file includes a header with function name, address range, and byte size.

11. export_function_graphs() -- Per-function control flow graphs via idaapi.FlowChart(). For each basic block: block ID, start/end address, size, and full instruction listing. Block-to-block edges (fall-through and branch targets) are extracted via block.succs(). Output as both JSON (structured blocks + edges) and DOT (for Graphviz visualization).

12. export_decompilation() -- Calls idaapi.init_hexrays_plugin() to initialize the Hex-Rays decompiler, then iterates all functions and calls idaapi.decompile(func_ea). On success, the pseudocode string (str(cfunc)) is written to a .c file with a header comment containing the function name and address. Failures are silently caught via a bare except Exception and skipped.

The script is invoked via IDA's headless batch mode or interactive scripting console. It does not call qexit() at the end, allowing the IDA database to remain open for further interactive analysis after extraction. Total extraction time is approximately 30-45 minutes on a workstation-class machine, dominated by the 6,483 decompilation calls in pass 12.

Source Attribution Technique

The single most powerful technique in this analysis is source attribution via __FILE__ strings. The EDG C++ frontend uses C-style assertions throughout its codebase. When an assertion fires, the handler receives the source file path, line number, and function name as compile-time string constants embedded by the __FILE__, __LINE__, and __func__ macros. Because the binary is stripped (no .symtab), these assertion strings are the only surviving link to the original source tree.

The Assert Handler

The central assert handler is sub_4F2930, located in error.c. It is a __noreturn function that formats and emits an internal compiler error message, then terminates the process. A total of 2,139 functions in the binary call sub_4F2930, with 5,178 total call sites (many functions have multiple assertion points throughout their bodies).

The highest-density callers are the 235 assert stubs in the region 0x403300--0x408B40. Each stub is exactly 29 bytes: three register loads (source file path via lea rdi, line number via mov esi, function name via lea rdx) followed by a call to sub_4F2930:

sub_403300:         ; assert stub for is_aliasable (attribute.c:10897)
  lea  rdi, aAttributeC    ; "/dvs/p4/.../EDG_6.6/src/attribute.c"
  mov  esi, 10897           ; line number (integer, not string)
  lea  rdx, aIsAliasable   ; "is_aliasable"
  call sub_4F2930           ; internal_error(__FILE__, __LINE__, __func__)

Of the 235 stubs, 200 reference .c file paths and 35 reference .h file paths (inlined assertions from header files). The stubs are sorted approximately by source file name within the stub region -- the linker grouped them from all 52 .c compilation units into one contiguous block.

Beyond the dedicated stubs, 1,904 additional functions contain inline assertion checks: the lea rdi, <file_path> instruction appears within the function body at the assertion site, not in a separate stub. These inline assertions provide the same source-file attribution as the stubs.

The Attribution Chain

The attribution chain works in three steps:

1. String discovery. Extract all strings matching the EDG build path prefix /dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/drivers/compiler/edg/EDG_6.6/src/. This yields one string per source file, each cross-referenced by the assert stubs that load it.

2. Xref tracing. For each assert stub, follow XrefsTo() to find which main-body functions call it. A function at 0x40DFD0 that calls the attribute.c:5108 stub was compiled from attribute.c. This attributes the caller to the source file.

3. Range extension. Assert stubs are sparse -- not every function contains an assertion. Once a set of functions in a contiguous address range are attributed to the same source file, the entire range is assigned to that file. This works because the linker places all object code from a single .c file contiguously, and the files are arranged roughly alphabetically by filename.

This technique attributed 2,209 functions (34.1% of the binary) to specific source files. The remaining 4,274 functions fall into three categories: C++ runtime code (1,085 functions from libstdc++/glibc, identifiable by address range), PLT/init stubs (283 functions), and unmapped EDG functions (2,906 functions that contain no assertions and cannot be confidently attributed).

Build Path

The full build path embedded in the binary is:

/dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/drivers/compiler/edg/EDG_6.6/src/

This reveals the NVIDIA internal Perforce depot structure (/dvs/p4/), the release branch (r13.0), and the EDG version (EDG_6.6). It confirms the binary was built from EDG C++ Front End version 6.6, licensed from Edison Design Group.

Confidence Levels

Every identification in the raw sweep reports and wiki pages carries one of four confidence levels:

In practice, approximately 34% of functions are CONFIRMED (via assert strings), ~20% are HIGH (via distinctive strings or unique callgraph positions), ~25% are MEDIUM, and ~21% are LOW or unattributed.

Call Graph Analysis

The complete call graph contains 67,756 edges connecting the 6,483 functions. This graph is the primary tool for understanding system architecture -- which subsystems call which, where the hot paths are, and how NVIDIA's additions integrate with the EDG base.

Hub Identification

Hub functions -- those with exceptionally high in-degree (many callers) or out-degree (many callees) -- reveal the architectural spine of the compiler:

Graph Structure

The call graph exhibits a layered structure typical of compiler frontends:

1. Entry layer. main() at 0x408950 calls exactly 8 stage functions in sequence.
2. Stage layer. Each stage function (init, CLI, parse, wrapup, backend) fans out to dozens of subsystem entry points.
3. Core layer. The parser (expr.c, decls.c, statements.c) calls into the type system (types.c, exprutil.c), IL builder (il.c, il_alloc.c), and name lookup (lookup.c, scope_stk.c).
4. Leaf layer. Memory management (mem_manage.c), error reporting (error.c), and type queries form the bottom of the call hierarchy, referenced from almost every subsystem.

NVIDIA's nv_transforms.c sits as a lateral extension at the core layer: it is called from class_decl.c, cp_gen_be.c, and statements.c (via nv_transforms.h inlines), but does not itself call back into the EDG parser. This clean separation suggests NVIDIA modifies the EDG source minimally, preferring to hook into existing EDG extension points rather than fork the core.

String-Based Discovery

The binary contains 52,489 strings in .rodata. These strings are the second most important evidence source after the assertion paths. Major categories:

String Mining Techniques

Three string mining techniques are used throughout the analysis:

1. Error message tracing. CUDA-specific error messages (e.g., "calling a __host__ function from a __device__ function is not allowed") are grepped from the string table, their xrefs traced to the emitting function, and the emitting function's callers analyzed to understand the validation logic that triggers the error.

2. Keyword enumeration. The keyword initialization function (sub_5863A0) loads 200+ string constants in sequence. By reading the strings in load order, the complete CUDA keyword vocabulary is recovered -- including internal-only keywords not documented in the CUDA C++ Programming Guide.

3. Format string analysis. Format strings in the backend (cp_gen_be.c) reveal the exact syntax of .int.c output. A string like "static void __device_stub__%s(" tells us the precise naming convention for device stub wrapper functions.

Decompilation Quality

Hex-Rays produces readable pseudocode for the vast majority of functions, but several systematic limitations affect the analysis:

Control Flow Artifacts

Hex-Rays occasionally introduces control flow constructs that do not exist in the original source. The most prominent example is the while(1) loop in main() (sub_408950): the decompiler wraps the entire function body in an infinite loop because a setjmp-based error recovery mechanism creates a backward edge in the CFG. In reality, main() executes linearly and returns -- the while(1) is a decompiler artifact, not a real loop.

Similar artifacts appear in functions with complex switch statements (EDG uses computed gotos for performance), where Hex-Rays may produce nested if-else chains instead of the flat dispatch table the original code uses.

Lost Preprocessor Logic

The original EDG source makes heavy use of preprocessor conditionals (#if CUDA_SUPPORT, #ifdef FRONT_END_CPFE, etc.). The compiled binary contains only the taken branch -- the preprocessor evaluated all conditions at build time. This means the decompiled code shows the CUDA-enabled configuration only; any host-only or non-CUDA EDG behavior is invisible.

Similarly, C macros that wrap common patterns (assertion macros, IL access macros, type query macros) are fully expanded in the binary. The decompiled output shows the expanded form -- a sequence of struct field accesses and conditional jumps -- rather than the concise macro invocation the original source used.

Unnamed Variables

The binary is stripped. All local variable names are lost. Hex-Rays assigns synthetic names (v1, v2, a1, a2) based on register allocation and stack slot positions. Function parameters are named a1 through aN in declaration order. During analysis, meaningful names are sometimes manually applied in the IDA database, but most decompiled output uses the synthetic names.

Structure field accesses appear as byte-offset expressions (*((_BYTE *)a1 + 182)) rather than named fields (entity->execution_space). Reconstructing the structure layouts from these offset patterns is a core part of the analysis -- see the Entity Node Layout page for the most extensively reconstructed structure.

Decompilation Failures

The IDAPython extraction log reports 6,343 successful decompilations out of 6,487 attempts (140 failures). Due to filename collisions in the output directory (functions with identical sanitized names at different addresses overwrite each other), the actual output directory contains 6,202 unique .c files. The 281 "missing" files break down as:

The 140 true decompilation failures are concentrated in the C++ runtime region (0x7DF400--0x829722), particularly in the libstdc++ locale facet implementations (complex template instantiations with deeply nested virtual dispatch) and Berkeley SoftFloat 3e functions (pure arithmetic with non-standard calling conventions). For these functions, analysis relies on the raw disassembly output in disasm/ instead.

Phase 1: Address-Range Sweeps

The first phase of analysis consists of 20 address-range sweeps that collectively cover the entire .text section from 0x403000 to 0x82A000. Each sweep examines a contiguous address range of 128--256 KB, documenting every function within that range.

Sweep Index

The P1.11 sweep was subdivided into six sub-sweeps (11a through 11f) because the il.c region is dense and complex, containing the core IL node creation and manipulation functions that are referenced from nearly every other source file.

Sweep Report Format

Each sweep report follows a consistent format:

================================================================================
P1.XX SWEEP: Address range 0xNNNNNN - 0xMMMMMM
================================================================================
Range: 0xNNNNNN - 0xMMMMMM
Functions found: N
EDG source files:
  - file.c (assert stub range, main body range)
  ...

### 0xAAAAAA -- sub_AAAAAA (NN bytes / NN lines)
**Identity**: function_name (source_file.c:NNNN)
**Confidence**: CONFIRMED / HIGH / MEDIUM / LOW
**EDG Source**: source_file.c
**Notes**: Additional observations about behavior, callers, callees

Every function in the sweep range gets an entry. Functions are documented in address order. The identity field records the inferred function name and source location. The confidence field uses the four-level system defined above. Notes capture anything unusual -- unexpected callers, CUDA-specific behavior, undocumented error codes, or connections to other subsystems.

Phase 2: Targeted Deep Dives

After the Phase 1 sweep establishes the complete function map and identifies all source files, Phase 2 produces the detailed wiki pages. Each wiki page corresponds to one W-series work report that focuses on a specific subsystem or topic.

Deep Dive Methodology

Each W-series report follows a consistent process:

1. Scope definition. Identify the set of functions relevant to the topic. For example, W012 (Execution Spaces) requires the CUDA attribute application handlers in attribute.c, the execution space checking functions in nv_transforms.c, and the virtual override validator in class_decl.c.

2. Decompilation review. Read the full Hex-Rays pseudocode for every function in scope. For complex functions, also review the raw disassembly to catch decompiler artifacts.

3. String evidence collection. Grep the string table for all strings referenced by the in-scope functions. Error messages reveal validation rules; format strings reveal output patterns; keyword strings reveal accepted syntax.

4. Call graph traversal. Starting from the in-scope functions, walk callers and callees to understand the full data flow. Who calls apply_nv_global_attr? What does it call? How does data arrive and where does it go?

5. Struct layout reconstruction. When decompiled code accesses struct fields via byte offsets, reconstruct the field layout by collecting all access patterns across all functions that touch the same struct. Cross-validate offsets across multiple functions.

6. Pseudocode reconstruction. Translate the Hex-Rays output into readable C-like pseudocode with meaningful variable names, proper control flow, and comments explaining the logic. This reconstructed pseudocode appears in the wiki pages.

7. Cross-reference synthesis. Link findings to other wiki pages and W-series reports. Every page should situate itself within the overall architecture.

W-Series Report Index

As of this writing, 28 W-series reports have been produced, each backing one or more wiki pages:

Numerical Summary

Limitations and Caveats

What This Analysis Cannot Determine

1. Preprocessor-disabled code. Any EDG code behind #if 0, #ifndef CUDA_SUPPORT, or similar guards was compiled out. The binary reflects only the CUDA-enabled, Linux x86-64, EDG 6.6 configuration. Other EDG frontend features (e.g., Fortran support, Windows target, older C++ standards) are not present.

2. Inlined function boundaries. When the compiler inlines a function, its code merges with the caller. The binary may contain hundreds of inlined instances of small EDG utility functions (type queries, IL accessors) that are invisible as separate entities. The 6,483 function count represents only the non-inlined functions.

3. Original variable names. All local and most global variable names are lost. The wiki uses reconstructed names based on semantics (e.g., execution_space_byte for *((_BYTE *)entity + 182)), but these are analyst-assigned, not original.

4. Exact source line mapping. While assertion strings encode line numbers, these are the assertion site's line number, not the calling function's line number. The analyst can determine that is_aliasable in attribute.c has an assertion at line 10897, but cannot determine the start line of is_aliasable itself.

5. NVIDIA-internal documentation. Any design documents, code comments, commit messages, or internal wikis that informed the original development are unavailable. All behavioral descriptions in this wiki are inferred from the binary alone.

Reproducibility

Every finding in this wiki can be reproduced by:

1. Obtaining cudafe++ from CUDA Toolkit 13.0 (version string embedded in binary as the build path prefix r13.0).
2. Loading it into IDA Pro 9.0 (64-bit) with default x86-64 analysis settings. Wait for auto-analysis to complete (5-10 minutes).
3. Running analyze_cudafe++.py via File > Script File to extract all raw data (30-45 minutes).
4. Querying the exported JSON files with jq to trace cross-reference chains, string lookups, and callgraph paths.
5. Reading the decompiled .c files and raw .asm files for behavioral analysis.

No proprietary tools beyond IDA Pro + Hex-Rays are required. The analysis does not depend on NVIDIA source code access, NDA-protected documentation, or insider knowledge. Every claim is derived from the publicly distributed binary.

Pipeline Overview

cudafe++ is a source-to-source compiler. It reads a .cu file, parses it as C++ with CUDA extensions using a modified EDG 6.6 frontend, then emits a transformed .int.c file where device code is suppressed and host-side stubs replace kernel launch sites. The entire binary is a single-threaded, single-pass-per-stage pipeline controlled from main() at 0x408950.

Pipeline Diagram

  input.cu
     |
     v
 [1] fe_pre_init          sub_585D60   fe_init.c
     9 subsystem pre-initializers
     |
     v
     * sub_5AF350(v7) ---- capture "Total compilation time" start
     |
     v
 [2] proc_command_line     sub_459630   cmd_line.c
     276 CLI flags parsed, mode selection
     |
     v
 [3] fe_one_time_init      sub_585DB0   fe_init.c
     38 subsystem initializers + keyword registration
     |--- fe_init_part_1 (sub_585EE0): per-unit inits, output file open
     |--- keyword_init + fe_translation_unit_init (sub_5863A0)
     |
     v
     * sub_5AF350(v8) ---- capture "Front end time" start
     |
     v
 [4] reset_tu_state        sub_7A4860   trans_unit.c
     Zero all TU globals
     |
     v
 [5] process_trans_unit    sub_7A40A0   trans_unit.c
     Allocate 424-byte TU descriptor, parse source,
     build EDG IL tree, CUDA attribute propagation
     |
     v
 [6] fe_wrapup             sub_588F90   fe_wrapup.c
     5-pass IL finalization: needed-flags, keep-in-IL marking,
     dead entity elimination, scope cleanup
     |
     v
     * sub_5AF350(v9) ---- capture "Front end time" end
     * sub_5AF390("Front end time", v8, v9)
     |
     v
     * sub_5AF350(v10) --- capture "Back end time" start
     |
     v
 [7] Backend entry         sub_489000   cp_gen_be.c
     Walk source sequence, emit .int.c, device stubs,
     lambda wrappers, registration tables
     |
     v
     * sub_5AF350(v11) --- capture "Back end time" end
     * sub_5AF390("Back end time", v10, v11)
     |
     v
     * sub_5AF350(v12) --- capture "Total compilation time" end
     * sub_5AF390("Total compilation time", v7, v12)
     |
     v
 [8] exit_with_status      sub_5AF1D0   host_envir.c
     Map internal status to exit code, terminate

     |----- "Front end time" covers stages 4-6 ----------|
     |----- "Back end time" covers stage 7 ---------------|
     |----- "Total compilation time" covers stages 2-8 ---|

Call Hierarchy from main()

The decompiled main() at 0x408950 calls the pipeline stages in this exact order:

void main(int argc, char **argv, char **envp)
{
    sub_585D60(argc, argv, envp);      // [1] fe_pre_init
    sub_5AF350(v7);                     //     capture_time (total start)
    sub_459630(argc, argv);             // [2] proc_command_line
    // [stack limit adjustment via setrlimit]
    sub_585DB0();                       // [3] fe_one_time_init
    if (dword_106C0A4)
        sub_5AF350(v8);                 //     capture_time (frontend start)
    sub_7A4860();                       // [4] reset_tu_state
    sub_7A40A0(qword_126EEE0);         // [5] process_translation_unit
    sub_588F90(v5, 1);                  // [6] fe_wrapup
    if (dword_106C0A4) {
        sub_5AF350(v9);
        sub_5AF390("Front end time", v8, v9);
    }
    // --- error-recovery re-compilation loop ---
    if (qword_126ED90) {               //     errors present?
        dword_106C254 = 1;             //     skip backend
    }
    while (1) {
        sub_6B8B20(0);                  //     reset file state
        sub_589530();                   //     write signoff + cleanup
        // exit code computation
        if (dword_106C0A4)
            sub_5AF390("Total compilation time", ...);
        sub_5AF1D0(exit_code);          // [8] exit
        // --- if dword_106C254 == 0, backend runs ---
        if (!dword_106C254) {
            if (dword_106C0A4)
                sub_5AF350(v10);        //     capture_time (backend start)
            sub_489000();               // [7] process_file_scope_entities
            if (dword_106C0A4) {
                sub_5AF350(v11);
                sub_5AF390("Back end time", v10, v11);
            }
        }
    }
}

The while(1) loop with sub_5AF1D0 (which calls exit() / abort()) never actually iterates -- the call to sub_5AF1D0 is __noreturn. The compiler just arranged the basic blocks this way: the backend stage at label LABEL_16 falls through from a goto at the top of the loop when dword_106C254 == 0 (no errors).

Stage Details

Stage 1: fe_pre_init -- sub_585D60 (0x585D60)

Source: fe_init.c

Performs absolute minimum initialization before anything else can run. Called with the raw argc, argv, envp from the OS.

Sets dword_126C5E4 = -1 (current scope index = "none") and dword_126C5C8 = -1 (secondary scope index = "none").

Data flow: No input beyond process args. Output: global state zeroed and ready for CLI parsing.

Stage 2: proc_command_line -- sub_459630 (0x459630)

Source: cmd_line.c (4105 decompiled lines)

Parses all 276 CLI flags. Populates global configuration variables that control every subsequent stage. Key outputs:

The parser builds four hash tables for macro defines (qword_106C248), include paths (qword_106C240), and system includes (qword_106C238, qword_106C228). It also suppresses a default set of diagnostic numbers (1257, 1373, 1374, 1375, 1633, 2330, 111, 185, 175).

Data flow: Input: argv. Output: ~150+ global configuration variables populated.

Stage 3: fe_one_time_init -- sub_585DB0 (0x585DB0)

Source: fe_init.c

The heaviest initialization stage. Calls 38 subsystem initializers in dependency order, then validates the function pointer dispatch table (a sentinel check: off_D560C0 must equal the address of nullsub_6). After validation, calls sub_585EE0 (fe_init_part_1) which:

1. Records compilation timestamp via time()/ctime() into byte_106B5C0
2. Runs 26 per-compilation-unit initializers
3. Opens the output file (qword_106C280 = stdout or file)
4. Writes the output file header via sub_5AEDB0
5. Calls the keyword registration function sub_5863A0 which registers 200+ C/C++ keywords plus NVIDIA CUDA-specific type traits (__nv_is_extended_device_lambda_closure_type, etc.)

38 subsystem initializers (in call order):

Data flow: Input: populated config globals. Output: all subsystems initialized, keyword table built, output file open.

Stage 4: reset_tu_state -- sub_7A4860 (0x7A4860)

Source: trans_unit.c

Zeroes all translation unit tracking globals to prepare for processing:

qword_106BA10 = 0;   // current_translation_unit
qword_106B9F0 = 0;   // primary_translation_unit
qword_12C7A90 = 0;   // tu_chain_tail
dword_106B9F8 = 0;   // has_module_info
qword_106BA18 = 0;   // tu_stack_top
dword_106B9E8 = 0;   // tu_stack_depth

Data flow: No input. Output: TU state clean-slated.

Stage 5: process_translation_unit -- sub_7A40A0 (0x7A40A0)

Source: trans_unit.c

The main frontend workhorse. This single call parses the entire .cu source file into the EDG intermediate language. Workflow:

1. Debug trace: "Processing translation unit %s"
2. Clean up any previous TU state (sub_7A3A50)
3. Reset error state (sub_5EAEC0)
4. Allocate 424-byte TU descriptor via sub_6BA0D0
5. Initialize TU scope state (offsets 24..192 via sub_7046E0)
6. Set as primary TU (qword_106B9F0) if first
7. Link into TU chain
8. Call sub_586240 -- parse the source file (this enters the EDG parser, which handles all of C++ plus CUDA extensions: __device__, __host__, __global__, __shared__, __managed__, etc.)
9. Depending on mode:
   • Module compilation: sub_6FDDF0
   • Standard compilation: sub_6F4AD0 (header-unit) + sub_4E8A60 (standard)
10. Post-processing: sub_588E90 (translation_unit_wrapup -- scope closure, template wrapup, IL output)
11. Debug trace: "Done processing translation unit %s"

At the end of this stage, the EDG IL tree is fully built. Every declaration, type, expression, and statement from the source has been parsed into IL nodes. CUDA execution-space attributes (__device__, __host__, __global__) have been recorded on entity nodes at byte offset +182 (bit 6 = device/global, bits 4-5 = execution space).

Data flow: Input: source filename from qword_126EEE0. Output: complete EDG IL tree anchored at qword_106BA10 (TU descriptor), source sequence list at *(qword_106BA10 + 8).

Stage 6: fe_wrapup -- sub_588F90 (0x588F90)

Source: fe_wrapup.c

Five-pass finalization over all translation units. Each pass iterates the TU chain (qword_106B9F0). Passes 2-4 are per-TU error-gated (skip TUs with qword_126ED90 != 0); passes 1 and 5 run unconditionally.

Between Pass 1 and Pass 2, if no errors have occurred, sub_796C00 runs cross-TU entity marking.

Post-pass operations:

• Cross-TU consistency (sub_796BA0, error-gated)
• Scope renumbering (sub_707480 double-loop)
• Template validation (sub_765480)
• File index cleanup (sub_6B8B20 for indices 2..dword_126EC80)
• Output flush + close three output files (IDs 1513, 1514, 1515)
• Memory statistics: sums 10 space_used() callbacks
• State teardown

Data flow: Input: fully built IL tree. Output: finalized IL with dead entities eliminated and device-needed entities marked. The source sequence list (qword_1065748) is the ordered list of top-level declarations the backend will walk.

Stage 7: Backend Code Generation -- sub_489000 (0x489000)

Source: cp_gen_be.c (723 decompiled lines, the largest single function in the backend)

This is the host-side C++ code generator. It walks the EDG source sequence and emits the .int.c file that the host compiler (gcc/cl.exe/clang) will compile. The backend is gated by dword_106C254: if set to 1 (errors occurred), stage 7 is skipped entirely.

Initialization:

1. Zeros output state: dword_1065834 (indent level), stream handle, counters
2. Clears four 512KB hash tables (memset 0x7FFE0 bytes each)
3. Sets up gen_be_info callback table (xmmword_1065760..10657B0)
4. Creates output file: <input>.int.c (or stdout for "-")

Boilerplate emission:

• #pragma GCC diagnostic push/pop blocks for suppressing host compiler warnings
• __nv_managed_rt initialization boilerplate (for __managed__ variables)
• Lambda type-trait macro definitions

Main processing loop:

• Walks qword_1065748 (global source sequence list)
• For each entry: dispatches to sub_47ECC0 (gen_template/process_source_sequence)
• Kind 57 entries are pragma interleavings (handled inline)

CUDA-specific transformations performed:

1. Device stub generation: For __global__ kernels, emit __wrapper__device_stub_<name>() forwarding, wrap original body in #if 0/#endif
2. Device-only suppression: Device-only declarations wrapped in #if 0/#endif
3. Lambda wrappers: __nv_dl_wrapper_t<> for device lambdas, __nv_hdl_create_wrapper_t<> for host-device lambdas
4. Runtime header injection: #include "crt/host_runtime.h" at first CUDA entity
5. Registration tables: sub_6BCF80 called 6 times for device/host/managed/constant combinations
6. Anonymous namespace: _NV_ANON_NAMESPACE macro for unique global symbols

Trailer:

• Empty-file guard: int __dummy_to_avoid_empty_file;
• Re-inclusion of original source via #include "<original_file>"
• #undef _NV_ANON_NAMESPACE

Data flow: Input: finalized source sequence from stage 6. Output: .int.c file on disk.

Stage 8: exit_with_status -- sub_5AF1D0 (0x5AF1D0)

Source: host_envir.c

Maps internal compilation status to process exit codes:

In SARIF mode (dword_106BBB8), text messages are suppressed but exit codes remain the same.

Key Global Variables Controlling Flow

Timing Regions

When dword_106C0A4 is set (via --timing or equivalent flag), three timing regions are printed:

Front end time                     12.34 (CPU)     15.67 (elapsed)
Back end time                       3.45 (CPU)      4.56 (elapsed)
Total compilation time             15.79 (CPU)     20.23 (elapsed)

Format string: "%-30s %10.2f (CPU) %10.2f (elapsed)\n"

The timing is implemented via sub_5AF350 (capture_time: records clock() as CPU milliseconds and time() as wall seconds) and sub_5AF390 (report_timing: computes deltas and prints).

Error Recovery Loop

The main() function contains a while(1) loop that appears to support re-compilation (the TU processing infrastructure has a dword_106BA08 "is_recompilation" flag and sub_7A40A0 checks an a2 recompilation parameter). In practice, for the standard CUDA compilation flow, this loop executes exactly once: sub_5AF1D0 is __noreturn and terminates the process.

The loop body:

1. sub_6B8B20(0) -- reset file state for the source file manager
2. sub_589530() -- write output signoff (sub_5AEE00) + close source manager (sub_6B8DE0)
3. Compute exit code from qword_126ED90 (errors) and qword_126ED88 (additional status)
4. Print total timing if enabled
5. Restore stack limit if it was raised
6. sub_5AF1D0(exit_code) -- terminate

Cross-References

• Entry Point & Initialization -- detailed breakdown of stages 1-3
• CLI Processing -- all 276 flags parsed in stage 2
• Frontend Invocation -- stage 5 (parse + IL build) in depth
• Frontend Wrapup -- 5-pass architecture of stage 6
• Backend Code Generation -- stage 7 (.int.c emission) in depth
• Timing & Exit -- stage 8 and exit code mapping
• Device/Host Separation -- how the backend filters device vs host code
• Kernel Stub Generation -- __wrapper__device_stub_ pattern
• Extended Lambda Overview -- lambda wrapper generation in backend
• .int.c File Format -- structure of the backend output

Entry Point & Initialization

main() at 0x408950 is a 488-byte __noreturn function that orchestrates the entire cudafe++ compilation pipeline. It takes the standard POSIX signature (int argc, char **argv, char **envp), performs two phases of subsystem initialization, optionally raises the process stack limit, then runs the frontend, backend, and exit sequence in a linearized loop that executes exactly once. The function has 22 direct callees (including getrlimit, setrlimit, and library calls) and never returns -- sub_5AF1D0 at the bottom of the loop calls exit() or abort().

Key Facts

Annotated Decompilation

void __noreturn main(int argc, char **argv, char **envp)
{
    rlim_t original_stack;
    bool stack_was_raised;
    uint8_t exit_code;
    struct rlimit rlimits;
    timestamp_t t_total_start, t_fe_start, t_fe_end, t_be_start, t_be_end, t_total_end;

    // --- Redirect diagnostic output to stderr ---
    s = stderr;                            // 0x126EDF0 alias
    qword_126EDF0 = stderr;                // diagnostic stream

    // === PHASE 1: Pre-initialization (9 subsystem calls) ===
    sub_585D60(argc, argv, envp);          // fe_pre_init

    // --- Capture total compilation start time ---
    sub_5AF350(&t_total_start);            // capture_time

    // === PHASE 2: Command-line parsing ===
    sub_459630(argc, argv);                // proc_command_line (276 flags)

    // === Stack limit adjustment ===
    if (dword_106C064                      // --modify-stack-limit (default: ON)
        && !getrlimit(RLIMIT_STACK, &rlimits))
    {
        original_stack = rlimits.rlim_cur;
        rlimits.rlim_cur = rlimits.rlim_max;  // raise to hard limit
        stack_was_raised = (setrlimit(RLIMIT_STACK, &rlimits) == 0);
    }

    // === PHASE 3: Heavy initialization (38 subsystem calls + validation) ===
    sub_585DB0();                          // fe_one_time_init
    //   └─ sub_585EE0()  fe_init_part_1  (33 per-unit inits, output file, keywords)

    if (dword_106C0A4)                     // --timing enabled?
        sub_5AF350(&t_fe_start);           // capture frontend start

    // === PHASE 4: Translation unit setup ===
    sub_7A4860();                          // reset_tu_state (zero 6 TU globals)

    // === PHASE 5: Frontend parse + IL build ===
    sub_7A40A0(qword_126EEE0);            // process_translation_unit

    // === PHASE 6: Frontend wrapup (5-pass IL finalization) ===
    sub_588F90(qword_126EEE0, 1);         // fe_wrapup

    if (dword_106C0A4) {
        sub_5AF350(&t_fe_end);
        sub_5AF390("Front end time", &t_fe_start, &t_fe_end);
    }

    // --- Error gate: skip backend if frontend had errors ---
    if (!qword_126ED90) goto backend;     // no errors → run backend
    dword_106C254 = 1;                    // skip-backend flag

    // === Linearized exit loop (executes once) ===
    while (1) {
        exit_code = 8;                    // default: warnings
        sub_6B8B20(0);                    // reset file state
        sub_589530();                     // write signoff + close source mgr

        if (!qword_126ED90)               // re-check after wrapup
            exit_code = qword_126ED88 ? 5 : 3;  // success codes

        if (dword_106C0A4) {
            sub_5AF350(&t_total_end);
            sub_5AF390("Total compilation time", &t_total_start, &t_total_end);
        }

        if (stack_was_raised) {           // restore original stack limit
            rlimits.rlim_cur = original_stack;
            setrlimit(RLIMIT_STACK, &rlimits);
        }

        sub_5AF1D0(exit_code);            // __noreturn: exit() or abort()

    backend:
        if (!dword_106C254) {             // backend not skipped
            if (dword_106C0A4)
                sub_5AF350(&t_be_start);
            sub_489000();                 // process_file_scope_entities (backend)
            if (dword_106C0A4) {
                sub_5AF350(&t_be_end);
                sub_5AF390("Back end time", &t_be_start, &t_be_end);
            }
        }
    }
}

The while(1) never actually loops. The call to sub_5AF1D0 is __noreturn (it calls exit() or abort() internally), so control never reaches the second iteration. The compiler arranged the basic blocks this way because the backend code at backend: is reached via a goto from the error-gate check, placing it logically "after" the exit call in the CFG.

Phase 1: fe_pre_init -- sub_585D60 (0x585D60)

The first thing main() does after redirecting stderr is call sub_585D60, which performs the absolute minimum initialization needed before command-line parsing can proceed. This function lives in fe_init.c and makes 9 sequential calls to subsystem pre-initializers, plus two inline global assignments.

Pre-Init Call Table

host_envir_early_init (sub_5B1E70) Detail

This is the most substantial pre-init call. It initializes the host environment interface layer from host_envir.c:

Signal handlers (one-time, guarded by dword_E6E120):

After signal setup, dword_E6E120 is set to 0 so handlers are registered only once.

Locale: Calls newlocale(LC_NUMERIC, "C", 0) then uselocale() to force the C locale for numeric output. If either call fails, asserts with "could not set LC_NUMERIC locale" at host_envir.c:264.

Working directory: Iteratively calls getcwd() with a growing buffer (starting at 256 bytes, expanding by 256 on ERANGE) until it fits, then copies the result into qword_126EEA0 via permanent allocation.

Environment variables:

• EDG_BASE -- read into qword_126EE38 (base path for EDG data files; empty string if unset)
• EDG_SUPPRESS_ASSERTION_LINE_NUMBER -- if set and not "0", sets dword_126ED40 = 1 (suppress line numbers in internal assertion messages)

CPU time limit: Calls getrlimit(RLIMIT_CPU) then setrlimit() with rlim_cur = RLIM_INFINITY to disable the CPU time limit.

Global zeroing: Zeros ~50 host-environment globals including file descriptors, path buffers, platform flags, output filename pointers.

Language mode: Sets dword_126EFB4 = 2 (default to C++ mode -- this is later overridden by CLI parsing if -x c is specified).

Sentinel validation: Checks off_E6E0E0 against the string "last" to verify that the predef_macro_mode_names table was properly initialized at link time. On mismatch, asserts with "predef_macro_mode_names not initialized properly" at host_envir.c:6927.

Stack Limit Adjustment

Between CLI parsing and heavy initialization, main() conditionally raises the process stack limit:

if (dword_106C064 && !getrlimit(RLIMIT_STACK, &rlimits)) {
    original_stack = rlimits.rlim_cur;
    rlimits.rlim_cur = rlimits.rlim_max;   // raise soft to hard limit
    stack_was_raised = (setrlimit(RLIMIT_STACK, &rlimits) == 0);
}

The flag dword_106C064 is set to 1 by default in sub_45EB40 (cmd_line_pre_init) and can be disabled via the --modify_stack_limit=false CLI flag. The purpose is to prevent stack overflow during deep recursion in the C++ parser, template instantiation engine, and constexpr interpreter. After compilation completes (just before exit), main() restores the original rlim_cur value.

Phase 3: fe_one_time_init -- sub_585DB0 (0x585DB0)

This is the heaviest initialization stage. It zeroes the token state (qword_126DD38 -- 6 bytes packed as a dword + word), optionally calls sub_5AF330 for profiling init if dword_106BD4C is set, then makes 38 sequential calls to subsystem one-time initializers.

One-Time Init Call Table

The call order reflects dependency constraints: types before scopes, scopes before declarations, declarations before expressions, expressions before the parser, etc. Template initialization is split across three calls (#16, #33, #37) because different phases of template support depend on different subsystems being initialized first.

Function Pointer Table Validation

After all 38 initializers complete, sub_585DB0 performs a critical integrity check:

if (funcs_6F71AE || off_D560C0 != nullsub_6)
    sub_4F21C0("function_pointers is incorrectly initialized");

This validates two conditions:

1. funcs_6F71AE must be zero. This global acts as a "dirty flag" -- if any initializer wrote a nonzero value here, the table was not properly zeroed during static initialization.

2. off_D560C0 must point to nullsub_6 (0x585B00). The address off_D560C0 is the last entry in a function pointer dispatch table in .rodata. The empty function nullsub_6 acts as a sentinel -- its known address is compared against the table's last slot to verify that the table was correctly populated at link time. If the linker reordered or dropped entries, the sentinel would not match.

If either check fails, sub_4F21C0 emits a fatal diagnostic ("function_pointers is incorrectly initialized") and then falls through to sub_585EE0 (fe_init_part_1) regardless -- this is a non-recoverable error that will likely cause crashes later, but the code attempts to continue.

On successful validation, sub_585DB0 returns without calling sub_585EE0. However, sub_585EE0 is actually called from a different path: the normal flow is that sub_585DB0 returns, and main() proceeds. The sub_585EE0 call on the error path in sub_585DB0 appears to be a fallthrough from the panic handler.

Correction from the sweep report: Examination of the actual decompiled code shows that sub_585EE0 (fe_init_part_1) is called only on the error path of the sentinel check within sub_585DB0. On the normal (no-error) path, sub_585DB0 returns sub_7DF400()'s return value directly. This means fe_init_part_1 is called from the sentinel-check error handler, not from the main success path of sub_585DB0. The actual invocation of fe_init_part_1 in the normal flow must occur elsewhere in the pipeline (likely called from within one of the subsystem initializers or from sub_7A40A0).

fe_init_part_1 -- sub_585EE0 (0x585EE0)

This function performs per-compilation-unit initialization. It is identified by the debug trace string "fe_init_part_1" at level 5 and an assertion path fe_init.c:2007. Its responsibilities:

Compilation Timestamp

time(&timer);
char *t = ctime(&timer);
if (!t) t = "Sun Jan 01 00:00:00 1900\n";
if (strlen(t) > 127)
    assert("fe_init.c", 2007, "fe_init_part_1");  // buffer overflow guard
strcpy(byte_106B5C0, t);   // 128-byte timestamp buffer
dword_126EE48 = 1;         // init-complete flag

Per-Unit Initializer Call Table

After the timestamp, sub_585EE0 calls 33 per-compilation-unit initializers:

Compilation Mode Flags

After the per-unit initializers, sub_585EE0 copies global configuration values (set during CLI parsing) into the compilation-mode descriptor at 0x126EB88:

Output File Setup

if (dword_106C298) {                          // output enabled
    if (qword_106C278)                        // output path specified
        qword_106C280 = sub_4F48F0(path, 0, 0, 16, 1513);  // open file (ID 1513)
    else
        qword_106C280 = stdout;               // default to stdout
}
sub_5AEDB0();                                 // write output header

The output file ID 1513 is one of three output file slots used during compilation (1513, 1514, 1515).

Initialization Summary

The total initialization sequence before parsing begins involves 80+ subsystem init calls across three layers:

main()
 ├─ sub_585D60()  fe_pre_init           9 subsystem pre-inits
 │   ├─ sub_48B3C0   error              4 globals zeroed
 │   ├─ sub_6BB290   srcfile            10 globals zeroed
 │   ├─ sub_5B1E70   host_envir         signals, locale, CWD, env vars, ~50 globals
 │   ├─ sub_752C90   types              type table alloc, compiler defaults
 │   ├─ sub_45EB40   cmd_line           272-flag bitmap, ~350 config defaults
 │   ├─ sub_4ED530   declarations       error counters, diagnostic severity table (15KB)
 │   ├─ sub_6F6020   il                 3 globals zeroed
 │   ├─ [inline]     scope indices      dword_126C5E4 = dword_126C5C8 = -1
 │   ├─ sub_7A48B0   tu_tracking        13 globals zeroed
 │   └─ sub_7C00F0   templates          1 global zeroed
 │
 ├─ sub_459630()  proc_command_line     276 flags → ~150 config globals
 │
 ├─ [RLIMIT_STACK adjustment]           raise soft limit to hard limit
 │
 └─ sub_585DB0()  fe_one_time_init      38 subsystem one-time inits
     ├─ token state zeroing             qword_126DD38 = 0 (6 bytes)
     ├─ 38 subsystem calls              types → scopes → errors → ... → CUDA
     ├─ sentinel check                  funcs_6F71AE == 0 && off_D560C0 == nullsub_6
     └─ sub_585EE0()  fe_init_part_1    (on error path, or called from subsystem)
         ├─ compilation timestamp       byte_106B5C0 via ctime()
         ├─ 33 per-unit inits           declarations → overload → ... → templates
         ├─ compilation mode flags      copy CLI config into descriptor struct
         ├─ output file open            stdout or file (ID 1513)
         └─ sub_5AEDB0()               write output header

Global State Set Before Parsing

By the time sub_7A40A0 (process_translation_unit) is called, the following critical globals have been established:

The Error Gate

The transition from frontend to backend is controlled by a simple error check:

if (!qword_126ED90)          // qword_126ED90 = error count from frontend
    goto backend_label;      // no errors → run backend
dword_106C254 = 1;           // errors → set skip-backend flag

When dword_106C254 == 1, the backend stage (sub_489000) is skipped entirely. The process still writes a signoff trailer and exits with a nonzero status code. This means a cudafe++ compilation with frontend errors produces no .int.c output file -- the backend never runs.

Exit Code Mapping

The exit function sub_5AF1D0 at 0x5AF1D0 maps internal status codes to process exit codes:

In SARIF mode (dword_106BBB8 set), the text messages ("Compilation terminated.", "Compilation aborted.") are suppressed, but exit codes remain identical.

Cross-References

• Pipeline Overview -- complete 8-stage pipeline diagram
• CLI Processing -- detailed breakdown of sub_459630 and all 276 flags
• Frontend Invocation -- sub_7A40A0 (process_translation_unit) internals
• Frontend Wrapup -- 5-pass architecture of sub_588F90
• Backend Code Generation -- sub_489000 (.int.c emission)
• Timing & Exit -- sub_5AF350/sub_5AF390/sub_5AF1D0 details
• EDG Overview -- EDG 6.6 source tree and NVIDIA modifications
• EDG Lexer -- keyword registration performed during sub_5863A0

CLI Processing

proc_command_line (sub_459630) at 0x459630 is a 21,773-byte function (4,105 decompiled lines, 296 callees) in cmd_line.c that parses the entire cudafe++ command line. It registers 276 flags into a flat lookup table, iterates argv with prefix-matching against that table, dispatches each matched flag through a 275-case switch statement, then resolves language dialect settings and opens output files. This function is the second stage of the pipeline, called directly from main() at 0x408950 before any heavy initialization.

Nobody invokes cudafe++ directly. NVIDIA's driver compiler nvcc decomposes its own options and passes the appropriate low-level flags via -Xcudafe <flag>. The full flag inventory is in CLI Flag Inventory; this page documents the implementation mechanics of the parsing system itself.

Key Facts

Flag Table Layout

The flag table is a contiguous array starting at dword_E80060. Each of the 552 slots occupies 40 bytes. The current entry count is tracked in dword_E80058.

Offset   Field            Type       Access pattern
------   -----            ----       --------------
+0       case_id          int32      dword_E80060[idx * 10]
+8       name             char*      qword_E80068[idx * 5]
+16      short_char       int16      word_E80070[idx * 20]    (low byte = char, high byte = 1)
+17      is_valid         int8       (high byte of short_char word, always 1)
+18      takes_value      int8       byte_E80072[idx * 40]
+19      visible          int8       (part of dword_E80080[idx * 10] at +32)
+20      is_boolean       int8       byte_E80073[idx * 40]
+24      name_length      int64      qword_E80078[idx * 5]    (precomputed strlen)
+32      mode_flag        int32      dword_E80080[idx * 10]

The flag-was-set bitmap at byte_E7FF40 spans 0x110 bytes (272 slots). When a flag is matched during parsing, the corresponding byte is set to 1 to record that the user explicitly provided it. The bitmap is zeroed by default_init (sub_45EB40) before every compilation.

Registration: sub_452010 (init_command_line_flags)

sub_452010 at 0x452010 is a 30,133-byte function (3,849 decompiled lines) that populates the entire flag table. It is called once, at line 280 of proc_command_line, before the parsing loop begins.

register_command_flag (sub_451F80)

Each flag is registered through sub_451F80 (25 lines), called approximately 275 times from sub_452010:

void register_command_flag(
    int    case_id,       // dispatch ID for the switch (1-275)
    char*  name,          // flag name without dashes ("preprocess", "timing", etc.)
    char   short_opt,     // single-char alias ('E', '#', etc.), 0 for none
    char   takes_value,   // 1 if the flag requires =<value>
    int    mode_flag,     // visibility/classification (mode vs. action)
    char   enabled        // whether the flag is active (1 = registered, 0 = disabled)
);

The function writes into the next free slot at index dword_E80058, precomputes strlen(name) into name_length, always sets the is_valid byte to 1, then increments the counter. If the counter reaches 552, it panics via sub_40351D -- the table is statically sized.

Paired Toggle Registration

Approximately half of all flags are boolean toggles registered as pairs: --flag and --no_flag share the same case_id but differ in which value they write. Pairs are registered in two ways:

1. Two sequential register_command_flag calls -- both point to the same case_id; the parsing loop determines whether the matched name starts with no_ and sets the target global to 0 or 1 accordingly.

2. Inline table population -- seven additional paired flags (relaxed_abstract_checking, concepts, colors, keep_restrict_in_signatures, check_unicode_security, old_id_chars, add_match_notes) are written directly into the array without going through register_command_flag.

Parsing Loop

After flag registration, proc_command_line performs five sequential setup steps, then enters the main argv iteration.

Pre-Loop Setup

Step 1:  Initialize qword_126DD38, qword_126EDE8 (token state / source position)
Step 2:  Call sub_452010() -- register all 276 flags
Step 3:  Allocate 4 hash tables (16-byte header + 256-byte data each):
           qword_106C248  macro define/alias map
           qword_106C240  include path list
           qword_106C238  system include map
           qword_106C228  additional system include map
Step 4:  Suppress 9 diagnostic numbers by default:
           1257, 1373, 1374, 1375, 1633, 2330, 111, 185, 175
         Each via sub_4ED400(number, suppress_severity, 1)
Step 5:  Set dword_E7FF20 = 1 (argv index, skipping argv[0])

The default-suppressed diagnostics are EDG warnings that NVIDIA considers noise for CUDA compilation. Diagnostic 111 ("statement is unreachable"), 185 ("pointless comparison of unsigned integer with zero"), and 175 ("subscript out of range") are common false positives in CUDA template-heavy code.

argv Iteration

The loop processes argv[dword_E7FF20] through argv[argc-1]. For each argument:

1. Dash detection -- if the argument does not start with -, it is treated as the input filename (stored in qword_126EEE0). Only one non-flag argument is expected.

2. Short flag matching -- for single-dash arguments (-X), the parser scans the flag table for an entry whose short_char matches. If the flag takes_value, the next argv element is consumed as the value.

3. Long flag matching -- for double-dash arguments (--flag-name), the parser calls parse_flag_name_value (sub_451EC0) to split on =:

// sub_451EC0: split "--name=value" into name and value
// Respects backslash escapes and quoted strings
// If no '=' found: *name_out = src, *value_out = NULL
void parse_flag_name_value(char* src, char** name_out, char** value_out);

The name portion is then matched against the flag table using strncmp with each entry's precomputed name_length. The parser iterates all entries and counts exact and prefix matches:

• Exact match (length equals name_length and strncmp returns 0) -- dispatches immediately.
• Unique prefix match (only one entry's name starts with the given prefix) -- dispatches to that entry.
• Ambiguous prefix (multiple entries match the prefix) -- emits error 923 ("ambiguous command-line option").
• No match -- the argument is silently ignored or treated as input.

Conflict Detection

Before the main loop, check_conflicting_flags (sub_451E80, 15 lines) validates that mutually exclusive flags were not specified together. It checks byte_E7FFF2 || byte_E80031 || byte_E80032 || byte_E80033, corresponding to flags 3, 193, 194, and 195. If any conflict is detected, it emits error 1027 via sub_4F8480.

The Dispatch Switch (275 Cases)

After a flag is matched, its case_id indexes into a giant switch statement occupying the bulk of proc_command_line. The following sections document the most important cases grouped by function.

Preprocessor Control (Cases 3--9)

Compilation Mode (Cases 14, 20--26)

Diagnostic Control (Cases 39--44)

Cases 39--43 (diag_suppress, diag_remark, diag_warning, diag_error, diag_once) share the same value-parsing logic:

1. Read the value string (after '=')
2. Strip leading/trailing whitespace
3. Split on commas
4. For each token:
   a. Parse as integer (diagnostic number)
   b. Call sub_4ED400(number, severity, 1)

The severity values map to:

• Suppress = skip entirely
• Remark = informational (level 4)
• Warning = default warning (level 5)
• Error = hard error (level 7)
• Once = emit on first occurrence only

Case 44 (display_error_number / no_display_error_number) toggles whether error codes appear in diagnostic messages.

CUDA-Specific Flags (Cases 45--89)

Output File Paths

Data Model (Cases 65--66, 90--91)

Device Compilation Control

Template Instantiation (Case 16)

The instantiate flag takes a string value and sets dword_106C094:

Include and Macro Arguments (Cases 29--31)

Cases 29 (include_directory / -I) and 167 (sys_include) append entries to linked lists via sub_4595D0:

// sub_4595D0: append_to_linked_list
// Allocates a 24-byte node: {next_ptr, string_ptr, int_field}
// Appends to singly-linked list with head/tail pointers
void append_to_linked_list(list_head*, char* string, int type);

A special case: -I- (the literal string "-") sets a flag for stdin include mode rather than appending to the path list. It calls sub_5AD0A0 for the actual path registration.

Case 30 (define_macro / -D) builds a linked list of macro definitions via sub_4595D0. Case 31 (undefine_macro / -U) allocates the same 24-byte node but marks the int_field as 1 to indicate undefine.

Language Standard Selection (Cases 228, 240--252)

These cases set dword_126EF68 -- the internal value of __cplusplus or __STDC_VERSION__:

SM Architecture Target (Case 245)

case 245:  // --target=<sm_arch>
    dword_126E4A8 = sub_7525E0(value_string);

sub_7525E0 parses the SM architecture string (e.g., "sm_90", "sm_100") and returns the internal architecture code stored in dword_126E4A8. This value gates which CUDA features are available during compilation (see Architecture Feature Gating).

Host Compiler Compatibility (Cases 182--188)

Raw Flag Manipulation (Case 193)

case 193:  // --set_flag=<name>=<value>  or  --clear_flag=<name>
    // Looks up <name> in off_D47CE0 (a name-to-address lookup table)
    // Sets the corresponding global to <value> (integer)

This is a backdoor for nvcc to set arbitrary internal globals by name, used for flags that do not have dedicated case_id entries.

Output Mode (Case 274)

case 274:  // --output_mode=text  or  --output_mode=sarif
    if (strcmp(value, "text") == 0)
        output_mode = 0;     // plain text diagnostics (default)
    else if (strcmp(value, "sarif") == 0)
        output_mode = 1;     // SARIF JSON diagnostics

SARIF (Static Analysis Results Interchange Format) output is used by IDE integrations and CI pipelines. When enabled, diagnostic messages are emitted as structured JSON instead of traditional file:line: error: format.

Dump Options (Case 273)

case 273:  // --dump_command_options
    // Iterates the entire flag table
    // For each entry where is_valid == 1:
    //   printf("--%s ", name);
    // Then exits

This is a diagnostic/debug mode that prints every registered flag name and exits. Used by nvcc to discover the cudafe++ flag namespace.

Post-Parsing: Dialect Resolution

After the argv loop exits, proc_command_line enters a massive dialect resolution block (approximately 800 lines). This phase reconciles the various mode flags into a consistent configuration.

Input Filename Extraction

The last non-flag argv element is the input filename, stored in qword_126EEE0. This pointer is later passed to process_translation_unit (sub_7A40A0) in stage 5 of the pipeline.

Memory Region Initialization

Eleven memory regions (numbered 1--11) are initialized with default configurations. These correspond to CUDA memory spaces (global, shared, constant, local, texture, etc.) and are used by the frontend to track address space qualifiers.

GCC/Clang Feature Resolution

The resolver checks GCC version thresholds to decide which extensions to enable:

GCC version thresholds (stored as integer * 100):
  40299 (0x9D6B)  -- GCC 4.2.99 boundary
  40599 (0x9E97)  -- GCC 4.5.99 boundary
  40699 (0x9EFB)  -- GCC 4.6.99 boundary
  etc.

For each threshold, specific feature flags are conditionally enabled. For example, if GCC version >= 40599, rvalue references and variadic templates are enabled even if the language standard is technically C++03. This emulates how GCC provides extensions ahead of standards.

C++ Standard Feature Cascade

Based on the value of dword_126EF68 (__cplusplus), the resolver enables feature flags in a cascade:

199711 (C++98):  base features only
201103 (C++11):  + lambdas, rvalue_refs, auto_type, nullptr,
                   variadic_templates, unrestricted_unions,
                   delegating_constructors, user_defined_literals, ...
201402 (C++14):  + digit_separators, generic lambdas, relaxed_constexpr
201703 (C++17):  + exc_spec_in_func_type, aligned_new, if_constexpr,
                   structured_bindings, fold_expressions, ...
202002 (C++20):  + concepts, modules, coroutines, consteval, ...
202302 (C++23):  + deducing_this, multidimensional_subscript, ...

Conflict Validation

Post-dialect resolution performs consistency checks:

• If both gcc and clang modes are enabled, GCC takes precedence
• If cfront_2.1 or cfront_3.0 is set alongside modern C++ features, features are silently disabled
• If no_exceptions is set but coroutines is requested, coroutines are disabled (they require exceptions)

Output File Opening

After all flags are resolved:

• The output .int.c file is opened (path from case 45/gen_c_file_name, or stdout if path is "-")
• The error output file is opened if --error_output was specified (case 35)
• The listing file is opened if --list was specified (case 33)

Default Diagnostic Severity Overrides

Nine diagnostic numbers are suppressed by default before any user --diag_suppress flags are processed:

Users can override these defaults with explicit --diag_error=111 (or similar) on the command line, since user-specified severity always wins.

Key Helper Functions

Key Global Variables

Annotated Parsing Flow

int64_t proc_command_line(int argc, char** argv)
{
    // --- Phase 1: Global state init ---
    qword_126DD38 = 0;                         // zero token state
    qword_126EDE8 = 0;                         // zero source position

    // --- Phase 2: Register all flags ---
    init_command_line_flags();                  // sub_452010: 3849 lines, 276 flags

    // --- Phase 3: Allocate hash tables ---
    qword_106C248 = alloc_hash_table();        // macro defines/aliases
    qword_106C240 = alloc_hash_table();        // include paths
    qword_106C238 = alloc_hash_table();        // system includes
    qword_106C228 = alloc_hash_table();        // additional system includes

    // --- Phase 4: Default diagnostic suppressions ---
    set_diagnostic_severity(1257, SUPPRESS, 1);
    set_diagnostic_severity(1373, SUPPRESS, 1);
    set_diagnostic_severity(1374, SUPPRESS, 1);
    set_diagnostic_severity(1375, SUPPRESS, 1);
    set_diagnostic_severity(1633, SUPPRESS, 1);
    set_diagnostic_severity(2330, SUPPRESS, 1);
    set_diagnostic_severity(111,  SUPPRESS, 1);
    set_diagnostic_severity(185,  SUPPRESS, 1);
    set_diagnostic_severity(175,  SUPPRESS, 1);

    // --- Phase 5: Main parsing loop ---
    for (int i = 1; i < argc; i++) {
        char* arg = argv[i];
        if (arg[0] != '-') {
            qword_126EEE0 = arg;               // input filename
            continue;
        }

        // Split --name=value
        char *name, *value;
        parse_flag_name_value(arg + 2, &name, &value);  // sub_451EC0

        // Match against flag table
        int match_count = 0;
        int matched_id = -1;
        for (int f = 0; f < dword_E80058; f++) {
            if (strncmp(name, flag_table[f].name, strlen(name)) == 0) {
                if (strlen(name) == flag_table[f].name_length) {
                    matched_id = flag_table[f].case_id;  // exact match
                    break;
                }
                match_count++;
                matched_id = flag_table[f].case_id;
            }
        }

        if (match_count > 1) {
            error(923);  // "ambiguous command-line option"
            continue;
        }

        byte_E7FF40[matched_id] = 1;           // mark flag as set

        switch (matched_id) {
            case 3:   /* no_line_commands */ ...
            case 4:   /* preprocess */      ...
            ...
            case 274: /* output_mode */     ...
            case 275: /* incognito */       ...
        }
    }

    // --- Phase 6: Post-parsing dialect resolution ---
    // ~800 lines: resolve gcc/clang versions, cascade C++ features,
    // validate consistency, open output files

    // --- Phase 7: Memory region init (1-11) ---
    // Initialize CUDA memory space descriptors

    return 0;
}

Version Banner

Case 21 (--version / -v) prints the following banner to stdout (does not exit):

cudafe: NVIDIA (R) Cuda Language Front End
Portions Copyright (c) 2005, 2024-YYYY NVIDIA Corporation
Portions Copyright (c) 1988-2018, 2024 Edison Design Group Inc.
Based on Edison Design Group C/C++ Front End, version 6.6 (BUILD_DATE BUILD_TIME)
Cuda compilation tools, release 13.0, V13.0.88

Case 92 (--Version / -V) prints a different copyright format and then calls exit(1). This variant is used for machine-parseable version queries.

Relationship to Pipeline

proc_command_line is called as stage 2 of the pipeline, after fe_pre_init (sub_585D60) has initialized signal handlers, locale, working directory, and default config:

main()
  |-- sub_585D60()           [1] fe_pre_init (10 subsystem pre-initializers)
  |-- sub_5AF350()               capture_time (total start)
  |-- sub_459630(argc, argv) [2] proc_command_line  <-- THIS FUNCTION
  |-- setrlimit()                conditional stack raise (gated by dword_106C064)
  |-- sub_585DB0()           [3] fe_one_time_init (38 subsystem initializers)
  ...

By the time proc_command_line returns, every global configuration variable is set to its final value. The subsequent fe_one_time_init phase reads these globals to configure keyword tables, type system parameters, and per-translation-unit state.

Frontend Invocation

process_translation_unit (sub_7A40A0, 1267 bytes at 0x7A40A0, from EDG source file trans_unit.c) is the main frontend workhorse -- stage 5 of the pipeline. Called once from main(), it orchestrates the entire transformation from .cu source text to a fully-built EDG IL tree. The function allocates a 424-byte translation unit descriptor, opens the source file via the lexer, drives the C++ parser to completion, runs semantic analysis on the parsed declarations, and finally performs per-TU wrapup (stop-token verification, class linkage checking, module finalization). By the time it returns, every declaration, type, expression, and statement from the source has been parsed into IL nodes, CUDA execution-space attributes have been resolved, and the TU is linked into the global TU chain ready for the 5-pass fe_wrapup stage.

Key Facts

Annotated Decompilation

int process_translation_unit(char *filename,       // source file path
                              int is_recompilation, // nonzero on error-retry pass
                              void *module_info)    // non-NULL for C++20 module TUs
{
    bool is_primary = (module_info == NULL);

    // --- Debug trace on entry ---
    if (debug_verbosity > 0 || (debug_enabled && trace_category("trans_unit")))
        fprintf(stderr, "Processing translation unit %s\n", filename);

    // --- Module-mode state validation ---
    // If this is a primary TU (no module_info) but we've already seen a module TU,
    // that's an internal consistency error.
    if (is_recompilation)
        goto skip_validation;
    if (!is_primary) {
skip_validation:
        if (module_info)
            has_seen_module_tu = 1;                          // dword_12C7A88
        goto proceed;
    }
    if (has_seen_module_tu)
        assertion_failure("trans_unit.c", 696, "process_translation_unit", 0, 0);

proceed:
    // --- Save previous TU state if any ---
    if (current_translation_unit)                            // qword_106BA10
        save_translation_unit_state(current_translation_unit);  // sub_7A3A50

    // --- Reset per-TU compilation state ---
    current_source_position = 0;                             // qword_126DD38
    is_recompilation_flag = is_recompilation;                 // dword_106BA08
    current_filename = filename;                             // qword_106BA00
    has_module_info = (module_info != NULL);                  // dword_106B9F8

    // --- Initialize error/parser state ---
    reset_error_state();                                     // sub_5EAEC0
    if (is_recompilation)
        fe_init_part_1();                                    // sub_585EE0

    // ==========================================================
    //  PHASE 1: Allocate and initialize TU descriptor (424 bytes)
    // ==========================================================
    registration_complete = 1;                               // dword_12C7A8C
    tu_descriptor *tu = allocate_storage(424);               // sub_6BA0D0
    tu->next_tu = NULL;                                      // [0]
    ++tu_count;                                              // qword_12C7A78
    tu->storage_buffer = allocate_storage(per_tu_storage_size); // [16], sub_6BA0D0
    tu->tu_name = NULL;                                      // [8]
    init_scope_state(tu + 24);                               // sub_7046E0, offsets [24..192]
    tu->field_192 = 0;
    tu->field_352 = 0;
    tu->field_184 = 0;
    memset(&tu->scope_decl_area, 0, ...);                    // [200..360] zeroed
    tu->field_360 = 0;
    tu->field_368 = 0;
    tu->field_376 = 0;
    tu->flags = 0x0100;                                      // [392] = "initialized"
    tu->error_severity_count = 0;                            // [408]
    tu->field_416 = 0;

    // --- Copy registered variable defaults into per-TU storage ---
    for (reg = registered_variable_list; reg; reg = reg->next) {
        if (reg->offset_in_tu)
            *(tu + reg->offset_in_tu) = reg->variable_value;
    }

    // --- Set module info pointer and primary flag ---
    tu->module_info_ptr = module_info;                       // [376]
    tu->is_primary = is_primary;                             // [392] byte 0

    // ==========================================================
    //  PHASE 2: Link TU into global chains
    // ==========================================================

    // --- Set as primary TU if this is the first ---
    if (primary_translation_unit == NULL) {                   // qword_106B9F0
        primary_translation_unit = tu;
        if (!is_recompilation)
            assertion_failure("trans_unit.c", 725, "process_translation_unit", 0, 0);
    }

    // --- Push onto TU stack ---
    current_translation_unit = tu;                           // qword_106BA10
    // (stack entry allocated from free list or via permanent alloc)
    stack_entry = alloc_stack_entry();                        // 16 bytes
    stack_entry->tu_ptr = tu;
    stack_entry->next = tu_stack_top;
    if (tu != primary_translation_unit)
        ++tu_stack_depth;                                    // dword_106B9E8
    tu_stack_top = stack_entry;                               // qword_106BA18

    // --- Append to TU linked list ---
    if (tu_chain_tail)                                       // qword_12C7A90
        tu_chain_tail->next_tu = tu;
    tu_chain_tail = tu;

    // ==========================================================
    //  PHASE 3: Source file setup + parse
    // ==========================================================

    if (module_info) {
        // --- Module compilation path ---
        // Extract header info from module descriptor
        module_id = module_info[7];
        module_info[2] = tu;                                 // back-link TU into module
        current_module_id = module_id;                       // qword_106C0B0
        // ... copy include paths, source paths from module descriptor ...
        source_dir = intern_directory_path(filename, 1);     // sub_5ADC60
        set_include_paths(source_dir, &include_list, &sys_list); // sub_5AD120
        fe_translation_unit_init(source_dir, &include_list); // sub_5863A0
        import_module = module_info[3];
        tu->error_severity_count = current_error_severity;   // [408]
        set_module_id(import_module);                        // sub_5AF7F0
        if (preprocessing_only)                              // dword_106C29C
            goto compile;
        goto compile_module;
    }

    // --- Standard (non-module) path ---
    fe_translation_unit_init(0, 0);                          // sub_5863A0
    tu->error_severity_count = current_error_severity;
    if (preprocessing_only)
        goto compile;

    // --- PCH header processing (optional) ---
    if (pch_enabled && !pch_skip_flag) {                     // dword_106BF18, dword_106B6AC
        setup_pch_source();                                  // sub_5861C0
        precompiled_header_processing();                     // sub_6F4AD0
    }

compile:
    // --- Main compilation: parse + build IL ---
    compile_primary_source();                                // sub_586240
    semantic_analysis();                                     // sub_4E8A60 (standard path)
    goto wrapup;

compile_module:
    compile_primary_source();                                // sub_586240
    module_compilation();                                    // sub_6FDDF0 (module path)

wrapup:
    // ==========================================================
    //  PHASE 4: Per-TU wrapup + stack pop
    // ==========================================================
    translation_unit_wrapup();                               // sub_588E90

    // --- Pop TU stack (inlined pop_translation_unit_stack) ---
    top = tu_stack_top;
    popped_tu = top->tu_ptr;
    if (popped_tu != current_translation_unit)
        assertion_failure("trans_unit.c", 556,
                          "pop_translation_unit_stack", 0, 0);
    if (popped_tu != primary_translation_unit)
        --tu_stack_depth;
    tu_stack_top = top->next;
    // (return stack entry to free list)
    if (tu_stack_top)
        switch_translation_unit(tu_stack_top->tu_ptr);       // sub_7A3D60

    // --- Debug trace on exit ---
    if (debug_verbosity > 0 || (debug_enabled && trace_category("trans_unit")))
        fprintf(stderr, "Done processing translation unit %s\n", filename);
}

Execution Flow

process_translation_unit (sub_7A40A0)
  |
  |-- [1] Debug trace: "Processing translation unit %s"
  |-- [2] Module-state validation (assert at trans_unit.c:696)
  |-- [3] Save previous TU state (sub_7A3A50)
  |-- [4] Reset error state (sub_5EAEC0)
  |-- [5] If recompilation: re-run fe_init_part_1 (sub_585EE0)
  |
  |-- [6] Allocate 424-byte TU descriptor (sub_6BA0D0)
  |       |-- Allocate per-TU storage buffer (sub_6BA0D0(per_tu_storage_size))
  |       |-- Initialize scope state at [24..192] (sub_7046E0)
  |       |-- Zero remaining fields [192..416]
  |       |-- Copy registered variable defaults
  |       |-- Set module_info_ptr [376] and flags [392]
  |
  |-- [7] Set as primary TU if first (assert at trans_unit.c:725)
  |-- [8] Push onto TU stack, link into TU chain
  |
  |-- [9] Module path? (module_info != NULL)
  |       |-- YES: Extract module header info
  |       |        sub_5ADC60 (intern_directory_path)
  |       |        sub_5AD120 (set_include_paths)
  |       |        sub_5863A0 (fe_translation_unit_init)
  |       |        sub_5AF7F0 (set_module_id)
  |       |
  |       |-- NO:  sub_5863A0 (fe_translation_unit_init) with NULL args
  |                sub_5861C0 + sub_6F4AD0 (PCH processing, if enabled)
  |
  |-- [10] sub_586240  -- compile_primary_source (parser entry)
  |
  |-- [11] Post-parse semantic analysis:
  |        |-- Module path: sub_6FDDF0 (module_compilation)
  |        |-- Standard path: sub_4E8A60 (translation_unit / semantic analysis)
  |
  |-- [12] sub_588E90 -- translation_unit_wrapup
  |
  |-- [13] Pop TU stack (assert at trans_unit.c:556)
  |-- [14] Debug trace: "Done processing translation unit %s"

Phase 1: Error State Reset -- sub_5EAEC0

Before any parsing begins, sub_5EAEC0 resets the parser's error recovery state. This is a tiny function (22 bytes) that configures the error-recovery token scan depth based on whether this is a recompilation pass:

void reset_error_state(void) {
    if (is_recompilation) {            // dword_106BA08
        error_scan_depth = 8;          // dword_126F68C -- shallower scan on retry
        error_scan_mode = 0;           // dword_126F688
        error_recovery_kind = 16;
    } else {
        error_recovery_kind = 24;      // full recovery on first pass
    }
    error_token_limit = error_recovery_kind;  // dword_126F694
    error_count_local = 0;                     // dword_126F690
}

The different error_recovery_kind values (16 vs 24) control how aggressively the parser attempts to resynchronize after encountering a syntax error. On recompilation (error-retry), the compiler uses a smaller recovery window to avoid cascading errors.

Phase 2: TU Descriptor Allocation

The 424-byte TU descriptor is the central data structure tracking a single translation unit's state during compilation. It is allocated from EDG's permanent storage pool via sub_6BA0D0 and linked into two separate data structures: the TU linked list and the TU stack.

Translation Unit Descriptor Layout (424 bytes)

Registered Variable Mechanism

EDG's multi-TU infrastructure requires certain global variables to be saved and restored when switching between translation units (e.g., during relocatable device code compilation). The mechanism works as follows:

1. Registration phase (during initialization, before any TU processing): Subsystem initializers call f_register_trans_unit_variable (sub_7A3C00) to register global variables that need per-TU state. Each registration creates a 40-byte entry:

2. Accumulated size tracking: Each registration pads the variable's size to 8-byte alignment and adds it to qword_12C7A98 (per-TU storage size). The linked list head is qword_12C7AA8, tail is qword_12C7AA0.

3. TU creation: When a TU descriptor is allocated, a storage buffer of per_tu_storage_size bytes is allocated alongside it at offset [16]. Default values from the field_offset_in_tu entries are copied into the TU descriptor's own fields.

4. TU switching: save_translation_unit_state (sub_7A3A50) iterates the registered variable list, copying each variable's current value from its global address into the outgoing TU's storage buffer. switch_translation_unit (sub_7A3D60) does the reverse: copies from the incoming TU's storage buffer back to the global addresses.

Three core variables are always registered (by sub_7A4690):

Additional variables are registered by other subsystem initializers (trans_corresp registers 3 more via sub_7A3920).

Phase 3: TU Linking and Stack Management

TU Linked List

Translation units are linked in processing order through the next_tu field at offset [0]:

qword_106B9F0 (primary_translation_unit)
  |
  v
  TU_0 --[next_tu]--> TU_1 --[next_tu]--> TU_2 --[next_tu]--> NULL
                                                       ^
                                                       |
                                          qword_12C7A90 (tu_chain_tail)

qword_106B9F0 always points to the first (primary) TU. qword_12C7A90 always points to the last. The chain is walked by fe_wrapup (sub_588F90) during its 5-pass finalization.

TU Stack

The TU stack tracks the active compilation context. Each stack entry is a 16-byte structure:

Stack entries are allocated from a free list (qword_12C7AB8); when the free list is empty, a new 16-byte block is allocated via sub_6B7340 (permanent allocator).

qword_106BA18 (tu_stack_top)
  |
  v
  entry_N: [next=entry_N-1, tu_ptr=current_tu]
  entry_N-1: [next=entry_N-2, tu_ptr=prev_tu]
  ...
  entry_0: [next=NULL, tu_ptr=primary_tu]

The stack depth counter dword_106B9E8 tracks how many non-primary TUs are stacked. It is incremented on push (if tu != primary_tu) and decremented on pop.

The pop operation at the end of process_translation_unit includes an assertion (at trans_unit.c:556) verifying that the top-of-stack TU matches current_translation_unit. This guards against mismatched push/pop sequences, which would corrupt the multi-TU state:

if (stack_top->tu_ptr != current_translation_unit)
    assertion_failure("trans_unit.c", 556, "pop_translation_unit_stack", 0, 0);

Phase 4: Source File Setup

The source file setup differs between standard compilation and C++20 module compilation.

Standard Path (module_info == NULL)

1. sub_5863A0 (fe_translation_unit_init / keyword_init, 1113 lines, fe_init.c): The largest initialization function in the binary. Performs two tasks in sequence:

   • Token state reset: Zeros qword_126DD38 (6-byte source position) and qword_126EDE8 (mirror).
   • Per-TU subsystem reinit: Calls 15+ subsystem re-initializers to prepare for a new compilation unit (source file manager, scope system, preprocessor, diagnostics, etc.).
   • Keyword registration: Registers 200+ C/C++ keywords via sub_7463B0 (enter_keyword), including all C89/C99/C11/C23 keywords, C++ keywords through C++26, GNU extensions, MSVC extensions, Clang extensions, 60+ type traits, and three NVIDIA CUDA-specific type trait keywords (__nv_is_extended_device_lambda_closure_type, __nv_is_extended_host_device_lambda_closure_type, __nv_is_extended_device_lambda_with_preserved_return_type). Keyword registration is version-gated by the language mode (dword_126EFB4) and C++ standard version (dword_126EF68).
   • File scope creation: Calls sub_7047C0(0) to push the initial file scope onto the scope stack.
   • C++ builtins: For C++ mode, registers namespace std, operator new/operator delete allocation functions, std::align_val_t.

2. PCH processing (optional, if dword_106BF18 is set): Calls sub_5861C0 to open the source file with minimal setup (same as sub_586240 but without the recompilation logic), followed by sub_6F4AD0 (precompiled_header_processing, 721 lines, pch.c) which searches for an applicable .pch file, validates memory allocation history, and restores saved variable state from the precompiled header.

Module Path (module_info != NULL)

When compiling a C++20 module unit, the module descriptor (passed as a3) provides pre-computed configuration:

module_info[2] = tu;              // back-link TU into module descriptor
qword_106C0B0 = module_info[7];  // module identifier
qword_126EE98 = module_info[4];  // include path list
qword_126EE78 = module_info[6];  // system include path list
qword_126EE90 = module_info[5];  // additional path list

The module path then calls:

• sub_5ADC60(filename, 1) -- intern the source directory path (cached allocation)
• sub_5AD120(source_dir, &include_list, &sys_list) -- configure include search paths from the module descriptor
• sub_5863A0(source_dir, &include_list) -- fe_translation_unit_init with module-specific paths
• sub_5AF7F0(module_info[3]) -- set the module identifier for this TU (asserts not already set)

Phase 5: Compilation Driver -- sub_586240

sub_586240 (fe_init.c, 63 lines) is the compilation driver that opens the source file and launches the parser. It is called for both standard and module compilation paths.

void compile_primary_source(void) {
    // If recompilation: reset file-scope scope pointer
    if (is_recompilation)
        *(uint64_t *)&xmmword_126EB60 = 0;

    // Allocate mutable copy of filename for the lexer
    char *fn_copy = temp_allocate(strlen(current_filename) + 1);  // sub_5E0460
    strcpy(fn_copy, current_filename);

    // --- Open source file and push onto input stack ---
    open_file_and_push_input_stack(fn_copy, 0, 0, 0, 0, 0, 0, 0, 0, 0);  // sub_66E6E0

    // Record source file descriptor in TU
    current_tu->source_file_entry = *(source_file_descriptor + 64);  // [184]

    // --- Scope handling ---
    if (!pch_mode) {                                         // dword_106B690
        init_global_scope_flag = 1;                          // dword_126C708
        global_scope_decl_list = global_decl_chain;          // qword_126C710
        finalize_scope();                                    // sub_66E920
    }
    open_scope(1, 0);                                        // sub_6702F0

    // --- PCH recompilation metadata ---
    if (is_recompilation) {
        // Allocate 4-byte version marker (3550774 = "6.6\0")
        char *ver = temp_allocate(4);
        *(uint32_t *)ver = 3550774;                          // EDG 6.6 version tag
        edg_version_ptr = ver;                               // qword_126EB78
        // Copy compilation timestamp
        char *ts = temp_allocate(strlen(byte_106B5C0));
        compilation_timestamp_copy = strcpy(ts, byte_106B5C0); // qword_126EB80
        dialect_version_snapshot = dialect_version;           // dword_126EBF8
    }

    // --- PCH header loading ---
    if (pch_mode) {
        load_precompiled_header(byte_106B5C0);               // sub_6B5C10
        pch_header_loaded = 1;                               // dword_106B6B0
    }
}

Parser Entry: sub_66E6E0 (open_file_and_push_input_stack)

sub_66E6E0 (lexical.c, 95 lines) is the gateway from file-level compilation into the EDG lexer/parser. It takes 10 parameters controlling how the source file is opened:

The function delegates to sub_66CBD0 which resolves the file path, opens the file handle, and creates the file descriptor. Then sub_66DFF0 pushes the opened file onto the lexer's input stack, making it the active source for tokenization. The lexer reads from this stack via get_next_token (sub_676860, 1995 lines).

At debug verbosity > 3, it prints: "open_file_and_push_input_stack: skipping guarded include file %s\n" when an include guard causes the file to be skipped.

Phase 6: Semantic Analysis -- sub_4E8A60

After parsing completes, sub_4E8A60 (translation_unit, decls.c, 77 lines) performs semantic analysis on the parsed declarations. This function is called only on the standard (non-module) compilation path.

void translation_unit(void) {
    // PCH mode: additional scope finalization
    if (pch_mode)
        finalize_pch_scope();                                // sub_6FC900
    if (global_decl_chain)
        process_pending_declarations();                      // sub_6FDD60

    // --- Main declaration processing loop ---
    declaration_processing_active = 1;                       // dword_126C704
    parse_declaration_seq();                                  // sub_676860 (get_next_token)
    declaration_processing_active = 0;

    // Header-unit stop detection
    if (header_unit_mode)
        finalize_header_unit();                              // sub_6F4A10

    // --- Top-level declaration loop ---
    // Repeatedly processes declarations until token 9 (EOF) is reached.
    // For C++ (dword_126EFB4 == 2) with C++14+ (dword_126EF68 > 201102):
    //   calls sub_6FBCD0 (deferred template processing)
    //   then sub_4E6F80(1, 0) (process next declaration)
    while (current_token != 9) {  // 9 = EOF token
        if (is_cpp && (cpp_version > 201102 || has_cpp20_features))
            process_deferred_templates();                    // sub_6FBCD0
        if (declaration_enabled)
            process_declaration(1, 0);                       // sub_4E6F80
    }

    // --- Post-parse validation ---
    if (!header_unit_mode) {
        if (is_cpp && (cpp_version > 201102 || has_cpp20_features))
            process_deferred_templates();                    // sub_6FBCD0 final pass
        finalize_module_interface();                         // sub_6F81D0
    } else {
        // Header-unit mode assertion: stop position must be found
        assertion_failure("decls.c", 23975, "translation_unit",
                          "translation_unit:", "header stop position not found");
    }
}

The C++ standard version checks (dword_126EF68 > 201102) gate C++14+ features like deferred template instantiation. The value 201102 corresponds to C++11 (__cplusplus value). For C++14 and later, sub_6FBCD0 handles deferred template processing between declaration groups.

Phase 7: Translation Unit Wrapup -- sub_588E90

sub_588E90 (translation_unit_wrapup, fe_wrapup.c, 36 lines) performs per-TU finalization after parsing and semantic analysis are complete. It is the last step before the TU stack is popped.

void translation_unit_wrapup(void) {
    if (debug_enabled)
        trace_enter(1, "translation_unit_wrapup");

    // [1] Stop-token verification
    check_all_stop_token_entries_are_reset(                  // sub_675DA0
        file_scope_stop_tokens + 8);                         // qword_126DB48 + 8

    // [2] Class linkage checking (conditional)
    if (!preprocessing_only) {
        if (rdc_enabled || rdc_alt_enabled)                  // dword_106C2BC, dword_106C2B8
            check_class_linkage();                           // sub_446F80
    }

    // [3] Module import finalization
    finalize_module_imports();                                // sub_7C24D0

    // [4] IL output
    complete_scope();                                        // sub_709250

    // [5] Close file scope
    close_file_scope(1);                                     // sub_7047C0

    // [6] Module correspondence finalization (non-preprocessing)
    if (!preprocessing_only)
        process_verification_list();                         // sub_7A2FE0

    // [7] Write compilation unit boundary
    make_module_id(0);                                       // sub_5AF830

    // [8] Namespace cleanup (C++ only, non-PCH, non-preprocessing)
    if (is_cpp && !is_recompilation && !preprocessing_only)
        namespace_cleanup();                                 // sub_76C910

    if (debug_enabled)
        trace_leave();                                       // sub_48AFD0
}

Sub_675DA0: check_all_stop_token_entries_are_reset

Iterates all 357 entries in the stop-token array. If any nonzero entry is found, logs "stop_tokens[\"%s\"] != 0\n" (using off_E6D240 as the token name table) and asserts with "stop token array not all zero" at lexical.c:17680. This catches lexer state corruption where a stop-token (used during error recovery and tentative parsing) was not properly cleared.

Sub_446F80: check_class_linkage

Called only when relocatable device code (RDC) compilation is enabled (dword_106C2BC or dword_106C2B8). Iterates file-scope type entities looking for class/struct/union types (kind 9-11) and scoped enums (kind 2, bit 3 of +145) that need external linkage for cross-TU visibility. For qualifying types, calls sub_41F800 (make_class_externally_linked) to set the linkage bits at offset +80 to 0x20 (external linkage flag). The function performs a two-pass scan:

1. Pass 1: Identify types needing external linkage. Checks whether the type is used by externally-visible definitions, has nested types with external linkage requirements, or has member functions with non-inline definitions.

2. Pass 2: If any types were promoted, propagates linkage to member functions and nested class template instantiations via sub_41FD90.

Sub_7A2FE0: process_verification_list (Module Finalization)

sub_7A2FE0 (trans_corresp.c, 69 lines) processes the deferred correspondence verification list for multi-TU compilation. This is the mechanism EDG uses to verify that declarations shared across translation units are structurally compatible (One Definition Rule checking for RDC).

void process_verification_list(void) {
    if (is_recompilation || error_count != saved_error_count)
        goto skip;  // skip if new errors appeared

    correspondence_active = 1;                               // dword_106B9E4
    source_seq = *(current_tu + 8);                          // TU source sequence

    prepare_correspondence(source_seq);                      // sub_79FE00
    verify_correspondence(source_seq);                       // sub_7A2CC0

    // Process pending verification items
    while (pending_list) {                                   // qword_12C7790
        pending_list_snapshot = pending_list;
        pending_list = NULL;
        for (item = pending_list_snapshot; item; item = next) {
            next = item->next;
            switch (item->kind) {                            // byte at [8]
                case 0:  break;                              // no-op
                case 2:  verify_typedef_correspondence(item->data);          // sub_7986A0
                case 6:  verify_friend_correspondence(item->data);           // sub_7A1830
                case 7:  verify_nested_class_correspondence(item->data);     // sub_798960
                case 8:  verify_enum_member_correspondence(item->data);      // sub_798770
                case 11: verify_member_function_correspondence(item->data);  // sub_7A1DB0
                case 28: verify_using_declaration_correspondence(item->data);// sub_7982C0
                case 58: verify_base_class_correspondence(item->data);       // sub_7A27B0
                default: assertion_failure("trans_corresp.c", 7709, ...);
            }
            // Return item to free list
            item->next = corresp_free_list;
        }
    }

    correspondence_active = 0;
    correspondence_complete = 1;                             // dword_106B9E0

skip:
    correspondence_complete = 1;
}

The kind codes (0, 2, 6, 7, 8, 11, 28, 58) correspond to EDG declaration kinds: typedef (2), friend (6), nested class (7), enum member (8), member function (11), using declaration (28), base class (58).

Module vs Standard Compilation Path

The control flow diverges based on dword_106C29C (preprocessing-only mode) and the presence of module_info:

                          module_info?
                         /            \
                       YES             NO
                        |               |
                  sub_5ADC60         sub_5863A0(0,0)
                  sub_5AD120              |
                  sub_5863A0       PCH enabled?
                  sub_5AF7F0        /        \
                        |         YES         NO
                        |          |           |
                        |     sub_5861C0       |
                        |     sub_6F4AD0       |
                        |          |           |
                        +-----+----+-----+-----+
                              |          |
                         sub_586240  sub_586240
                              |          |
                      preprocessing_only?
                        /            \
                      YES             NO
                       |               |
                  sub_6FDDF0      sub_4E8A60
                  (module comp)   (standard comp)
                       |               |
                       +-------+-------+
                               |
                         sub_588E90
                    (translation_unit_wrapup)

Note: sub_6FDDF0 is the module compilation driver (59 lines, lower_il.c). It enters a loop calling sub_676860 (get_next_token) until EOF (token 9), processing module import/export declarations. Between module units, it calls sub_66EA70 to close the current input source and advance to the next module partition.

Global State Variables

Translation Unit Tracking

Registration Infrastructure

Statistics Counters

These counters are reported by sub_7A45A0 (print_trans_unit_statistics), which prints formatted memory usage:

trans. unit corresps          N x 24 bytes
translation units             N x 424 bytes
trans. unit stack entry       N x 16 bytes
variable registration         N x 40 bytes

Assertions

The function contains three assertion checks, each producing a fatal diagnostic via sub_4F2930:

Callee Reference Table