Scope Entry

The scope entry is the 784-byte record that forms the elements of the scope stack, the central data structure in cudafe++ for tracking nested lexical scopes during C++ parsing and semantic analysis. The scope stack is a flat array at qword_126C5E8, indexed by dword_126C5E4 (current depth). Every time the parser enters a new scope -- file, block, function body, class definition, template declaration, namespace -- a new 784-byte entry is pushed onto this stack. When the scope closes, the entry is popped and all associated cleanup runs: symbol table housekeeping, using-directive deactivation, name collision discriminator assignment, template parameter restoration, and memory region disposal.

This page documents the scope stack entry layout, the scope kind enum, the key flag bytes, the CUDA-specific additions (device/host scope context), the template instantiation depth counters, and the major push/pop functions.

Key Facts

Scope Stack Global Variables

Full Offset Map

The table documents every field observed in the 784-byte scope stack entry. These are scope stack entry fields, not IL scope node fields (the IL scope is a separate 288-byte structure pointed to from offset +192).

Scope Stack Kind Enum

The scope stack kind byte at +4 uses a different, larger enum than the IL scope kind (sck_*) at IL scope node +28. The scope stack enum includes additional entries for reactivation states, template instantiation context, and module scopes. The mapping is derived from scope_kind_to_string (sub_7000E0, 77 lines) which contains display string literals for each enum value, and from display_scope (sub_5F2140) in il_to_str.c.

Scope Stack Kind Values

Relationship to IL Scope Kinds

The IL scope node (288 bytes, allocated by alloc_scope at sub_5E7D80) uses a smaller sck_* enum at its +28 field. The scope stack entry at +192 points to the IL scope that persists after the stack entry is popped. Not all scope stack kinds produce an IL scope -- reactivation kinds (5, 7) and context kinds (9, 10) reuse existing IL scopes.

CUDA-Specific Fields

NVIDIA added two device/host scope tracking bits to the scope entry, grafted onto the EDG base structure.

Byte +6, Bit 1: Device Scope Context

scope_entry+6, bit 1 (0x02):

  When set: code in this scope is compiled for the device execution space.
  When clear: code in this scope is compiled for the host.

This bit is tested by CUDA-specific code paths to determine whether the current compilation context targets device or host. It affects:

• Whether __device__-only functions suppress certain diagnostics (e.g., missing return value warning at check_void_return_okay, sub_719D20)
• Whether device-specific type validation applies
• Severity overrides via byte_126ED55 (default diagnostic severity for device mode)

The bit is set when entering __device__ or __global__ function scopes and cleared when entering __host__ scopes. This allows mixed host/device compilation to track which context is active at any nesting depth.

Byte +14, Bit 2: In Device Code

scope_entry+14, bit 2 (0x04):

  Secondary device-code marker. Set when the parser is inside a device
  function body. Used in conjunction with dword_106C2C0 (CUDA device
  compilation mode flag).

Template Instantiation Depth Counters

Three fields at offsets +576, +580, and +584 form the template instantiation depth tracking system. These fields enable the scope stack to handle nested template instantiations without fully pushing/popping scope entries at every nesting level.

Mechanism

When push_template_instantiation_scope (sub_709DE0) sets up a template instantiation, it writes the current scope stack depth into +580 (orig_depth) and the saved global depth into +584 (saved_scope_depth). The +576 counter starts at 0.

If the same template scope is re-entered for a nested instantiation (e.g., recursive template), +576 is incremented rather than pushing a full new scope entry. On pop, pop_template_instantiation_scope (sub_708EE0) checks +576:

if (scope_entry[576] > 0) {
    scope_entry[576]--;     // just decrement, don't pop
    return;
}
// Full pop: restore scope stack to orig_depth
pop_scopes_to(scope_entry[580]);
restore(dword_126C5AC, scope_entry[584]);

This optimization avoids deep scope stack growth during deeply recursive template instantiations (e.g., std::tuple<T1, T2, ..., TN> with large N).

Validation

pop_template_instantiation_scope_with_check (sub_708E90) validates that +576 matches the expected depth before calling the actual pop. The assertion is at scope_stk.c line 5593. A mismatch triggers an internal error.

Push Scope: push_scope_full (sub_700560)

The core scope push function (1476 lines, 13 parameters, located at 0x700560). Called directly or via thin wrappers for each scope kind.

Parameters

The 13-parameter signature handles all scope kinds through a single entry point:

1. Scope kind
2. Associated entity pointer (class type, namespace entity, routine descriptor, etc.)
3. Region number
4. Additional flags 5-13. Kind-specific parameters (template info, reactivation data, etc.)

