Memory Spaces

Every CUDA variable that resides in GPU memory belongs to one of four memory spaces: __device__ (global memory), __shared__ (per-block scratchpad), __constant__ (read-only broadcast memory), or __managed__ (unified memory). cudafe++ encodes memory space as a two-byte bitfield at offsets +148 and +149 of the variable entity node. These two bytes are the variable-side analog of the execution space byte at +182 used for functions -- the two systems are complementary but independent.

The memory space bitfield passes through three processing stages. First, attribute handlers in attribute.c set the appropriate bits and enforce mutual exclusion constraints (no __shared__ + __constant__, no thread_local, no grid_constant conflict). Second, declaration processing in decls.c applies additional validation: VLA restrictions for __shared__, constexpr and external-linkage restrictions for __constant__/__device__, and structured binding constraints for all spaces. Third, symbol reference recording in symbol_ref.c checks whether host code illegally accesses device-side variables at reference time.

Memory spaces apply exclusively to variables (entity kind 7). __shared__ and __constant__ have no function-side meaning -- only __device__ (kind 'W', 87) doubles as a function execution space attribute.

Key Facts

The Memory Space Bitfield (Entity +148 / +149)

Byte +148: Primary Memory Space

Byte at entity+148:

  bit 0  (0x01)   __device__       Variable in device global memory
  bit 1  (0x02)   __shared__       Variable in per-block shared memory
  bit 2  (0x04)   __constant__     Variable in constant memory
  bit 3  (0x08)   type_member      Set when variable inherits space from type context
  bit 4  (0x10)   device_at_file   __device__ at file scope (no enclosing function)
  bit 7  (0x80)   weak_odr         Set by apply_nv_weak_odr_attr (sub_40AD80)

Bits 3, 4, and 7 are set by decl_variable (sub_4CA6C0) during declaration processing, not by the attribute handlers. Bit 3 is set via *(_BYTE *)(v33 + 148) |= 8u when the variable inherits its memory space from a type context (such as a static member of a class with a device annotation). Bit 4 is set via *(_BYTE *)(v43 + 148) = v73 | 0x10 when a __device__ variable is declared at file scope (dword_126C5D8 == -1, meaning no enclosing function).

Byte +149: Extended Memory Space

Byte at entity+149:

  bit 0  (0x01)   __managed__    Unified memory (host + device accessible)
  bits 1-7        (reserved)

Word-Level Access

Some validation code reads bytes +148 and +149 together as a 16-bit word. The __grid_constant__ conflict check in apply_nv_managed_attr tests:

// sub_40E0D0, line 26 (apply_nv_managed_attr)
if ( (a2[164] & 4) != 0 && (*((_WORD *)a2 + 74) & 0x102) != 0 )

Here (_WORD *)(a2 + 148) (offset 74 in 16-bit units) is tested against 0x0102. In little-endian layout, 0x0102 means byte +148 bit 1 (__shared__) OR byte +149 bit 0 (__managed__). This catches the case where a __grid_constant__ parameter also carries __shared__ or __managed__.

Mutual Exclusion

In valid CUDA programs, at most one of __device__, __shared__, and __constant__ should be set. However, __managed__ always implies __device__ -- the handler sets both +149 bit 0 and +148 bit 0. The validation logic permits __device__ + __managed__ but rejects combinations like __shared__ + __constant__.

The mutual exclusion check appears identically in both apply_nv_managed_attr and apply_nv_device_attr:

// From sub_40EB80 (apply_nv_device_attr), variable path:
v9 = *(_BYTE *)(a2 + 148) | 1;     // set __device__ bit
*(_BYTE *)(a2 + 148) = v9;
if ( ((v9 & 2) != 0) + ((v9 & 4) != 0) == 2 )
    sub_4F81B0(3481, a1 + 56);      // error: conflicting spaces

The expression ((v9 & 2) != 0) + ((v9 & 4) != 0) == 2 is true only when both __shared__ (bit 1) and __constant__ (bit 2) are set simultaneously. This means:

