Keyboard shortcuts

Press or to navigate between chapters

Press S or / to search in the book

Press ? to show this help

Press Esc to hide this help

IL Display

The IL display subsystem produces a human-readable text dump of the entire Intermediate Language graph. It is compiled from EDG's il_to_str.c (source path /dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/drivers/compiler/edg/EDG_6.6/src/il_to_str.c, confirmed by an assertion at line 6175 in form_float_constant). The display code occupies address range 0x5EB290--0x603A00 in the binary (roughly 90KB), with the main dispatch functions at 0x5EC600--0x5F7FD0 and formatting helpers continuing through 0x6039E0.

Activation is via the il_display CLI flag (flag index 10 in the boolean flag table), which triggers display_il_file after the frontend completes parsing. The output goes to stdout through an indirectable callback mechanism (qword_126F980). When active, every IL entry in every memory region is printed with labeled fields, 25-column-aligned formatting, and scope/address annotations.

Key Facts

PropertyValue
Source fileil_to_str.c (EDG 6.6)
Address range0x5EB290--0x6039E0
Top-level entry pointsub_5F7DF0 (display_il_file), 56 lines
Header + file-scopesub_5F76B0 (display_il_header), 174 lines
Main dispatchersub_5F4930 (display_il_entry), 1,686 lines
Single-entity displaysub_5F7D50 (display_single_entity), 38 lines
CLI flagil_display (index 10, boolean)
Output callbackqword_126F980 (function pointer, default sub_5EB290 = fputs(s, stdout))
Display-active flagbyte_126FA16 (set to 1 during display)
Scope context flagdword_126FA30 (1 = file-scope region, 0 = function-scope)
Entry kind name tableoff_E6DD80 (~84 entries, indexed by entry kind byte)

Top-Level Control Flow

display_il_file (sub_5F7DF0)
│
│  printf("Display of IL file \"%s\", produced by the compilation of \"%s\"\n",
│         il_file_name, source_file_name)
│
├── display_il_header (sub_5F76B0)
│   │  dword_126FA30 = 1                              // file-scope mode
│   │  puts("\n\nIntermediate language for memory region 1 (file scope):")
│   │  puts("\nil_header:")
│   │  ... 30+ header fields ...
│   │
│   └── walk_file_scope_il(display_il_entry, ...)      // sub_60E4F0
│       └── display_il_entry (sub_5F4930)              // callback per entity
│
└── for region = 2 .. dword_126EC80:
        dword_126FA30 = 0                              // function-scope mode
        // lookup function name from scope table
        printf("\n\nIntermediate language for memory region %ld (function \"%s\"):\n",
               region, function_name)
        walk_routine_scope_il(region, display_il_entry, ...)   // sub_610200
        └── display_il_entry (sub_5F4930)              // callback per entity

Memory region 1 is always file-scope (global declarations, types, templates). Regions 2+ correspond to individual function bodies. The scope table at qword_126EB90 maps each region index to its owning scope entry; the display code checks scope.kind == 17 (sck_function) and extracts the routine name for the banner.

IL Header Fields

display_il_header (sub_5F76B0) prints the translation-unit-level metadata stored in BSS at 0x126EB60--0x126EBF8:

FieldTypeNotes
primary_source_fileIL pointerSource file entry for the main .cu file
primary_scopeIL pointerFile-scope scope entry
main_routineIL pointermain() routine entry, if present
compiler_versionstringEDG compiler version string
time_of_compilationstringBuild timestamp
plain_chars_are_signedboolDefault char signedness
source_languageenumsl_Cplusplus (0) or sl_C (1), from dword_126EBA8
std_versionintegerC/C++ standard version (e.g., 201703 for C++17), from dword_126EBAC
pcc_compatibility_modeboolPCC compatibility
enum_type_is_integralboolWhether enum underlying type is integral
default_max_member_alignmentintegerDefault structure packing alignment
gcc_modeboolGCC compatibility mode
gpp_modeboolG++ compatibility mode
gnu_versionintegerGNU compatibility version number
short_enumsbool-fshort-enums behavior
default_nocommonboolDefault -fno-common
UCN_identifiers_usedboolUniversal character names in identifiers
vla_usedboolVariable-length arrays present
any_templates_seenboolWhether any templates were parsed
prototype_instantiations_in_ilboolTemplate prototype instantiations included
il_has_all_prototype_instantiationsboolAll prototypes included (ALL_TEMPLATE_INFO_IN_IL=1)
il_has_C_semanticsboolC-language semantics active
nontag_types_used_in_exception_or_rttiboolNon-tag types in EH/RTTI
seq_number_lookup_entriesintegerCount of source sequence entries
target_configuration_indexintegerTarget configuration selector

After printing the header, display_il_header calls walk_file_scope_il (sub_60E4F0) with display_il_entry as the per-entity callback. This iterates every IL entry in file-scope region 1.

The Main Dispatcher: display_il_entry

display_il_entry (sub_5F4930, 1,686 lines) is the central per-entity display function. It receives an entry pointer and an entry kind byte, and dispatches to the appropriate per-kind display function.

Transparent (Inline) Kinds

The first switch handles kinds that are displayed inline by their parent and should not appear as standalone entries. These return immediately without output:

Transparent kinds (early return):
  4   routine_type_supplement      15  (reserved)
  5   routine_type_extra           19  try_supplement
  14  (reserved)                   20  asm_supplement
  27  template_parameter_suppl     34  (reserved)
  37  (reserved)                   38  (reserved)
  46  (reserved)                   47  (reserved)
  48  (reserved)                   75  (reserved)
  76  (reserved)                   77  (reserved)
  78  (reserved)                   81  (reserved)

Entry Header Line

For non-transparent kinds, the dispatcher prints a scope-annotated header:

file-scope type@7f3a4b200100
func-scope variable@7f3a4b300200

The scope prefix comes from dword_126FA30 (1 = "file-scope", 0 = "func-scope"). The kind name is looked up from off_E6DD80[kind_byte]. The address is the raw entry pointer value. For entries in function-scope regions while dword_126FA30 == 1, a warning "**NON FILE SCOPE PTR**" is emitted.

Dispatch Table

The second switch dispatches to specialized display functions:

KindHexNameDisplay functionLines
10x01source_file_entryinline in dispatcher~40
20x02constantsub_5F2720 (display_constant)605
30x03param_typeinline in dispatcher~30
60x06typesub_5F06B0 (display_type)1,033
70x07variablesub_5EE500 (display_variable)614
80x08fieldinline in dispatcher~80
90x09exception_specificationinline~20
100x0Aexception_spec_typeinline~10
110x0Broutinesub_5EF1A0 (display_routine)1,160
120x0Clabelinline~30
130x0Dexpr_nodesub_5ECFE0 (display_expr_node)534
160x10switch_case_entryinline~15
170x11switch_infoinline~10
180x12handlerinline~15
210x15statementsub_5EC600 (display_statement)328
220x16object_lifetimeinline~20
230x17scopesub_5F2140 (display_scope)177
280x1Cnamespaceinline~20
290x1Dusing_declarationinline~20
300x1Edynamic_initsub_5F37F0 (display_dynamic_init)248
310x1Flocal_static_variable_initinline~15
320x20vla_dimensioninline~10
330x21overriding_virtual_funcinline~15
350x23derivation_pathinline~10
360x24base_classinline~25
390x27class_infosub_5F4030 (display_class_supplement)366
410x29constructor_initinline~15
420x2Aasm_entryinline~25
430x2Basm_operandinline~15
440x2Casm_clobberinline~10
500x32source_sequence_entryinline~15
510x33full_entity_decl_infoinline~15
520x34instantiation_directiveinline~10
530x35src_seq_sublistinline~10
540x36explicit_instantiation_declinline~10
550x37orphaned_entitiesinline~10
560x38hidden_nameinline~10
570x39pragmainline~20
580x3Atemplateinline~20
590x3Btemplate_declinline~15
600x3Crequires_clauseinline~10
610x3Dtemplate_paraminline~15
620x3Ename_referencesub_5EBC60 (display_name_reference)84
630x3Fname_qualifierinline~15
640x40seq_number_lookupinline~10
650x41local_expr_node_refinline~10
660x42static_assertinline~10
670x43linkage_specinline~10
680x44scope_refinline~10
700x46lambdainline~15
710x47lambda_captureinline~15
720x48attributeinline~20
730x49attribute_argumentinline~10
740x4Aattribute_groupinline~10
790x4Ftemplate_infoinline~15
800x50subobject_pathinline~10
820x52module_infoinline~10
830x53module_declinline~10

Per-Kind Display Functions

source_file_entry (Kind 1)

Displayed inline in the dispatcher. Fields:

FieldTypeNotes
file_namestringShort file name
full_namestringFull path
name_as_writtenstringAs-written in #include
first_seq_numberintegerFirst source sequence number in this file
last_seq_numberintegerLast source sequence number
first_line_numberintegerFirst line number
child_filesIL pointer listIncluded files
is_implicit_includeboolImplicitly included
is_include_fileboolIs an #included file (not the primary TU)
top_level_fileboolTop-level compilation unit

source_corresp (Shared Prefix)

All named entities (variable, routine, type, field, label, namespace, template_param) share a source_corresp sub-record, printed by display_source_corresp (sub_5EDF40, 170 lines). This is the first thing displayed for each such entity:

source_corresp:
  name:                    foo
  unmangled_name_or_mangled_encoding: _Z3foov
  decl_position.seq:       42
  decl_position.column:    5
  name_references:         name_reference@7f3a...
  is_class_member:         TRUE
  access:                  public
  parent_scope:            file-scope scope@7f3a...
  enclosing_routine:       NULL
  referenced:              TRUE
  needed:                  TRUE
  name_linkage:            external

Fields displayed by display_source_corresp:

FieldTypeLookup table
namestringDirect string
unmangled_name_or_mangled_encodingstringDirect string
decl_positionpositionseq + column sub-fields
name_referencesIL pointername_reference entry
is_class_memberbool--
accessenumoff_A6F760 (4 entries: public/protected/private/none)
parent_scopeIL pointerScope entry
enclosing_routineIL pointerRoutine entry
referencedbool--
neededbool--
is_local_to_functionbool--
parent_via_local_scope_refIL pointer--
name_linkageenumoff_E6E040 (none/internal/external/C/C++)
has_associated_pragmabool--
is_decl_after_first_in_comma_listbool--
copied_from_secondary_trans_unitbool--
same_name_as_external_entity_in_secondary_trans_unitbool--
member_of_unknown_basebool--
qualified_unknown_base_memberbool--
marked_as_gnu_extensionbool--
is_deprecated_or_unavailablebool--
externalizedbool--
maybe_unusedbool[[maybe_unused]] attribute
source_sequence_entryIL pointer--
attributesIL pointerAttribute list

type (Kind 6)

display_type (sub_5F06B0, 1,033 lines) handles all 22 type kinds. After calling display_source_corresp, it prints common type fields then switches on the type kind byte at offset +132:

Common type fields:

FieldLookup table
nextIL pointer
based_typesLinked list, kind from off_A6F420 (6 entries)
sizeInteger
alignmentInteger
incompletebool
used_in_exception_or_rttibool
declared_in_function_prototypebool
alignment_set_explicitlybool
variables_are_implicitly_referencedbool
may_aliasbool
autonomous_primary_tag_declbool
is_builtin_va_listbool
is_builtin_va_list_from_cstdargbool
has_gnu_abi_tag_attributebool
in_gnu_abi_tag_namespacebool
type_kindEnum from off_A6FE40 (22 entries)

Type kind switch (offset +132):

KindNameKey sub-fields
2integerint_kind (via sub_5F9110), explicitly_signed, wchar_t_type, char8_t_type, char16_t_type, char32_t_type, bool_type; for enums: is_scoped_enum, packed, originally_unnamed, is_template_enum, ELF_visibility, base_type, assoc_template
3/4/5float/double/ldoublefloat_kind (via sub_5F93D0)
6pointertype_pointed_to, is_reference, is_rvalue_reference
7functionreturn_type, param_type_list, assoc_routine, has_ellipsis, prototyped, trailing_return_type, value_returned_by_cctor, does_not_return, result_should_be_used, is_const, explicit_calling_convention, calling_convention (from off_E6CDA0), this_class, qualifiers, ref_qualifiers, prototype_scope, exception_specification
8arrayelement_type, qualifiers, is_static, is_variable_size_array, is_vla, element_count, bound_constant
9/10/11class/struct/unionfield_list, extra_info (class supplement via sub_5F4030), final, abstract, any_virtual_base_classes, any_virtual_functions, originally_unnamed, is_template_class, is_specialized, is_empty_class, is_packed, max_member_alignment
12typereftyperef_type, template_arg_list, assoc_template, typeref_kind (from off_A6F640, 28 entries), qualifiers, predeclared, has_variably_modified_type, is_nonreal
13member pointerclass_of_which_a_member, type
14template paramkind (tptk_param/tptk_member/tptk_unknown), is_pack, is_generic_param, is_auto_param, class_type, coordinates
15vectorelement_type, size_constant, is_boolean_vector, vector_kind
16tupleelement_type, tuple_elements

variable (Kind 7)

display_variable (sub_5EE500, 614 lines) is one of the most field-heavy display functions. After display_source_corresp, it prints:

FieldLookup table / Notes
nextIL pointer
typeIL pointer
storage_classoff_A6FE00 (7 entries: none/auto/register/static/extern/mutable/thread_local)
declared_storage_classSame table
asm_name or regoff_A6F480 (53 register kind entries)
alignmentInteger
ELF_visibilityoff_A6F720 (5 entries)
init_priorityInteger
cleanup_routineIL pointer
container / bindingsSelected by bits at offset +162
sectionString (ELF section name)
aliased_variableIL pointer
declared_typeIL pointer
template_infoIL pointer

CUDA-specific variable fields:

FieldNotes
shared__shared__ memory space
constant__constant__ memory space
device__device__ memory space

Boolean flags (approximately 50 flags spanning bytes 144--208):

is_weak, is_weakref, is_gnu_alias, has_gnu_used_attribute, has_gnu_abi_tag_attribute, is_not_common, is_common, has_internal_linkage_attribute, asm_name_is_valid, used, address_taken, is_parameter, is_parameter_pack, is_pack_element, is_enhanced_for_iterator, initializer_in_class, constant_valued, is_thread_local, extends_lifetime, is_template_param_object, compiler_generated, is_in_class_specialization, is_handler_param, is_this_parameter, referenced_non_locally, modified_within_try_block, is_template_variable, is_prototype_instantiation, is_nonreal, is_specialized, specialized_with_old_syntax, explicit_instantiation, class_explicitly_instantiated, explicit_do_not_instantiate, param_value_has_been_changed, param_used_more_than_once, is_anonymous_parent_object, is_member_constant, is_constexpr, declared_constinit, is_inline, suppress_inline_definition, superseded_external, has_variably_modified_type, is_vla, is_compound_literal, has_explicit_initializer, has_parenthesized_initializer, has_direct_braced_initializer, has_flexible_array_initializer, declared_with_auto_type_specifier, declared_with_decltype_auto, declared_with_class_template_placeholder

routine (Kind 11)

display_routine (sub_5EF1A0, ~1,160 lines) is the single largest per-kind display function. After display_source_corresp:

FieldLookup table / Notes
nextIL pointer
typeIL pointer (function type)
function_def_numberInteger
memory_regionInteger (region index for function body)
storage_classoff_A6FE00 (7 entries)
declared_storage_classSame table
special_kindoff_A6FC00 (13 entries: none/constructor/destructor/conversion/operator/lambda_call_operator/...)
opname_kindoff_A6FC80 (47 entries)
builtin_function_kindInteger
ELF_visibilityoff_A6F720
virtual_function_numberInteger
constexpr_intrinsic_numberInteger
sectionString
aliased_routineIL pointer
inline_partnerIL pointer
ctor_priority / dtor_priorityInteger
asm_nameString
declared_typeIL pointer
generating_using_declIL pointer
befriending_classesIL pointer list
assoc_templateIL pointer
template_arg_listVia display_template_arg_list

CUDA-specific routine flags (byte 182):

FlagBitMeaning
nvvm_intrinsicbit 4NVVM intrinsic function
devicebit 5__device__ execution space
globalbit 6__global__ execution space
hostbit 4 (byte 183)__host__ execution space

C99-specific fields (displayed when dword_126EBA8 == 1 and std_version > 199900):

fp_contract, fenv_access, cx_limited_range -- pragma state values from off_A6F460 (4 entries).

Boolean flags (approximately 60 flags spanning bytes 176--191):

address_taken, is_virtual, overrides_base_member, pure_virtual, final, override, covariant_return_virtual_override, is_inline, is_declared_constexpr, is_constexpr, is_constexpr_intrinsic, compiler_generated, defined, called, is_explicit_constructor, is_explicit_conversion_function, is_trivial_default_constructor, is_trivial_copy_function, is_trivial_destructor, is_initializer_list_ctor, is_delegating_ctor, is_inheriting_ctor, assignment_to_this_done, is_prototype_instantiation, is_template_function, is_specialized, specialized_with_old_syntax, explicit_instantiation, class_explicitly_instantiated, explicit_do_not_instantiate, has_nodiscard_attribute, never_throws, is_in_class_specialization, never_inline, is_pure, is_initialization_routine, is_finalization_routine, is_weak, is_weakref, is_gnu_alias, is_ifunc, has_gnu_used_attribute, has_gnu_abi_tag_attribute, in_gnu_abi_tag_namespace, allocates_memory, no_instrument_function, no_check_memory_usage, always_inline, gnu_c89_inline, implicit_alias, has_internal_linkage_attribute, contains_try_block, contains_local_class_type, superseded_external, defined_in_friend_decl, contains_statement_expression, inline_in_class_definition, is_lambda_body, is_defaulted, is_deleted, contains_local_static_variable, is_raw_literal_operator, is_tls_init_routine, has_deducible_return_type, has_deduced_return_type, contains_generic_lambda, is_coroutine, is_top_level_in_mem_region, friend_defined_in_instantiation, is_ineligible, definition_needed, defined_outside_of_parent, trailing_requires_clause

expr_node (Kind 13)

display_expr_node (sub_5ECFE0, 534 lines) handles 36 expression node kinds. Common expression fields are printed first:

FieldNotes
typeIL pointer (expression type)
orig_lvalue_typeIL pointer
nextIL pointer
is_lvaluebool
is_xvaluebool
result_is_not_usedbool
is_pack_expansionbool
is_parenthesizedbool
compiler_generatedbool
volatile_fetchbool
do_not_interpretbool
type_definition_neededbool

Expression kind switch (offset +24):

KindNameKey sub-fields
0enk_error(none)
1enk_operationoperation.kind from off_A6F840 (120 operator kinds), operation.type_kind from off_A6FE40 (22 type kinds), 20+ boolean flags for cast semantics, ADL suppression, virtual call properties, evaluation order
2enk_constantConstant value reference
3enk_variableVariable reference
4enk_fieldField access
5enk_temp_initTemporary initialization
6enk_lambdaLambda expression
7enk_new_deleteis_new, placement_new, aligned_version, array_delete, global_new_or_delete, deducible_type, type, routine, arg, dynamic_init, number_of_elements
8enk_throwThrow expression
9enk_conditionConditional expression
10enk_object_lifetimeObject lifetime tracking
11enk_typeidtypeid expression
12enk_sizeofsizeof expression
13enk_sizeof_packsizeof... (pack)
14enk_alignofalignof expression
15enk_datasizeof__datasizeof
16enk_address_of_ellipsisAddress of ...
17enk_statementStatement expression
18enk_reuse_valueValue reuse
19enk_routineFunction reference
20enk_type_operandType operand
21enk_builtin_operationBuilt-in op from off_E6C5A0
22enk_param_refParameter reference
23enk_braced_init_listBraced initializer
24enk_c11_generic_Generic selection
25enk_builtin_choose_expr__builtin_choose_expr
26enk_yieldco_yield
27enk_awaitco_await
28enk_foldFold expression
29enk_initializerInitializer
30enk_concept_idConcept ID
31enk_requiresrequires expression
32enk_compound_reqCompound requirement
33enk_nested_reqNested requirement
34enk_const_eval_deferredConsteval deferred
35enk_template_nameTemplate name

Every expression case ends with dump_source_position("position", ...) to record the source location.

