Lambda Restrictions

Extended lambdas are CUDA's most constraint-heavy feature. Before a lambda can be wrapped in __nv_dl_wrapper_t or __nv_hdl_wrapper_t for device transfer, cudafe++ must verify that the closure type is serializable: no reference captures (device memory cannot hold host-side pointers), no function-local types in the public interface (device compiler has no access to them), no unnamed parent classes (the wrapper tag requires a mangleable name), and dozens of other structural invariants. The restriction checker runs as Phase 4 of scan_lambda (sub_447930, lines 626--866 of the 2113-line function) and continues through per-capture validation in make_field_for_lambda_capture (sub_42EE00) and recursive type walks in sub_41A3E0 / sub_41A1F0. Together, these functions enforce 39 distinct diagnostic tags covering 35+ error categories and approximately 45 unique error code call sites.

All restrictions apply only when dword_106BF38 (--extended-lambda / --expt-extended-lambda) is set and the lambda has an explicit __device__ or __host__ __device__ annotation. Standard C++ lambdas and lambdas defined inside __device__ / __global__ function bodies are exempt.

Key Facts

Restriction Categories

The tables below list every restriction enforced by the extended lambda validator, organized by the phase of validation in which each check occurs. The Error column gives the internal error index (displayed to users as 20000-series with the renumbering formula code + 16543). The Tag column gives the diagnostic tag name usable with --diag_suppress / #pragma nv_diag_suppress.

Category 1: Capture Restrictions

These checks run in two phases. The per-lambda checks (3593, 3595) occur in scan_lambda Phase 4 and in sub_4F9F20 (capture count finalization). The per-capture checks (3596--3599, 3616) run inside make_field_for_lambda_capture (sub_42EE00), which calls sub_41A1F0 for array dimension and constructibility analysis.

Category 2: Type Restrictions

Type restrictions enforce that every type visible in the lambda's public interface (captures, parameters, return type, and parent function template arguments) is accessible to the device compiler. Three contexts are checked, each with two sub-checks (function-local types and private/protected class member types). Additionally, the parent function's template arguments are checked for private/protected template members.

Type Walk Dispatch via dword_E7FE78

The callback sub_41B420 uses a global discriminator dword_E7FE78 to select between the three contexts. Each context is called with a different value:

The dispatch formula in sub_41B420 is 4 * (dword_E7FE78 != 1) + base_error. For local types, base is 3603; for private types, base is 3604. When dword_E7FE78 == 0, the multiplier is 41 = 4, yielding 3603+0 / 3604+0. When dword_E7FE78 == 1, the multiplier is 40 = 0, yielding 3603+3 = 3606 / 3604+3 = 3607. When dword_E7FE78 == 2 (and != 1), the multiplier is 4*1 = 4, yielding 3603+4 = (incorrect -- the actual formula uses a conditional). In practice the decompiled code shows:

// For function-local type check:
v2 = 3603;
if (dword_E7FE78)
    v2 = 4 * (unsigned int)(dword_E7FE78 != 1) + 3606;
// dword_E7FE78=0 -> 3603
// dword_E7FE78=1 -> 4*0 + 3606 = 3606
// dword_E7FE78=2 -> 4*1 + 3606 = 3610

// For private/protected type check:
v4 = 3604;
if (dword_E7FE78)
    v4 = 4 * (unsigned int)(dword_E7FE78 != 1) + 3607;
// dword_E7FE78=0 -> 3604
// dword_E7FE78=1 -> 4*0 + 3607 = 3607
// dword_E7FE78=2 -> 4*1 + 3607 = 3611

The tree walk itself is invoked via sub_7B0B60(type_node, sub_41B420, error_base). The error_base parameter (792 or 795) is stored in a global and used by the walker to control recursion behavior, not error selection.

Category 3: Enclosing Parent Function Restrictions

The parent function (the function in whose body the extended lambda is defined) must satisfy several naming and linkage constraints. These exist because the device compiler must be able to instantiate the wrapper template at a globally-unique mangled name derived from the parent function's signature.