• __device__ + __shared__ is allowed (the bits coexist)
• __device__ + __constant__ is allowed
• __shared__ + __constant__ triggers error 3481

Attribute Handlers

apply_nv_managed_attr -- sub_40E0D0

The __managed__ handler is the simplest and most thoroughly documented. It demonstrates the full validation pattern that all memory space handlers share.

Entry point: Called from apply_one_attribute (sub_413240) when attribute kind is 'f' (102).

Decompiled logic (47 lines, attribute.c:10523):

// sub_40E0D0 -- apply_nv_managed_attr
// a1: attribute node, a2: entity node, a3: entity kind

// Gate: only applies to variables
if ( a3 != 7 )
    internal_error("attribute.c", 10523, "apply_nv_managed_attr");

// Step 1: Set managed flag AND device flag
v3 = a2[148];           // save old memory space byte
a2[149] |= 1;           // set __managed__ bit
a2[148] = v3 | 1;       // set __device__ bit (managed implies device)

// Step 2: Mutual exclusion check
if ( ((v3 & 2) != 0) + ((v3 & 4) != 0) == 2 )
    error(3481, ...);    // __shared__ + __constant__ conflict

// Step 3: Thread-local check
if ( (char)a2[161] < 0 )
    error(3482, ...);    // __managed__ on thread_local

// Step 4: Local variable check
if ( (a2[81] & 4) != 0 )
    error(3485, ...);    // __managed__ on local variable

// Step 5: __grid_constant__ conflict
if ( (a2[164] & 4) != 0 && (*(WORD*)(a2 + 148) & 0x102) != 0 )
{
    // Determine which space string to display
    v4 = a2[148];
    v5 = "__constant__";
    if ( (v4 & 4) == 0 ) {
        v5 = "__managed__";
        if ( (a2[149] & 1) == 0 ) {
            v5 = "__shared__";
            if ( (v4 & 2) == 0 ) {
                v5 = "__device__";
                if ( (v4 & 1) == 0 )
                    v5 = "";
            }
        }
    }
    error(3577, ..., v5);   // incompatible with __grid_constant__
}

The space-name selection cascade (__constant__ > __managed__ > __shared__ > __device__ > empty) is used in error messages to show which memory space conflicts with __grid_constant__. The cascade tests bits in priority order, matching the most "restrictive" space first.

apply_nv_device_attr -- sub_40EB80

The __device__ handler is dual-purpose: it handles both variables (a3 == 7) and functions (a3 == 11).

Entry point: Called from apply_one_attribute when attribute kind is 'W' (87).

Variable path (entity kind 7):

// sub_40EB80, variable branch
*(_BYTE *)(a2 + 148) |= 1;          // set __device__ bit

// Validation (identical to __managed__):
// 1. Error 3481 if __shared__ + __constant__ both set
// 2. Error 3482 if thread_local (byte +161 bit 7)
// 3. Error 3485 if local variable (byte +81 bit 2)
// 4. Error 3577 if __grid_constant__ conflict

Function path (entity kind 11):

// sub_40EB80, function branch
// Check: not an implicitly-deleted function
if ( (*(_QWORD *)(a2 + 184) & 0x800001000000LL) != 0x800000000000LL
     || (*(_BYTE *)(a2 + 176) & 2) != 0 )
{
    // Conflict with __global__
    if ( !dword_106BFF0 && (*(_BYTE *)(a2 + 182) & 0x40) != 0 )
        error(3481, ...);

    *(_BYTE *)(a2 + 182) |= 0x23;    // set device execution space

    // Local function with __global__ conflict
    if ( (*(_BYTE *)(a2 + 81) & 4) != 0 && (*(_BYTE *)(a2 + 182) & 0x40) != 0 )
        error(3688, ...);

    // __device__ on main()
    if ( a2 == qword_126EB70 && (*(_BYTE *)(a2 + 182) & 0x20) != 0 )
        warning(3538, ...);
}
else
{
    // Implicitly-deleted function: just warn
    v14 = get_entity_display_name(a2);
    error(3469, ..., "__device__", v14);
}