Key Operations

1. Stack growth: Increments dword_126C5E4. If the stack exceeds its allocation, it is reallocated (the base pointer qword_126C5E8 may change).

2. Entry initialization: Zeros the 784-byte entry, then sets:

   • +0 = scope number (from a global counter)
   • +4 = scope kind
   • +240 = region number
   • +192 = IL scope pointer (newly allocated via alloc_scope or reused from a reactivated entity)
   • +560 = parent scope index

3. Kind-specific setup:

   • File (0): Sets dword_126C5DC, initializes file-level state.
   • Block (2): Links to enclosing function scope.
   • Namespace (3, 4, 5): Sets +232 to namespace entity. For extensions (4), reuses existing IL scope. For reactivation (5), calls add_active_using_directives_for_scope.
   • Class (6, 7): Sets +216 to class type pointer, dword_126C5C8 to current index. For reactivation (7), walks the class hierarchy to restore template context.
   • Template declaration (8): Sets template-related bits in +8.
   • Function (17): Sets dword_126C5D8, qword_126C5D0, +224.

4. Parent scope linkage: Calls set_parent_scope_on_push to establish the scope tree.

5. Memory region: Calls get_enclosing_memory_region to determine the memory arena for allocations within this scope.

Push Wrappers

Pop Scope: pop_scope (sub_7076A0)

The core scope pop function (1142 lines, at 0x7076A0). Complement to push_scope_full. Performs all scope cleanup in a specific order.

Cleanup Sequence

1. Debug trace: If dword_126EFC8 is set, prints "pop_scope: number = %ld, depth = %d".

2. Scope wrapup: Calls wrapup_scope (sub_706710, 381 lines) which:

   • Iterates all symbols in the scope
   • Runs end_of_scope_symbol_check (sub_705440, 781 lines) for consistency validation
   • Emits needed definitions
   • Reports unreferenced entities

3. Using-directive deactivation: Clears active using-directives for this scope via sub_6FEC10 (debug: clear using-directive).

4. Template parameter restoration: If leaving a template scope, calls restore_default_template_params (sub_6FEEE0) to undo template parameter symbol bindings.

5. Name collision discriminators: Assigns ABI discriminator values to entities with the same name in this scope via assign_discriminators_to_entities_list (sub_7036E0).

6. C99 inline definitions: Checks check_c99_inline_definition (sub_703AD0) for C99-mode inline function rules.

7. Module/pragma state: Adjusts STDC pragma state (byte_126E558/559/55A) and module context if applicable.

8. Stack decrement: Decrements dword_126C5E4. Restores previous scope's global state (function scope index, class scope index, etc.).

9. Memory region disposal: Frees the memory arena associated with this scope if the scope kind has one.

Pop Variants

Template Instantiation Scope

The template instantiation scope push and pop are separate from the generic scope push/pop. They handle the complex process of binding template parameters to arguments, setting up the correct translation unit context, and managing pack expansions.

Push: push_template_instantiation_scope (sub_709DE0)

The largest function in the scope_stk.c range at 1281 lines. Takes 8 parameters: template pointer, association info, and various flags.

Key operations:

1. Translation unit check: Calls sub_7418D0 to verify that the template being instantiated belongs to the same translation unit, or that cross-TU instantiation is explicitly allowed (flag & 0x1000). Failure triggers the assert "push_template_instantiation_scope: wrong translation unit".

2. Template parameter binding: Iterates the template parameter list and the instantiation argument list in parallel, creating bindings. For each template parameter:

   • Type parameters: binds to the supplied type argument
   • Non-type parameters: binds to the supplied expression/value
   • Template template parameters: binds to the supplied template

3. Pack expansion: For variadic templates, handles parameter pack expansion. Creates pack instantiation descriptors via create_pack_instantiation_descr (sub_70CF50, 772 lines).

4. Scope entry setup: Writes +576 = 0, +580 = current depth, +584 = dword_126C5AC. Sets +216 to the translation unit pointer.

5. State save/restore: Saves dword_126C5B8 (is_member_of_template), dword_126C5D8 (function scope), qword_126C5D0 (routine descriptor).

6. Reactivation flags: Flag bits & 0x84000 control class template reactivation behavior. When set, the function enters class reactivation mode via sub_70BB60.

Pop: pop_template_instantiation_scope (sub_708EE0)

66 lines. Reverse of the push.