statement (Kind 21)

display_statement (sub_5EC600, 328 lines) handles 26 statement kinds. Common fields first:

FieldNotes
positionSource position
nextIL pointer
parentIL pointer (enclosing scope/block)
attributesIL pointer
has_associated_pragmabool
is_initialization_guardbool
is_lowering_boilerplatebool
is_fallthrough_statementbool
is_likelybool
is_unlikelybool

Statement kind switch (offset +32):

KindNameKey sub-fields
0stmk_exprExpression statement
1stmk_ifif
2stmk_constexpr_ifif constexpr
3stmk_if_constevalif consteval (C++23)
4stmk_if_not_constevalif !consteval
5stmk_whilewhile loop
6stmk_gotogoto
7stmk_labellabel
8stmk_returnreturn
9stmk_coroutineCoroutine body (see below)
10stmk_coroutine_returnCoroutine return
11stmk_blockBlock/compound: statements, final_position, assoc_scope, lifetime, end_of_block_reachable, is_statement_expression
12stmk_end_test_whiledo-while
13stmk_forfor loop
14stmk_range_based_forRange-for: iterator, range, begin, end, ne_call_expr, incr_call_expr
15stmk_switch_caseswitch case
16stmk_switchswitch
17stmk_initInitialization
18stmk_asmInline assembly
19stmk_try_blocktry block
20stmk_declDeclaration
21stmk_set_vla_sizeVLA size
22stmk_vla_declVLA declaration
23stmk_assigned_gotoComputed goto
24stmk_emptyEmpty statement
25stmk_stmt_expr_resultStatement expression result

Coroutine statement (case 9) displays the full C++20 coroutine lowering structure:

traits, handle, promise, init_await_resume, this_param_copy,
paramter_copies, final_suspend_label, initial_suspend_call,
final_suspend_call, unhandled_exception_call, get_return_object_call,
new_routine, delete_routine, ...

The field name "paramter_copies" (missing the 'e' in "parameter") is a typo preserved verbatim from the EDG source. This confirms the display strings originate from Edison Design Group's own il_to_str.c -- a reimplementation would spell it correctly.

scope (Kind 23)

display_scope (sub_5F2140, 177 lines) handles 9 scope kinds:

KindNameExtra fields
0sck_fileTop-level file scope
1sck_func_prototypeFunction prototype scope
2sck_blockassoc_handler
3sck_namespaceassoc_namespace
6sck_class_struct_unionassoc_type
8sck_template_declarationTemplate declaration scope
15sck_conditionassoc_statement
16sck_enumassoc_type
17sck_functionroutine.ptr, parameters, constructor_inits, lifetime_of_local_static_vars, this_param_variable, return_value_variable

Common scope fields: next, parent, kind

Boolean flags: do_not_free_memory_region, is_constexpr_routine, is_stmt_expr_block, is_placeholder_scope, needed_walk_done

Child entity lists: assoc_block, lifetime, constants, types, variables, nonstatic_variables, labels, routines, asm_entries, scopes

Conditional lists (controlled by bitmask tests on scope kind):

// Bitmask 0x20044 = bits 2+6+17 = sck_block + sck_class_struct_union + sck_function
// Bitmask 0x9     = bits 0+3    = sck_file + sck_namespace

if ((1LL << kind) & 0x20044) {
    // display: namespaces, using_declarations, using_directives
}
if ((1LL << kind) & 0x9) {
    // display: namespaces, using_declarations, using_directives
}
// Also: dynamic_inits, local_static_variable_inits (function/block scopes)
//       expr_node_refs, scope_refs, vla_dimensions (function scope + C mode)
//       pragmas, hidden_names, templates, source_sequence_list, src_seq_sublist_list

constant (Kind 2)

display_constant (sub_5F2720, 605 lines) handles 16 constant kinds. After display_source_corresp, common fields include next, type, orig_type, expr, and approximately 25 boolean flags.

Constant kind switch (offset +148):

KindNameKey sub-fields
0ck_error(none)
1ck_integerInteger value via sub_602F20
2ck_stringcharacter_kind (char/wchar_t/char8_t/char16_t/char32_t), length, literal_kind (see below)
3ck_floatFloat value via sub_5FCAF0
4ck_complexComplex value
5ck_imaginaryImaginary value
6ck_addressSub-kind: abk_routine/variable/constant/temporary/uuidof/typeid/label; subobject_path, offset
7ck_ptr_to_membercasting_base_class, name_reference, cast_to_base, is_function_ptr
8ck_label_differencefrom_address, to_address
9ck_dynamic_initdynamic_init pointer
10ck_aggregatefirst_constant, last_constant, has_dynamic_init_component
11ck_init_repeatconstant, count, multidimensional_aggr_tail_not_repeated
12ck_template_paramSub-kinds: tpck_param/expression/member/unknown_function/address/sizeof/datasizeof/alignof/uuidof/typeid/noexcept/template_ref/integer_pack/destructor
13ck_designatoris_field_designator, is_generic, uses_direct_init_syntax
14ck_void(none)
15ck_reflectionentity, local_scope_number

dynamic_init (Kind 30)

display_dynamic_init (sub_5F37F0, 248 lines) handles 9 dynamic initialization kinds:

KindNameKey sub-fields
0dik_none(none)
1dik_zeroZero-initialization
2dik_constantConstant initialization
3dik_expressionExpression initialization
4dik_class_result_via_ctorClass result through constructor
5dik_constructorroutine, args, is_copy_constructor_with_implied_source, is_implicit_copy_for_copy_initialization, value_initialization
6dik_nonconstant_aggregateNon-constant aggregate
7dik_bitwise_copysource
8dik_lambdalambda, constant, non_constant

Common fields: next, variable, destructor, lifetime, next_in_destruction_list, unordered, init_expr_lifetime, and approximately 20 boolean flags including static_temp, follows_an_exec_statement, inside_conditional_expression, has_temporary_lifetime, is_constructor_init, is_freeing_of_storage_on_exception, overlaps_temps_in_inner_lifetime, is_reused_value, is_creation_of_initializer_list_object, master_entry.

class_info (Kind 39)

display_class_type_supplement (sub_5F4030, 366 lines) is not dispatched directly from the kind table but called by display_type when the type kind is class/struct/union (kinds 9/10/11). It prints the class supplement record:

FieldNotes
base_classesIL pointer list
direct_base_classesIL pointer list
preorder_base_classesIL pointer list
primary_base_classIL pointer
size_without_virtual_base_classesInteger
alignment_without_virtual_base_classesInteger
highest_virtual_function_numberInteger
virtual_function_info_offsetInteger
virtual_function_info_base_classIL pointer
ELF_visibilityoff_A6F720
is_lambda_closure_classbool
is_generic_lambda_closure_classbool
has_lambda_conversion_functionbool
is_initializer_listbool
has_initializer_list_ctorbool
has_anonymous_union_memberbool
anonymous_union_kindenum (auk_none/auk_variable/auk_field)
is_va_list_tagbool
has_nodiscard_attributebool
has_field_initializerbool
removed_from_ilbool
contains_errorbool
befriending_classesLinked list (checks kind bytes 9/10/11 for class/struct/union)
friend_routinesIL pointer list
friend_classesIL pointer list
assoc_scopeIL pointer
assoc_templateIL pointer
template_arg_listVia display_template_arg_list
lambda_parent.variable / .field / .routineSelected by bits in byte 86
proxy_of_typeIL pointer

Formatting Infrastructure

25-Column Field Labels

dump_field_label (sub_5EB2A0, 22 lines) is the universal field label formatter. It prints "field_name:" then pads with spaces to column 25. If the label plus colon exceeds 24 characters, it prints a newline first to avoid misalignment:

storage_class:           static
alignment:               16
is_constexpr:            TRUE

This produces the consistent columnar output visible in all IL dumps.

Boolean Fields

dump_field_bool (sub_5EB450, 25 lines) prints a label and "TRUE" or "FALSE":

is_virtual:              TRUE
pure_virtual:            FALSE

Source Position Fields

dump_source_position (sub_5EB4E0, 82 lines) prints position as two sub-fields when the position is non-zero (seq != 0 or column != 0):

position.seq:            42
position.column:         5

Reads a 32-bit sequence number at *position and a 16-bit column at *(position + 4).

IL Pointer Annotations

dump_il_entity_pointer (sub_5EB8B0, 99 lines) is the most comprehensive pointer formatter. For each IL entity pointer, it prints:

  1. Scope prefix: "file-scope" or "func-scope" (from bit 0 of the entry prefix byte at entry_ptr - 8)
  2. Kind name: from off_E6DD80[kind_byte]
  3. Hex address: @%lx
  4. Entity name (kind-dependent):
    • Kinds with name at offset +8 (bitmask 0x2000000010001984): prints the name string
    • Kind 12 (label): prints "label " prefix + name
    • Kind 6 (type): calls qualified name formatter
    • Kind 2 (constant): calls type display
    • Kind 0x40 (seq_number_lookup): prints qualified name from offset +0
    • Kind with bit 36 set: prints qualified name from offset +40, plus "in" context from +56
primary_source_file:     file-scope source_file_entry@7f3a4b100020 "test.cu"
main_routine:            file-scope routine@7f3a4b200100 "main"

The variant dump_il_string_pointer (sub_5EB670) prints the same format but includes the string value from the pointed-to entry. A scope mismatch (e.g., function-scope pointer found during file-scope display) triggers a "**NON FILE SCOPE PTR**" warning.

Entity List Display

display_entity_list (sub_5EC450, 87 lines) walks a linked list of entity pointers and prints each with scope/kind/address annotations:

entities:                file-scope variable@7f3a... "x"
                         func-scope variable@7f3a... "y"

It follows the next link at offset 0 of each list node until NULL.

String Literal Display