// Check function parameters for missing default initializers
// (error 3669 for parameters without defaults in device context)

The function path is documented in Execution Spaces -- here we focus on the variable path.

shared and constant Handlers

The __shared__ and __constant__ attribute handlers are dispatched through apply_one_attribute (sub_413240) when attribute kind codes 'Z' (90) and '[' (91) are encountered. Their variable-path logic mirrors __device__ and __managed__:

The __shared__ and __constant__ keywords apply only to variables (kind 7). Unlike __device__, they do not have a function-path branch -- there is no __shared__ or __constant__ function execution space.

Variable Declaration Processing

sub_4DEC90 -- variable_declaration

The top-level declaration processor (decls.c) performs additional CUDA-specific validation after attribute handlers have set the memory space bits. This function is 1098 lines and handles both normal variable declarations and static data member definitions.

CUDA-specific checks in variable_declaration:

Memory space string selection (used in error messages):

// sub_4DEC90, line ~357: selecting display name for the memory space
v50 = "__constant__";
if ( (v49 & 4) == 0 ) {
    v50 = "__managed__";
    if ( (*(_BYTE *)(v15 + 149) & 1) == 0 ) {
        v50 = "__host__ __device__" + 9;   // pointer arithmetic: = "__device__"
        if ( (v49 & 2) != 0 )
            v50 = "__shared__";
    }
}

The string "__device__" is produced by taking the string "__host__ __device__" and advancing by 9 bytes, skipping past "__host__ ". This is a binary-level optimization -- the compiler shares string storage between the combined "__host__ __device__" literal and the standalone "__device__" reference.

sub_4CA6C0 -- decl_variable

The core variable declaration function (1090 lines, decls.c:7730) handles CUDA memory space propagation during symbol table entry creation. Key behaviors:

Storage class mapping: When declaration state byte at offset +269 equals 5, it indicates a CUDA memory space storage class. The function performs a scope walk to determine the correct namespace scope for the variable. If a prior declaration exists at the same scope (dword_126C5DC == dword_126C5B4), the CUDA storage class is reset to allow redeclaration.

Scope walk: Traverses the scope chain (784-byte scope entries at qword_126C5E8, indexed by dword_126C5E4) upward through class scopes (scope_kind 4) and template scopes (bit 0x20 at scope entry +9), until reaching a non-class, non-template scope. This determines whether the variable is at namespace scope, class scope, or block scope.

Error 3483 -- memory space in non-device function: When a variable with a device memory space bit (+148 bit 0 set) is declared inside a function body, and the enclosing routine is NOT device-only (+182 & 0x30 != 0x20), the function emits error 3483 with the storage kind and space name:

// From sub_4CA6C0, ~line 886-910
if (!at_namespace_scope) {
    char space = entity->byte_148;
    if (storage_class != 1 && (space & 0x01)) {
        routine_descriptor = qword_126C5D0;
        if (routine_descriptor) {
            entity_ptr = *(routine_descriptor + 32);
            if (entity_ptr && (entity_ptr[182] & 0x30) != 0x20) {
                const char *name = get_space_name(entity);  // priority cascade
                const char *kind = (storage_class == 2) ? "a static" : "an automatic";
                error(3483, source_loc, kind, name);
            }
        }
    }
}

File-scope device flag: When a __device__ variable is at file scope (dword_126C5D8 == -1), the function sets bit 4 of +148:

if ((entity->byte_148 & 0x01) && dword_126C5D8 == -1)
    entity->byte_148 |= 0x10;   // bit 4: device_at_file_scope

Redeclaration checking: When a variable is redeclared, the function compares memory space encoding at offset +136 (the attribute byte) between the existing and new entity. Error 1306 is emitted for mismatched CUDA memory spaces.

Memory space propagation: Calls sub_4C4750 (set_variable_attributes) for final attribute propagation, and sub_4CA480 (check_variable_redeclaration) for prior-declaration compatibility.