Category 4: Template Parameter Restrictions

The parent function's template parameter list must satisfy naming and variadic constraints to ensure the wrapper tag type can be uniquely instantiated.

Category 5: Nesting and Context Restrictions

Category 6: Specifier and Annotation Restrictions

Category 7: Enclosing Scope Miscellaneous

Validation Architecture

Phase 4 of scan_lambda: Per-Lambda Validation

After parsing the capture list and annotations (Phases 1--3), scan_lambda enters the extended lambda validation block. This block is guarded by dword_106BF38 (extended lambda mode) and the annotation bits at lambda_info + 25. The validation proceeds as:

sub_447930 (scan_lambda), Phase 4 entry:
  |
  +-- Call sub_6BCDD0 (nv_find_parent_lambda_function)
  |     Returns: parent function node, sets is_device/is_template flags
  |
  +-- If parent == NULL:  emit error 3617
  +-- If ambiguous:       emit error 3618 or 3619
  |
  +-- Validate parent function properties:
  |     entity+81 bit 0  -> error 3592 (deprecated/deducible)
  |     entity+191 bit 0 -> error 3613 (noexcept spec)
  |     entity+166 == 1|2 -> error 3600 (deleted/defaulted)
  |     entity+81 bit 2  -> unnamed scope check -> error 3601
  |     entity+80 bits 0-1 -> address-taken / access check -> error 3608
  |
  +-- Walk parent scope chain for unnamed classes:
  |     entity+8 == NULL -> error 3609
  |     Non-trivial copy   -> error 3605
  |
  +-- Check capture-default conflicts:
  |     byte+24 bits 4+5 both set -> error 3593 (& and = conflict)
  |
  +-- OptiX gate: dword_106BDD8 -> error 3689
  |
  +-- Lambda numbering via red-black tree:
        Lookup failure -> error 3690

Per-Capture Validation: sub_42EE00

For each captured variable, make_field_for_lambda_capture runs targeted checks:

sub_42EE00 (make_field_for_lambda_capture):
  |
  +-- If byte+25 bit 3 set (device wrapper):
  |     |
  |     +-- Check init-capture for array type
  |     |     type_node+48 == 3 && sub_kind == 21 -> error 3596
  |     |
  |     +-- Call sub_41A1F0 (walk_type_for_hd_violations)
  |     |     Counts array dimensions, checks element type
  |     |     dim > 7 -> error 3597
  |     |     Not default-constructible -> error 3598
  |     |     Not assignable -> error 3599
  |     |
  |     +-- Check for pack expansion
  |           sub_7A8C00 returns true -> error 3616
  |
  +-- (Later) If byte+25 bit 4 set (HD wrapper):
        |
        +-- Call sub_7B0B60 with sub_41B420 callback
              Walks entire type tree, fires 3603/3604 for
              function-local and private/protected types

Type Hierarchy Walker: sub_41A3E0 / sub_41A1F0

sub_41A3E0 is the outer wrapper that validates the per-capture annotation context. sub_41A1F0 performs the recursive array dimension walk and element-type validation.

sub_41A3E0 (validate_type_hd_annotation):
  |
  +-- Determine context string: "__device__" or "__host__ __device__"
  |     Based on a2 parameter (0 = HD, nonzero = device-only)
  |
  +-- Check annotation byte (a1+32):
  |     bit 0 set && a2==0 -> error 3615
  |     bit 3 set -> check parent visibility:
  |       entity+163 < 0 (private) -> check bit pattern
  |       Both bits 3+4 set with private parent -> error 3635
  |       Otherwise -> error 3593
  |     bit 5 set -> error 3594 (private/protected access)
  |
  +-- Unwrap CV-qualifiers on element type (kind==12 loop)
  |
  +-- Call sub_41A1F0 (walk_type_for_hd_violations):
  |     Recursive array walker:
  |       v6 = dimension counter
  |       Loop: while sub_7A8370(type) returns true
  |         increment v6, follow sub_7A9310 to element type
  |       If v6 > 7: error 3597
  |       Unwrap CV (kind==12 loop)
  |       If not in dependent context (dword_126C5C4 == -1):
  |         Check scope flags (byte+6 bits 1-2)
  |         sub_550E50(30, type, 0) -> error 3598 (not default-constructible)
  |         sub_5BD540 + sub_510860(60, ...) -> error 3599 (not assignable)
  |       Call sub_7B0B60(type, sub_41B420, 792) for deep type walk
  |
  +-- If a3 (third parameter) set:
        Check bit 4 of annotation byte -> error 3691

