EDG 6.6 Overview

cudafe++ is built on top of Edison Design Group's (EDG) commercial C++ frontend, version 6.6. EDG provides the complete C++ language implementation -- lexer, preprocessor, parser, semantic analysis, type system, template instantiation, overload resolution, constant evaluation, and Itanium ABI name mangling. NVIDIA licenses this frontend and compiles it from source with CUDA-specific modifications injected at three distinct integration levels: a dedicated NVIDIA source file (nv_transforms.c), surgical modifications to EDG source files that call into NVIDIA headers, and a large layer of CUDA property-query leaf functions that permeate every compilation phase.

The build path embedded in the binary is:

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

Source Tree

The binary contains debug path references to 52 .c files and 13 .h files. Together these constitute the entire EDG frontend plus NVIDIA's single dedicated source file.

Source Files (.c)

Header Files (.h)

Code Breakdown

The binary contains approximately 6,300 identifiable functions in the EDG portion of the code:

Top 10 Source Files by Function Count

Architecture: Classic Frontend Pipeline

EDG implements a textbook multi-pass compiler frontend. cudafe++ drives it in a single-threaded, sequential pipeline from main() at 0x408950:

  source.cu
     |
     v
  +-----------+     lexical.c, macro.c, preproc.c, literals.c
  |  Lexer /  |     357 token kinds, trigraph handling, raw string
  |  Preproc  |     adjustment, __CUDA_ARCH__ macro injection
  +-----------+
     |  token stream
     v
  +-----------+     expr.c, declarator.c, decl_spec.c, statements.c,
  |  Parser   |     class_decl.c, disambig.c, func_def.c, extasm.c
  |           |     Recursive-descent with disambiguation
  +-----------+
     |  parse tree
     v
  +-----------+     overload.c, exprutil.c, lookup.c, templates.c,
  | Semantic  |     types.c, attribute.c, const_ints.c, folding.c
  | Analysis  |     Type checking, overload resolution, template
  |           |     instantiation, constexpr evaluation
  +-----------+
     |  annotated AST
     v
  +-----------+     il.c, il_alloc.c, il_walk.c, scope_stk.c,
  |  IL Build |     symbol_tbl.c, src_seq.c, trans_unit.c
  |           |     Scope-linked graph of all declarations, types,
  |           |     expressions, statements, templates
  +-----------+
     |  IL graph
     v
  +-----------+     fe_wrapup.c, lower_name.c, trans_corresp.c
  |  Wrapup   |     5-pass finalization: dead code marking,
  |           |     name lowering, cross-TU correspondence (RDC)
  +-----------+
     |  finalized IL
     v
  +-----------+     cp_gen_be.c, nv_transforms.c, host_envir.c
  |  Backend  |     Walk source sequence, emit .int.c file,
  | Emission  |     inject CUDA stubs, lambda wrappers, host
  |           |     reference arrays, managed variable boilerplate
  +-----------+
     |
     v
  output.int.c

The process_translation_unit function (sub_7A40A0 in trans_unit.c) is the main entry point for compilation. It allocates a 424-byte TU descriptor, opens the source file, and orchestrates the parse-to-IL sequence. For the main compilation path, it calls:

1. sub_586240 -- parse the translation unit (drives lexer + parser)
2. sub_4E8A60 -- standard compilation finalization (IL completion)
3. sub_588F90 -- fe_wrapup (5-pass IL finalization)
4. sub_489000 -- backend entry (.int.c emission, "Back end time")

NVIDIA Modifications

NVIDIA's CUDA integration is organized in three layers, from most isolated to most pervasive.

Level 1: NVIDIA-Authored Source (nv_transforms.c + nv_transforms.h)

A single dedicated NVIDIA source file at address range 0x6BAE70--0x6BE4A0, containing approximately 34 functions in ~14KB of code. This file implements all CUDA-specific AST transformations:

The companion header nv_transforms.h declares the API surface that EDG source files call into. This is the primary NVIDIA integration point -- EDG code never calls nv_transforms.c functions directly; it calls through the header's declarations.

Key data structures managed by nv_transforms.c:

Level 2: NVIDIA-Modified EDG Files

Three EDG source files contain direct calls into nv_transforms.h functions, making them the "NVIDIA-aware" EDG files:

cp_gen_be.c -- The backend code generator. When it encounters a type named __nv_lambda_preheader_injection during source sequence walking, it calls nv_emit_lambda_preamble (sub_6BCC20) to inject the entire __nv_* template library. It also calls NVIDIA functions for host reference array emission, managed variable boilerplate, and device stub generation.