sub_4DC200 -- mark_defined_variable

Post-declaration validation for device-memory variables with external linkage (26 lines):

// sub_4DC200 -- mark_defined_variable (decompiled)
void mark_defined_variable(entity_t *a1, int a2) {
    if (a1[164] & 0x10) {   // already marked as defined
        if (!dword_106BFD0                    // cross-space checking not overridden
            && (a1[148] & 3) == 1             // __device__ set, __shared__ NOT set
            && !is_compiler_generated(a1)     // not compiler-generated
            && (a1[80] & 0x70) != 0x10)       // not anonymous
        {
            warning(3648, a1 + 64);           // external linkage warning
        }
    } else if (!a2 && (*(byte*)(*(qword*)a1 + 81) & 2)) {
        error(1655, ...);   // tentative definition of constexpr
    } else {
        // Same 3648 check on first definition
        if (!dword_106BFD0 && (a1[148] & 3) == 1 && ...)
            warning(3648, a1 + 64);
        a1[164] |= 0x10;   // mark as defined
    }
}

The condition (a1[148] & 3) == 1 tests that bit 0 (__device__) is set AND bit 1 (__shared__) is NOT set. This catches __device__ variables (including __device__ __constant__ and __device__ __managed__, since those have bit 0 set) but excludes __shared__ variables (which have bit 1 set). The check is NOT about __constant__ alone -- a pure __constant__ variable (only bit 2 set, value 0x04) would yield (0x04 & 3) == 0, failing the test. The p1.06 report's characterization of error 3648 as "constant with external linkage" is misleading; the actual condition is "device-accessible (non-shared) variable with external linkage."

sub_4CC150 -- cuda_variable_fixup

Called from variable_declaration after CUDA constexpr-if detection. This function:

• Manipulates variable entity fields at offset +148 (memory space) and +162 (visibility flags)
• Adjusts scope chains using the 784-byte scope entry array
• Creates new type entries for CUDA-specific variable rewriting

Bit Assignment Resolution

Two sweep reports provided conflicting bit assignments for byte +148:

The decompiled code resolves this definitively in favor of the p1.01 assignment. Two independent functions confirm it:

1. sub_40E0D0 (apply_nv_managed_attr) sets a2[149] |= 1 (managed at +149) and a2[148] = v3 | 1 (device at +148 bit 0). The subsequent conflict check tests (v3 & 2) for __shared__ and (v3 & 4) for __constant__.

2. sub_40EB80 (apply_nv_device_attr) sets *(_BYTE *)(a2 + 148) | 1 (device at +148 bit 0), then uses the identical conflict test ((v9 & 2) != 0) + ((v9 & 4) != 0) == 2.

The canonical encoding is:

Byte +148:  bit 0 = __device__,  bit 1 = __shared__,  bit 2 = __constant__
Byte +149:  bit 0 = __managed__

The p1.06 report's alternative encoding is an analysis error, caused by mark_defined_variable (sub_4DC200) testing +148 & 3 == 1 in the context of error 3648. That test checks for __device__ set (bit 0) without __shared__ (bit 1) -- not for __constant__ at bit 0. The error was then characterized as "constant with external linkage" based on the error message text rather than the actual bit test.

Validation Constraints

managed Constraints

__managed__ has the strictest requirements among memory space annotations. All five checks occur in apply_nv_managed_attr (sub_40E0D0):

The __managed__ keyword requires compute capability >= 3.0. This is verified at compilation time via version threshold comparisons (qword_126EF90 > 0x78B3, where 0x78B3 = 30899 in the CUDA version encoding scheme). The specific error code for architecture-too-low is not captured in the decompiled attribute handler.

shared Constraints

__shared__ variables have restrictions enforced across multiple functions:

constant Constraints

__constant__ carries additional restrictions related to constexpr and type:

Note: Error 3648 (external linkage warning) is emitted by sub_4DC200 but the condition tests (byte+148 & 3) == 1, which checks for __device__ set without __shared__ -- not specifically __constant__. The check applies to any device-accessible non-shared variable, including __device__, __device__ __constant__, and __device__ __managed__.

Cross-Space Variable Access Checking

When host code references a device-side variable, the symbol reference recorder emits diagnostics. This checking occurs in record_symbol_reference_full (sub_72A650 / sub_72B510, symbol_ref.c) and is gated by global flags dword_106BFD0 and dword_106BFCC.

Gate Logic

1. Is cross-space checking enabled?
   → dword_106BFD0 != 0 OR dword_106BFCC != 0

2. Is the referenced entity a variable (kind == 7)?
   → Yes: proceed to nv_check_device_var_ref_in_host
   → No (kind 10/11/20 -- function): check nv_check_host_var_ref_in_device

3. Get current routine from scope stack (dword_126C5D8)
4. Check routine execution space at +182 (0x30 mask):
   → 0x00 or 0x10 (host): emit device-var-in-host errors
   → 0x20 (device): emit host-var-in-device errors

Device Variable Referenced from Host Code

The nv_check_device_var_ref_in_host path (assert string at symbol_ref.c:2347) checks memory space bits and produces specific errors based on which space the variable occupies:

Host Variable Referenced from Device Code

The nv_check_host_var_ref_in_device path (assert string at symbol_ref.c:2390) handles the reverse direction:

The error 3623 has two context strings:

• "outside the bodies of device functions" -- general case
• "from a constexpr or consteval __device__ function" -- constexpr context

Relaxation: dword_106BF40

When dword_106BF40 is set (corresponding to --expt-relaxed-constexpr), and the current routine at +182 has the device annotation pattern (& 0x30 == 0x20) with +177 bit 1 set (explicit __device__), cross-space variable access checks are suppressed. This allows constexpr device functions to reference host variables during constant evaluation.

Host Reference Arrays

When the backend emits host-side code, variables marked with __device__, __shared__, or __constant__ are registered in ELF section arrays so the CUDA runtime can discover them at load time. The emission function sub_6BCF80 (nv_emit_host_reference_array) writes entries into six separate sections:

Each array entry contains the mangled name of the device symbol as a byte array:

extern "C" {
    extern __attribute__((section(".nvHRDE")))
    __attribute__((weak))
    const unsigned char hostRefDeviceArrayExternalLinkage[] = {
        /* mangled name bytes */ 0x0
    };
}

Six global lists (at addresses unk_1286780 through unk_12868C0) accumulate symbols during compilation, one per section type. Note that __shared__ variables do NOT get host reference arrays -- they have no host-visible address.

Redeclaration Compatibility

When a variable is redeclared, decl_variable (sub_4CA6C0) compares the memory space bits between the prior declaration and the new one. Error 1306 is emitted for mismatched CUDA memory spaces:

Error 1306: CUDA memory space mismatch on redeclaration

The comparison tests byte +148 of both the existing entity and the new declaration's computed attributes. The CUDA memory space acts as an implicit storage class -- storage class value 5 in the declaration state (offset 269) indicates a CUDA-specific storage class that requires special scope-walking behavior.

String Table Usage

The memory space keywords appear in the binary's string table and are referenced by error message formatting code:

The pointer-arithmetic trick for "__device__" appears in both sub_4DEC90 (variable_declaration) and error message formatting throughout the attribute handlers. It saves binary space by reusing the combined "__host__ __device__" string constant.

Error Code Summary

Attribute Application Errors

Declaration Processing Errors

Cross-Space Reference Errors

Global State Variables

Function Map

See Also

• Execution Spaces -- function-level __host__ / __device__ / __global__ encoding at entity +182
• Cross-Space Call Validation -- full cross-space call checking algorithm
• Entity Node Layout -- complete entity node offset map
• __managed__ Variables -- __managed__ attribute system details
• __grid_constant__ -- __grid_constant__ parameter attribute
• Host Reference Arrays -- runtime device symbol discovery