Format Specifiers

The cudafe++ diagnostic system uses a custom format specifier language -- not printf -- to expand parameterized error messages. The expansion engine is process_fill_in (sub_4EDCD0, 1,202 decompiled lines in error.c), called by write_message_to_buffer (sub_4EF620, 159 lines) during template string expansion. Each diagnostic record carries a linked list of typed fill-in entries that supply the actual values -- type nodes, entity pointers, strings, integers, source positions -- which the format engine renders into the final message text.

This page documents the specifier syntax, the fill-in kind system, entity-kind dispatch, suffix options, numeric indexing, and the labeled fill-in mechanism.

Specifier Syntax

When write_message_to_buffer walks an error template string (looked up from off_88FAA0[error_code]), it recognizes three format constructs:

Positional Specifier Parsing

The parser (sub_4EF620, error.c:4703) processes %XY...Zn specifiers as follows:

// After seeing '%', read next char as specifier letter
char spec_letter = template[pos + 1];      // 'T', 'd', 'n', 'p', 'r', 's', 't', 'u'
pos += 2;

// Collect option characters (a-z, A-Z) into buffer, max 29
int opt_count = 0;
char options[30];
while (true) {
    char c = template[pos];
    if (c >= '0' && c <= '9') {
        // Trailing digit = fill-in index (1-based)
        fill_in_index = c - '0';
        break;
    }
    if ((c & 0xDF) < 'A' || (c & 0xDF) > 'Y') {
        // Not a letter -- end of specifier, index defaults to 1
        fill_in_index = 1;
        break;
    }
    options[opt_count++] = c;
    if (opt_count > 29)
        assertion_handler("error.c", 4739,
            "write_message_to_buffer",
            "construct_text_message:",
            "too many option characters");
    pos++;
}
options[opt_count] = '\0';

process_fill_in(diagnostic_record, spec_letter, options, fill_in_index);

The maximum of 29 option characters is enforced by an assertion. In practice, specifiers use 0--3 option characters.

Fill-In Kinds

The specifier letter maps to a fill-in kind value through a switch on (letter - 84) in process_fill_in (sub_4EDCD0, error.c:4297):

Any other letter triggers the assertion: "process_fill_in: bad fill-in kind" (error.c:4297).

Usage Frequency Across 3,795 Templates

Measured across all error message templates in off_88FAA0:

Fill-In Entry Layout

Each fill-in entry is a 40-byte node allocated from a pool (qword_106B490) or heap by alloc_fill_in_entry (sub_4F2DE0):

Payload Layout by Kind

Kind 0 (decimal, %d) / Kind 1 (unsigned, %u) / Kind 3 (string, %s) / Kind 5 (type, %t) / Kind 6 (type, %T):

Kind 2 (position, %p):

Kind 4 (entity name, %n):

Kind 7 (%r):

Fill-In Linked List

Fill-in entries attach to the diagnostic record as a singly-linked list:

• Head pointer: diagnostic record offset 184 (fill_in_list_head)
• Tail pointer: diagnostic record offset 192 (fill_in_list_tail)

When process_fill_in searches for a matching entry, it walks the list from head, looking for the first entry where node->kind == requested_kind. If the specifier includes an index (e.g., %t2), it skips index - 1 matching entries before consuming the target:

const __m128i *node = *(diagnostic + 184);   // fill_in_list_head
if (!node)
    goto fill_in_not_found;

while (node->kind != requested_kind || --index > 0) {
    node = node->next;                        // offset 8
    if (!node)
        goto fill_in_not_found;
}

node->used_flag = 1;                          // mark consumed (offset 4)
// proceed with kind-specific rendering

If no matching entry is found, process_fill_in triggers an assertion with a diagnostic message identifying the missing fill-in: "specified fill-in (%X, N) not found for error string: \"...\"" (error.c:4317).

After all format specifiers have been expanded, construct_text_message (sub_4EF9D0) iterates the entire fill-in list and asserts that every entry has used_flag == 1. An unconsumed fill-in triggers: "construct_text_message: not all fill-ins used for error string: \"...\"" (error.c:4781).

Numeric Indexing

When a template string must reference multiple fill-ins of the same kind, a trailing digit selects which one:

The index is a single digit 0--9. Index 0 behaves identically to index 1 (the counter is pre-decremented before comparison). In practice, most templates use indices 1 and 2; a few use up to 3.

Real template example (CUDA cross-space call, error 3499):

calling a __device__ function(%sq1) from a __host__ function(%sq2) is not allowed

Here %sq1 and %sq2 are both kind 3 (string) with option q (quoted), selecting the first and second string fill-ins respectively. The caller attaches two string fill-ins -- the called function's name and the calling function's name.

Suffix Options

String Options (%s)

The %s specifier accepts only one option character: q for quoted output.