class_decl.c -- The class/struct declaration processor. The scan_lambda function (sub_447930, 2113 lines) detects __host__/__device__ annotations on lambda expressions, validates CUDA-specific constraints (35+ error codes in range 3592--3690), and records capture counts in the bitmaps via nv_record_capture_count.

statements.c -- The statement parser. Calls NVIDIA transform functions for statement-level CUDA validation, such as checking that __syncthreads() is not called in divergent control flow within __global__ functions.

Level 3: CUDA Property Query Layer

The most pervasive integration layer consists of 104 small leaf functions clustered at addresses 0x7A6000--0x7AA000 (within types.c). These are type-system query functions that answer questions like "is this type a __device__ pointer?", "does this class have __shared__ storage?", "is this a kernel function type?".

Each follows a canonical pattern:

bool is_<property>_type(type_node *t) {
    while (t->kind == 12)       // 12 = tk_typedef
        t = t->referenced_type; // strip typedef layers
    return <check on underlying type>;
}

These 104 accessors account for 3,648 total call sites across the binary. The top callers by call-site count:

CUDA integration is pervasive because these tiny accessors are called from every phase of compilation -- the parser checks execution space during declaration, semantic analysis validates cross-space calls, the type system queries CUDA qualifiers during overload resolution, and the backend reads them during IL emission. There is no isolated "CUDA layer"; the CUDA awareness is distributed across the entire frontend through these leaf functions.

Type Kind Constants

The type query functions operate on a type_node structure (176 bytes, IL entry kind 6). The kind field at offset +132 encodes:

Memory Management

EDG uses a custom region-based arena allocator implemented in mem_manage.c (address range 0x6B5E40--0x6BA230). Key characteristics:

• Block size: 64KB (0x10000) per block
• Region model: Multiple numbered regions (file-scope = region 1, per-function = region N)
• Free list recycling: Freed blocks go to qword_1280730 for reuse before new allocation
• Trim threshold: Blocks with more than 1,887 unused bytes are split; remainder goes to free list
• Tracking: All allocations recorded for watermark monitoring (qword_1280718 = total, qword_1280710 = peak)
• Dual mode: Malloc-based (mode 0) or mmap-based (mode 1), selected by dword_1280728 from CLI flag

Block structure (48+ bytes header per 64KB block):

The free_fe function (sub_6BA230, 533 lines) implements a hash-table-based deduplicating allocator for front-end object deallocation, using open addressing with linear probing.

C++20 Modules (Stubs)

The modules.c file (address range 0x7C0C60--0x7C2560) contains approximately 20 functions implementing the C++20 module import/export interface. CUDA does not support C++20 modules, so most functions are stubs that return 0:

• has_pending_template_definition_from_module -- returns 0
• has_pending_template_specializations_from_module -- returns 0
• Seven additional stub functions at 0x7C2350--0x7C2410 -- all return 0

The non-stub functions handle the binary module interface file format (magic header {0x9A, 0x13, 0x37, 0x7D}) and basic module name matching, likely preserved from the EDG baseline for future CUDA module support.

Cross-TU Correspondence (RDC Mode)

When compiling with Relocatable Device Code (--rdc), multiple translation units are processed sequentially. The trans_corresp.c file (address range 0x7A00D0--0x7A38A0) implements structural equivalence checking between types from different TUs:

• verify_class_type_correspondence (sub_7A00D0, 703 lines) -- Deep comparison of class types: base classes, friend declarations, member functions, nested types, template parameters
• verify_enum_type_correspondence (sub_7A0E10) -- Enum underlying type and enumerator list comparison
• verify_function_type_correspondence (sub_7A1230) -- Parameter list and return type comparison
• set_type_correspondence (sub_7A1460) -- Links two corresponding types across TUs

The trans_unit.c file manages TU lifecycle with a stack-based model:

process_translation_unit (sub_7A40A0) allocates a 424-byte TU descriptor and drives the parse-to-completion sequence. switch_translation_unit (sub_7A3D60) saves/restores per-TU state (registered variables, scope stack, file scope) when switching between TUs during RDC compilation.

Cross-References

• Pipeline Overview -- How EDG stages map to the 7-stage pipeline
• IL Overview -- The 85 entry kinds that EDG produces
• Extended Lambda Overview -- The nv_transforms.c lambda pipeline in detail
• Type System -- Deep dive on 22 type kinds and class layout
• Template Engine -- Template instantiation worklist
• Name Mangling -- Itanium ABI encoding with CUDA extensions
• Lexer -- Tokenizer and keyword registration
• Overload Resolution -- Candidate evaluation and ICS ranking
• Diagnostics Overview -- The 3,795 error message system