dump_string_value (sub_5EB300, 41 lines) prints string values with proper escape handling:

  • NULL pointers print "NULL"
  • Non-printable characters are printed as octal escapes (\OOO)
  • Backslash and double-quote are backslash-escaped (\\, \")
  • The octal mask width is controlled by dword_126E49C (CHAR_BIT equivalent, typically 8)
file_name:               "test.cu"
full_name:               "/home/user/project/test.cu"

Float Constant Formatting

form_float_constant (sub_5F7FD0, 302 lines) handles float-to-string conversion with EDG-specific formatting. An assertion at line 6175 guards against buffer overflow (63-byte limit).

Float kind suffixes:

KindSuffixType
0(none)double
2f/Ffloat
3f32xextended float32
5f64xextended float64
6l/Llong double
7wfloat128/wide
8qquad
9bf16bfloat16
10f16float16
11f32float32
12f64float64
13f128float128

Special value handling:

  • NaN: __builtin_nanf(""), __builtin_nan(""), etc. (when compiler version > 30299)
  • Infinity: __builtin_huge_valf() or (__extension__ 0x1.0p<exp>f)
  • Division form: (f/0.0f) or (f/(0,0.0f)) (C++ vs C modes, selected by dword_126E1D8/dword_126E1E8)
  • User-defined literals: (funcname("string_value")) form

Data Tables Referenced

The display subsystem relies on approximately 20 string-to-enum lookup tables in the .rodata segment:

AddressNameEntriesUsed by
off_A6F000attr_arg_kind_names6Attribute argument display
off_A6F040attr_location_names24Attribute display
off_A6F100attr_family_names5Attribute display
off_A6F140attr_kind_names86Attribute display
off_A6F3F0class_kind_labels3befriending_classes display
off_A6F420based_type_kind_names6display_type based_types
off_A6F460pragma_state_names4fp_contract/fenv_access/cx_limited_range
off_A6F480register_kind_names53display_variable reg field
off_A6F640typeref_kind_names28display_type typeref
off_A6F720elf_visibility_kind_names5ELF visibility (all entity types)
off_A6F760access_specifier_names4public/protected/private/none
off_A6F840expr_operator_kind_names120display_expr_node operations
off_A6FC00special_function_kind_names13display_routine special_kind
off_A6FC80operator_name_kind_names47display_routine opname_kind
off_A6FE00storage_class_names7Storage class (variable + routine)
off_A6FE40type_kind_names22Type kind (all type displays)
off_E6C5A0builtin_operation_namesvariesdisplay_expr_node builtins
off_E6CDA0calling_convention_namesvariesdisplay_type calling conventions
off_E6CDE0pragma_kind_namesvariesPragma display
off_E6CF40asm_clobber_reg_namesvariesAsm clobber display
off_E6D240token_kind_namesvariesFold expression / attribute_arg tokens
off_E6DD80il_entry_kind_names~84All display functions (entry kind)
off_E6E040linkage_kind_namesvariesName linkage (source_corresp)

All tables use the same bounds-checking pattern:

const char *name = "**BAD STORAGE CLASS**";
if ((unsigned char)value <= 6u)
    name = storage_class_names[value];
puts(name);

Out-of-range values produce "**BAD <KIND>**" sentinel strings, which serve as diagnostic markers for corrupted IL.

Global State

AddressNameTypePurpose
dword_126FA30is_file_scope_regionint1 during file-scope display, 0 during function-scope
qword_126F980output_callbacksfunction ptrOutput function (default: sub_5EB290 = fputs(s, stdout))
byte_126FA16display_activebyteSet to 1 during display, prevents re-entrant calls
byte_126FA11pcc_compat_shadowbyteShadow of PCC compatibility mode during display
dword_126EBA8source_languageint0 = C++, 1 = C
dword_126EBACstd_versionintC/C++ standard version number
dword_126EC80total_region_countintNumber of memory regions (1 = file scope only)
qword_126EC88region_tablepointer arrayRegion index to memory block mapping
qword_126EB90scope_tablepointer arrayRegion index to scope entry mapping
qword_126EEE0source_file_namestring ptrName of the source file being compiled

Helper Functions (0x5F8000--0x6039E0)

The display subsystem includes approximately 50 additional helper functions in the address range beyond the main dispatchers:

AddressLinesIdentityPurpose
sub_5F85E078display_bool_fieldBoolean TRUE/FALSE output
sub_5F876097display_flags_wordFlags word display
sub_5F891088display_type_qualifiersconst/volatile/restrict qualifier flags
sub_5F8A8049display_storage_classStorage class enum
sub_5F8BD0139display_access_specifierAccess with indentation
sub_5F8DF0103display_linkage_kindLinkage kind enum
sub_5F904028init_output_contextInitialize display callback state
sub_5F9110149display_int_type_kindInteger type kind name
sub_5F93D070display_float_type_kindFloat type kind name
sub_5F950070display_int_type_sizeInteger type size name
sub_5F965099display_qualifier_flagsFull qualifier flags
sub_5F982018display_ref_qualifier& or &&
sub_5F986091display_calling_conventionCalling convention from off_E6CDA0
sub_5F99A0115display_attribute_targetAttribute target kind
sub_5F9BC020display_asm_keyword"asm" or "volatile"
sub_5F9C1026display_elaborated_typeElaborated type specifier
sub_5F9CA050display_struct_layoutStructure layout padding mode
sub_5F9D8089display_member_alignmentMember alignment field
sub_5F9F7057display_template_kindTemplate kind name
sub_5FA0D0283display_template_arg_listFull template argument list
sub_5FA660127display_constraint_exprConstraint expression (C++20)
sub_5FA8F0118display_deduction_guideDeduction guide info
sub_5FAB70333display_capture_listLambda capture list
sub_5FB270556display_expr_operator_nameExpression operator name (120 kinds)
sub_5FBCD0571display_expr_detailsOperator-specific expression details
sub_5FCAF01,319display_float_constantFloat/complex/imaginary formatting
sub_5FE7C055display_expr_flagExpression flag display
sub_5FE8B01,659display_expr_operatorExpression operator details (2nd largest)
sub_60074072display_for_rangeRange-based for details
sub_600870171display_coroutine_infoCoroutine info (C++20)
sub_600BF019display_designated_initDesignated initializer
sub_600C50107display_attribute_entryAttribute entry
sub_600E0055display_asm_operandAsm operand display
sub_600EF076display_asm_statementAsm statement details
sub_600FF029display_gcc_builtin_kindGCC built-in kind
sub_60107087display_pragma_infoPragma info
sub_6011F0155display_declspec_attribute__declspec attribute
sub_60146092display_thread_localThread-local info
sub_6015A073display_module_infoModule info (C++20)
sub_6016F0197display_concept_requiresConcept/requires expression
sub_601B1048display_pack_expansionPack expansion info
sub_601BE050display_structured_bindingStructured binding (C++17)
sub_601CB0562display_additional_exprAdditional expression info
sub_6027D0144display_deduced_classDeduced class info
sub_6029B0190display_decl_sequenceDeclaration sequence entry
sub_602DC074display_enum_underlyingEnum underlying type
sub_602F20306display_integer_constantInteger constant formatting
sub_603670134display_vendor_attributeVendor attribute details
sub_6038F026display_cleanup_handlerCleanup handler
sub_6039E078display_sequence_entryLast function in il_to_str region

The "paramter_copies" Typo

The coroutine statement display (case 9 in display_statement) prints the field label "paramter_copies" -- missing the 'e' in "parameter." This typo is present in the compiled binary's string table and originates from Edison Design Group's source code. It serves as strong provenance evidence: a clean-room reimplementation would not reproduce this exact spelling error, confirming that cudafe++ links genuine EDG il_to_str.c object code.

Complete Call Graph

display_il_file (sub_5F7DF0) ─── TOP LEVEL
├── display_il_header (sub_5F76B0)
│   ├── init_output_context (sub_5F9040)
│   ├── dump_il_entity_pointer (sub_5EB8B0) ×30+ for header fields
│   ├── dump_field_bool (sub_5EB450) ×15+ for header booleans
│   ├── dump_string (sub_5EB790)
│   └── walk_file_scope_il (sub_60E4F0)
│       └── display_il_entry (sub_5F4930) ─── callback per entity
│
└── [loop over regions 2..N]
    └── walk_routine_scope_il (sub_610200)
        └── display_il_entry (sub_5F4930) ─── callback per entity

display_il_entry (sub_5F4930) ─── MAIN DISPATCHER
├── display_source_corresp (sub_5EDF40) ─── shared by named entities
├── display_statement (sub_5EC600) ─── case 0x15
│   ├── display_coroutine_info (sub_600870)
│   └── display_for_range (sub_600740)
├── display_expr_node (sub_5ECFE0) ─── case 0x0D
│   ├── display_expr_operator (sub_5FE8B0)
│   ├── display_expr_operator_name (sub_5FB270)
│   └── display_expr_details (sub_5FBCD0)
├── display_variable (sub_5EE500) ─── case 0x07
│   └── display_init_kind (sub_5EBB50)
├── display_routine (sub_5EF1A0) ─── case 0x0B
│   └── display_template_arg_list (sub_5EBF60 / sub_5FA0D0)
├── display_type (sub_5F06B0) ─── case 0x06
│   ├── display_class_supplement (sub_5F4030)
│   ├── display_int_type_kind (sub_5F9110)
│   └── display_float_type_kind (sub_5F93D0)
├── display_scope (sub_5F2140) ─── case 0x17
├── display_constant (sub_5F2720) ─── case 0x02
│   ├── display_integer_constant (sub_602F20)
│   └── display_float_constant (sub_5FCAF0)
├── display_dynamic_init (sub_5F37F0) ─── case 0x1E
├── display_name_reference (sub_5EBC60) ─── case 0x3E
└── display_entity_list (sub_5EC450) ─── multiple cases

display_single_entity (sub_5F7D50) ─── TARGETED DISPLAY
├── entity_lookup (sub_73D400)
├── resolve_entity (sub_7377D0)
├── get_entity_kind (sub_5C64C0)
├── init_output_context (sub_5F9040)
└── display_il_entry (sub_5F4930)

Relationship to Other Subsystems

The IL display subsystem is read-only: it never modifies the IL graph. It shares the same entry walker functions used by the IL Tree Walking framework (walk_file_scope_il = sub_60E4F0, walk_routine_scope_il = sub_610200) and the Keep-in-IL mark phase, but passes display_il_entry as the callback instead of a transformation function.

The IL Allocation subsystem provides dump_il_table_stats (sub_5E99D0), which dumps allocation counters rather than IL content -- a complementary diagnostic activated separately.

The field offsets printed by the display functions serve as ground truth for the IL Overview entry kind table and the Entity Node Layout documentation.