The q option wraps the string in double-quote characters (") and applies colorization if enabled (quote category, code 6 = bold). Any other option character on %s triggers: "process_fill_in: bad option" (error.c:4364).

Multiple q characters are permitted syntactically (the parser loops over all option chars validating each is q) but have no additional effect -- only one layer of quoting is applied.

Entity Name Options (%n)

The %n specifier accepts a rich set of option suffixes that control how an entity is rendered. Options are processed left-to-right, setting flags on the fill-in entry's flag bytes (offsets 28--34):

Options can be combined. Common combinations from the error template table:

No Options for Other Kinds

The %d, %u, %p, %t, %T, and %r specifiers reject all option characters:

if (*options != '\0')
    assertion_handler("error.c", 4372,
        "process_fill_in",
        "process_fill_in: bad option", NULL);

Kind-Specific Rendering

Kind 0 -- Signed Decimal (%d)

Renders the 64-bit signed integer payload using snprintf(buf, 20, "%lli", value), then writes the result to the output buffer. The 20-character buffer accommodates the full range of int64_t values including the sign.

Kind 1 -- Unsigned Decimal (%u)

Formats the payload through sub_4F63D0, which renders the unsigned 64-bit value into a dynamically-sized string buffer.

Kind 2 -- Source Position (%p)

Calls sub_4F6820 (form_source_position) with the position cookie from the fill-in payload. The rendering includes:

• File name (via sub_5B15D0 for display formatting)
• Line number
• Contextual text supplied by the caller through three string arguments (prefix, suffix, end-of-source fallback)

The caller passes context strings like " (declared ", ")", "(at end of source)" to frame the position reference. When the position resolves to line 0 or the file is "-" (stdin), alternate formats are used.

Kind 3 -- String (%s / %sq)

Without the q option, writes the string pointer payload directly to the output buffer via strlen + sub_6B9CD0 (buffer append).

With the q option, wraps the string in double quotes with colorization:

if (colorization_active)
    emit_escape(buffer, 6);       // quote color (bold)
write_char(buffer, '"');
write_string(buffer, payload);
if (colorization_active)
    emit_escape(buffer, 1);       // reset
write_char(buffer, '"');

Kind 5 -- Type, Lowercase (%t)

Renders the type node through the type formatting subsystem. The rendering pipeline:

1. Set byte_10678FA = 1 (name lookup kind = type display mode)
2. Write opening "
3. Call sub_600740 (format type for display) with the type node and the entity formatter callback table (qword_1067860)
4. Write closing "
5. Check via sub_7BE9C0 if the type has an "aka" (also-known-as) desugared form
6. If yes, append ' (aka "desugared_type")' -- comparing the rendered forms to avoid redundant output when they are identical

The aka check compares the rendered text of the original type against the desugared type. If they produce identical strings (same length, same content via strncmp), the aka suffix is suppressed by truncating the buffer back to the pre-aka position.

Kind 6 -- Type, Uppercase (%T)

Renders a type template argument list in angle brackets:

write_string(buffer, "\"<");
// Walk the template argument linked list
for (arg = payload; arg != NULL; arg = arg->next) {
    if (arg->kind != 3)   // skip pack expansion markers
        format_template_argument(arg, &entity_formatter);
    if (arg->next && arg->next->kind != 3)
        write_string(buffer, ", ");
}
write_string(buffer, ">\"");

Template argument entries with kind == 3 (at byte offset +8) are pack-expansion markers and are skipped during rendering.

Kind 7 -- Template Parameter Reference (%r)

Renders a template parameter by looking up the parameter entity through sub_5B9EE0 (entity lookup by scope + index). If found and non-null, renders via sub_4F3970 (unqualified entity name). Otherwise, falls back to sub_6011F0 (generic template parameter formatting).

Entity Kind Dispatch (%n)

When processing %n specifiers, process_fill_in reads the entity kind byte at offset 80 of the entity node and dispatches to kind-specific rendering logic. The function first resolves through projection indirection: if entity_kind == 16 (typedef), it follows the pointer at entity->info_ptr->pointed_to; if entity_kind == 24 (resolved namespace alias), it follows entity->info_ptr.

The dispatch handles 25 entity kind values (0--24, with gaps at 14/15/16/24 handled as special cases):

Any entity kind value outside 0--24 (excluding the gaps that trigger assertions) hits the default case: "form_symbol_summary: unsupported symbol kind" (error.c:2023).

Entity Rendering Pipeline

For entity kinds that produce a fully-formatted name (most non-trivial cases), the rendering proceeds through these stages:

1. Write entity kind label string (e.g., "function ")
   └── sub_6B9EA0(buffer, kind_label_string)
   └── sub_6B9CD0(buffer, " ", 1)

2. Open quote
   └── Optional colorization: sub_4ECDD0(buffer, 6)   // quote color
   └── sub_6B9CD0(buffer, "\"", 1)

3. Render type prefix (if has_type_info and full_qualification)
   └── sub_5FE8B0(type_node, 0, 1, 0, 0, &entity_formatter)

4. Render qualified or unqualified name
   ├── With template context:  sub_737A00(entity, &entity_formatter)
   └── Without template context: sub_4F3970(entity)

5. Render function parameters (if applicable)
   ├── Full parameter types: sub_5FB270(type, 0, 0, &entity_formatter)
   └── Simple type suffix:   sub_6016F0(type, &entity_formatter)

6. Close quote
   └── sub_6B9CD0(buffer, "\"", 1)
   └── Optional colorization: sub_4ECDD0(buffer, 1)   // reset

7. Append accessibility info (if 'a' option)
   └── " based on template argument(s) "
   └── sub_5FA660(template_arg_list, 0, &entity_formatter)

8. Append declaration location (if 'd' option)
   └── sub_4F6820(position, diag, " (declared ", ")", "(at end of source)")

9. Append translation unit info (if 'T' option)
   └── " (from translation unit <filename>)"

The original_name flag (o option) suppresses steps 1 and 3, rendering only the bare quoted name without a kind prefix or type qualification. The full_qualification flag (f option) enables step 3 and uses sub_737A00 for fully-qualified name rendering in step 4. The parameter_list flag (p option) forces step 5 to include full parameter-type rendering.

Template Context in Entity Names

When dword_126E274 (show template arguments) is non-zero and the entity has template context, the renderer can walk up the template scope chain:

1. Access the entity's routine info (for functions: offset 88 → offset 192 → offset 16)
2. Check for the instantiated-from entity (offset 104 of scope info, guarded by !(offset_176 & 1))
3. If found, use the instantiated-from entity as the display target
4. For class templates (entity_kind == 20): walk the template parameter chain, rendering <param1, param2, ...> with pack-expansion markers (...) for variadic parameters

CUDA-Specific Entity Rendering

Several entity kinds have CUDA-aware rendering paths:

• Class/struct (kinds 4/5): When dword_126EFB4 == 2 (CUDA C++ mode) and the entity has an anonymous flag (offset 161, bit 0x80), rendering jumps to the anonymous-class handler (kind 19) instead
• Field (kind 8): In CUDA C++ mode, the kind label is "member" (index 1480); in C mode, it is "field" (index 1481)
• Class/struct label selection: In CUDA C++ mode, the kind label index is always 1467; in non-CUDA mode, it depends on whether the entity is class vs struct

Labeled Fill-Ins (%[label])

The %[label] syntax references a named fill-in from the label table at off_D481E0. This mechanism allows error templates to include conditional text fragments that vary based on language mode or compilation context.

Label Table Structure

off_D481E0 is an array of 24-byte entries (3 pointers per entry):

Label Lookup Algorithm

// write_message_to_buffer, error.c:4714
char *label_start = template + pos + 2;      // skip "%["
char *label_end = strchr(template + pos + 1, ']');
if (!label_end)
    assertion_handler("error.c", 4714, "write_message_to_buffer", NULL, NULL);

size_t label_len = label_end - label_start;

// Walk off_D481E0 table
struct label_entry *entry = off_D481E0;
while (entry->name) {
    if (strncmp(entry->name, label_start, label_len) == 0) {
        // Found matching label
        int string_index;
        if (*entry->condition_ptr)
            string_index = entry->true_index;
        else
            string_index = entry->false_index;

        if (string_index > 3794)
            error_text_invalid_code();     // sub_4F2D30

        // Expand the referenced string directly into the buffer
        const char *text = off_88FAA0[string_index];
        write_to_buffer(buffer, text, strlen(text));
        pos = label_end + 1;
        break;
    }
    entry++;   // advance by 24 bytes
}

if (!entry->name) {
    // Label not found -- fatal
    fprintf(stderr, "missing fill-in label: %.*s\n", label_len, label_start);
    assertion_handler("error.c", 430,
        "get_label_fill_in_entry",
        "get_label_fill_in_entry: no label fill-in found", NULL);
}

The label table entries reference string indices in the same off_88FAA0 table used for error messages. This allows a single error template to produce different text depending on compilation mode -- for example, using "class" vs "struct" based on a language-mode flag, or "virtual" vs "" based on a feature flag.

The label text is written directly to the output buffer without further format specifier processing -- labels cannot contain nested % specifiers.

Output Buffer

All rendering targets the global message text buffer at qword_106B488:

• Initial allocation: 0x400 bytes (1 KB) via sub_6B98A0
• Dynamic growth: sub_6B9B20 doubles the buffer when capacity is exceeded
• String append: sub_6B9CD0(buffer, data, length) -- the workhorse write function
• String write: sub_6B9EA0(buffer, string) -- convenience wrapper (calls strlen + sub_6B9CD0)

The entity display callback infrastructure at qword_1067860 allows the type/name formatting subsystem to write to the same buffer through an indirect call:

Colorization Interaction

When dword_126ECA4 (colorization active) is non-zero, the format engine inserts ANSI escape sequences around quoted names and type references:

The escape sequences are emitted by sub_4ECDD0(buffer, color_code). The color codes correspond to the categories parsed from EDG_COLORS / GCC_COLORS environment variables during initialization.

Function Map

Cross-References

• Diagnostic Overview -- 7-stage pipeline, severity levels, diagnostic record layout
• CUDA Error Catalog -- all 338 CUDA-specific error templates with specifier usage
• SARIF & Pragma Control -- SARIF JSON output and #pragma nv_diagnostic system