1. Reads +576 (depth counter). If > 0, decrements and returns early (nested instantiation shortcut).
2. If bit 0 of +9 is set (needs_cleanup), calls sub_67B4E0() to clean up template instantiation artifacts.
3. Pops scope entries back to orig_depth (+580) via sub_7076A0.
4. Restores dword_126C5AC from +584.
5. Calls sub_6FED20 (debug trace: set using-directive).

Related Functions

Using-Directive Activation

When entering a namespace scope that has active using namespace directives, those directives must be reactivated so that names from the nominated namespace are visible. The scope stack manages this through two functions:

add_active_using_directives_for_scope (sub_6FFCC0)

246 lines. Called during scope push when entering a namespace or block that may have inherited using-directives. Walks the using-directive list for the scope and calls add_active_using_directive_to_scope for each one.

Debug trace format: "adding using-dir at depth %d for namespace %s applies at %d".

Using-Directive Debug Traces

Name Collision Discriminators

When multiple local entities share the same name (e.g., two struct S in different blocks within the same function), the Itanium ABI requires a discriminator suffix in the mangled name. The scope stack manages this through:

Class and Template Reactivation

When defining an out-of-line member function (void MyClass::foo() { ... }), the parser must reactivate the class scope so that class member names are visible. For class templates, this also requires reactivating the template instantiation scope.

reactivate_class_context (sub_7029D0 / sub_70BE50)

Two implementations exist:

• sub_7029D0 (196 lines, in p1.16): Reactivates a class scope for out-of-line definition. Asserts "reactivate_class_context: class type has NULL assoc_info".
• sub_70BE50 (130 lines, in p1.17): Additional variant that handles nesting, template flags, and scope_entry +8 bit manipulation.

push_class_and_template_reactivation_scope_full (sub_7098B0)

261 lines. Handles the combined case of class template reactivation. Reads symbol flags at offsets +80, +81, +161, +162. Processes "specified template decl info" at +64 of assoc_info. Detects member templates via bit 0x10 at +81. When dword_106BC58 is set, enters class reactivation mode with sub_70BB60.

reactivate_local_context (sub_702670 / sub_70C0F0)

• sub_702670 (120 lines): Reactivates a previously saved local scope context. Calls push_scope_full.
• sub_70C0F0 (50 lines): Asserts "reactivate_local_context".

Pack Expansion Support

The scope stack provides infrastructure for variadic template parameter pack expansion during instantiation.

Scope Stack Query Functions

Debug and Statistics

Scope Statistics Dump (sub_702DC0)

95 lines. Prints scope stack statistics when debug tracing is enabled. Output format:

Scope stack statistics
Stack entry size: 784
Max. stack depth: <N>

Followed by per-scope-kind counts using all scope kind display names.

Scope Entry Dump (sub_700260 / sub_7002D0)

• sub_700260 (17 lines): Prints " scope %d" with scope kind name via scope_kind_to_string. Detects bad depth with "***BAD SCOPE DEPTH***".
• sub_7002D0 (111 lines): Detailed dump using format "%s%3ld %3d " with associated type/symbol information.

End-of-Scope Processing

wrapup_scope (sub_706710)

381 lines. Major scope cleanup function called from pop_scope. Processes all symbols in the scope, emits needed definitions, runs end_of_scope_symbol_check. Debug traces: "wrapup_scope", "Wrapping up ", " scope".

end_of_scope_symbol_check (sub_705440)

781 lines, 6 assertion sites. The largest validation function. Checks:

• Symbol-to-IL-entry parent-class consistency ("end_of_scope_symbol_check: sym/il-entry parent-class mismatch")
• Parameter-to-routine association ("end_of_scope_symbol_check: parameter with no assoc routine")
• Hash table statistics ("hash_stats", "Hash statistics for: ")

set_needed_flags_at_end_of_file_scope (sub_707040)

188 lines. Determines which entities need to be emitted at the end of the translation unit. Validates scope kind ("set_needed_flags_at_end_of_file_scope: bad scope kind"). Debug brackets: "Start of set_needed_flags_at_end_of_file_scope\n" / "End of set_needed_flags_at_end_of_file_scope\n".

finish_function_body_processing (sub_6FE2A0)

142 lines. Post-processes function bodies after the scope closes. Determines whether the function needs to be emitted ("routine_needed_even_if_unreferenced", "Not calling mark_as_needed for", "storage class is %s\n").

Cross-References

• Entity Node Layout -- entity kind enum, execution space byte at +182
• IL Overview -- IL scope kinds (sck_*), IL entry kind 23 (scope, 288 bytes)
• IL Allocation -- alloc_scope allocator for 288-byte IL scope nodes
• Template Instance Record -- template instantiation data structures
• Translation Unit Descriptor -- TU pointer stored at scope entry +216