sub_41B420: Type Walk Callback

This compact callback (33 lines decompiled) is invoked by sub_7B0B60 for every type node in the capture's type tree. It checks two properties:

1. Function-local type -- entity+81 bit 0 set: the type is defined inside a function body. Error selection uses dword_E7FE78 to pick between capture context (3603), operator() context (3606), and parent template-arg context (3610).

2. Private/protected member type -- entity+81 bit 2 set AND entity+80 bits 0-1 in range [1,2] (private or protected access specifier). Error selection parallels the local-type case: 3604, 3607, or 3611 depending on dword_E7FE78.

Special case: when entity+132 == 9 (template parameter dependent type) AND entity+152 points to a class with byte+86 bit 0 set AND entity+72 is non-null, the function-local check is suppressed. This handles template parameters that are not themselves local but instantiate with local types -- the error is deferred to instantiation time.

Diagnostic Tag Reference

Complete list of all 39 extended lambda diagnostic tags, sorted alphabetically. All tags can be used with --diag_suppress, --diag_warning, --diag_error on the command line, and with #pragma nv_diag_suppress, #pragma nv_diag_warning, #pragma nv_diag_error in source.

Bitmap Interaction

The capture count limit of 1023 derives from the bitmap architecture. Each wrapper type (device and host-device) uses a 128-byte bitmap (unk_1286980 / unk_1286900) storing 1024 bits. The bitmap setter sub_6BCBF0 performs:

result[capture_count >> 6] |= 1LL << capture_count;

Bit 0 is never emitted as a wrapper specialization (the zero-capture case uses the primary template). Bits 1--1023 map to generated partial specializations. The error check at capture count > 0x3FE (1022) fires before the bitmap set operation, so the effective maximum is 1023 captures. Attempting 1024 or more would overflow the 64-bit word boundary calculation, though in practice the error prevents this.

Operator() Annotation Derivation

Error 3621 enforces a fundamental design rule: the operator() function of an extended lambda must not carry explicit execution space annotations. Instead, the execution space is derived from the closure class. During scan_lambda Phase 5 (decl_call_operator_for_lambda), the code sets the call operator's execution space from lambda_info + 25:

// Propagate device/host from lambda_info to call operator
byte[operator+182] = (4 * byte[lambda+25]) & 0x10 | byte[operator+182] & 0xEF;
byte[operator+182] = (16 * byte[lambda+25]) & 0x20 | byte[operator+182] & 0xDF;

If the call operator already has execution space bits set (from explicit annotation by the user), error 3621 fires. The rationale is that the wrapper template's tag type already encodes the execution space; having the operator carry its own annotations would create an inconsistency that the device compiler cannot resolve.

Key Functions

Global State

Related Pages

• Extended Lambda Overview -- end-to-end pipeline, annotation bits, lambda_info layout
• Device Lambda Wrapper -- __nv_dl_wrapper_t template structure
• Host-Device Lambda Wrapper -- __nv_hdl_wrapper_t type-erased design
• Capture Handling -- __nv_lambda_field_type, __nv_lambda_array_wrapper
• Preamble Injection -- sub_6BCC20 emission pipeline
• CUDA Error Catalog -- complete error index with message templates
• Cross-Space Call Validation -- execution space checking infrastructure