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

CLI Flag Inventory

Quick Reference: 20 Most Important CUDA-Specific Flags

Flag (via -Xcudafe)nvcc EquivalentIDEffect
--diag_suppress=N--diag-suppress=N39Suppress diagnostic number N (comma-separated)
--diag_error=N--diag-error=N42Promote diagnostic N to error
--diag_warning=N--diag-warning=N41Demote diagnostic N to warning
--display_error_number--44Show #NNNNN-D error codes in output
--target=smXX--gpu-architecture=smXX245Set SM architecture target (parsed via sub_7525E0)
--relaxed_constexpr--expt-relaxed-constexpr104Allow constexpr cross-space calls
--extended-lambda--expt-extended-lambda106Enable __device__/__host__ __device__ lambdas in host code (dword_106BF38)
--device-c-rdc=true77Relocatable device code (separate compilation)
--keep-device-functions--keep-device-functions71Do not strip unused device functions
--no_warnings-w22Suppress all warnings
--promote_warnings-W23Promote all warnings to errors
--error_limit=N--32Maximum errors before abort (default: unbounded)
--force-lp64-m6465LP64 data model (pointer=8, long=8)
--output_mode=sarif--274SARIF JSON diagnostic output
--debug_mode-G82Full debug mode (sets 3 debug globals)
--device-syntax-only--72Device-side syntax check without codegen
--no-device-int128--52Disable __int128 on device
--zero_init_auto_vars--81Zero-initialize automatic variables
--fe-inlining--54Enable frontend inlining
--gen_c_file_name=path--45Set output .int.c file path

These are the flags most commonly passed through -Xcudafe for CUDA development. The full inventory of 276 flags follows below.


cudafe++ accepts 276 command-line flags registered in a flat table at dword_E80060. The flags are not parsed directly from the binary's argv -- NVIDIA's driver compiler nvcc decomposes its own options and invokes cudafe++ with the appropriate low-level flags. Users never run cudafe++ directly; instead, they pass options through nvcc -Xcudafe <flag>, which strips the -Xcudafe prefix and forwards the remainder as a bare argument to the cudafe++ process.

The flag system is implemented in three functions within cmd_line.c:

FunctionAddressLinesRole
register_command_flagsub_451F8025Insert one entry into the flag table
init_command_line_flagssub_4520103,849Register all 276 flags (called once)
proc_command_linesub_4596304,105Main parser: match argv against table, dispatch to 276-case switch (IDs 0..275, case 0 = default sink for unknown flags)
default_initsub_45EB40470Zero 350 global config variables + flag-was-set bitmap

Flag Table Structure

Each flag occupies a 40-byte entry in a contiguous array beginning at dword_E80060, with a maximum capacity of 552 entries (overflow triggers a panic via sub_40351D). The current count is tracked in dword_E80058.

struct flag_entry {                    // 40 bytes per entry
    int32_t   case_id;                 // dword_E80060[idx*10]    -- switch dispatch ID
    char*     name;                    // qword_E80068[idx*5]     -- long flag name string
    int16_t   short_char;              // word_E80070[idx*20]     -- single-char alias (0 if none)
    int8_t    is_valid;                // word_E80070[idx*20]+1   -- always 1
    int8_t    takes_value;             // byte_E80072[idx*40]     -- flag requires =<value> argument
    int32_t   visible;                 // dword_E80080[idx*10]    -- mode/action classification
    int8_t    is_boolean;              // byte_E80073[idx*40]     -- flag is on/off toggle
    int64_t   name_length;             // qword_E80078[idx*5]     -- strlen(name), precomputed
};

The flag-was-set bitmap at byte_E7FF40 spans 0x110 bytes (272 flag slots). When a flag is matched during parsing, the corresponding bit is set to record that the user explicitly provided it. This bitmap is zeroed by default_init before every compilation.

Registration Protocol

register_command_flag (sub_451F80) is called approximately 275 times from init_command_line_flags. Its prototype:

void register_command_flag(
    int    case_id,        // dispatch ID for the switch statement
    char*  name,           // "--name" (without the dashes)
    char   short_opt,      // single-letter alias, 0 for none
    char   takes_value,    // 1 if the flag requires =<value>
    int    mode_flag,      // visibility / classification
    char   enabled         // whether the flag is active
);

Some flags are registered as paired toggles -- --flag and --no_flag share the same case_id but set the target global to 1 or 0 respectively. These pairs are registered either by two calls to register_command_flag or by inline table population within init_command_line_flags.

Parsing Flow

proc_command_line (sub_459630) is the master CLI parser. It:

  1. Calls init_command_line_flags to populate the flag table (once)
  2. Allocates four hash tables for accumulating -D, -I, system include, and macro alias arguments
  3. Adjusts nine diagnostic severities by default via sub_4ED400: four are suppressed (severity 3: errors 1373, 1374, 1375, 2330) and five are demoted to remark (severity 4: errors 1257, 1633, 111, 185, 175)
  4. Enters the main loop over argv:
    • Scans for - prefix to identify flags
    • Handles -X short flags and --flag-name long flags
    • Handles --flag=value syntax via parse_flag_name_value (sub_451EC0)
    • Matches flag names against the registered table using strncmp against each entry's precomputed name_length
    • Dispatches to a giant switch(case_id) with 276 cases (IDs 0..275; case 0 is the default sink for unrecognized flag names)
  5. Executes post-parsing dialect resolution (described below)
  6. Opens output, error, and list files
  7. Treats the remaining non-flag argv entry as the input filename

The -Xcudafe Pass-Through

Users never invoke cudafe++ directly. The intended usage path is:

nvcc --some-option -Xcudafe --diag_suppress=1234 source.cu

nvcc strips -Xcudafe and passes --diag_suppress=1234 directly to the cudafe++ process as an argv element. Multiple -Xcudafe arguments accumulate. Because cudafe++ flags use -- long-form prefixes, there is no ambiguity with nvcc's own flag namespace.

Certain nvcc flags like --expt-extended-lambda and --expt-relaxed-constexpr are translated by nvcc into the corresponding cudafe++ internal flags (--extended-lambda, --relaxed_constexpr) before invocation. Users do not need to know the internal names.

Flag Catalog by Category

The 276 flags are grouped below by functional category. Each table lists:

  • ID -- the case_id used in the dispatch switch
  • Flag -- the --name as registered (paired flags shown as name / no_name)
  • Short -- single-character alias (dash required: -E, -C, etc.)
  • Arg -- whether the flag takes a =<value> argument
  • Effect -- what the flag does internally

Core EDG Flags (1--44)

These are standard Edison Design Group frontend options that predate NVIDIA's CUDA modifications.

IDFlagShortArgEffect
1strict-AnoEnable strict standards conformance mode
2strict_warnings-anoStrict mode with extra warnings
3no_line_commands-PnoSuppress #line directives in preprocessor output
4preprocess-EnoPreprocessor-only mode (output to stdout)
5comments-CnoPreserve comments in preprocessor output
6old_line_commands--noUse old-style # N "file" line directives
7old_c-KnoK&R C mode (calls set_c_mode(1))
8dependencies-MnoOutput #include dependency list (preprocessor-only)
9trace_includes-HnoPrint each #include file as it is opened
10il_display--noDump intermediate language after parsing
11anachronisms / no_anachronisms--noAllow/disallow anachronistic C++ constructs
12cfront_2.1-bnoCfront 2.1 compatibility mode
13cfront_3.0--noCfront 3.0 compatibility mode
14no_code_gen-nnoParse only, skip code generation
15signed_chars / unsigned_chars-snoDefault char signedness
16instantiate-tyesTemplate instantiation mode: none, all, used, local
17implicit_include / no_implicit_include-BnoEnable/disable implicit inclusion of template definitions
18suppress_vtbl / force_vtbl--noControl virtual table emission
19dollar-$noAllow $ in identifiers
20timing-#noPrint compilation phase timing
21version-vnoPrint version banner and continue
22no_warnings-wnoSuppress all warnings (sets severity threshold to error-only)
23promote_warnings-WnoPromote warnings to errors
24remarks-rnoEnable remark-level diagnostics
25c-mnoForce C language mode
26c++-pnoForce C++ language mode
27exceptions / no_exceptions-xnoEnable/disable C++ exception handling
28no_use_before_set_warnings-jnoSuppress "used before set" variable warnings
29include_directory-IyesAdd include search path (handles - for stdin)
30define_macro-DyesDefine preprocessor macro (builds linked list)
31undefine_macro-UyesUndefine preprocessor macro
32error_limit-eyesMaximum number of errors before abort
33list-LyesGenerate listing file
34xref-XyesGenerate cross-reference file
35error_output--yesRedirect error output to file
36output-oyesSet output file path
37db-dyesLoad debug database
38time_limit--yesSet compilation time limit
39diag_suppress--yesSuppress diagnostic numbers (comma-separated list)
40diag_remark--yesDemote diagnostics to remark severity
41diag_warning--yesSet diagnostics to warning severity
42diag_error--yesPromote diagnostics to error severity
43diag_once--yesEmit diagnostic only on first occurrence
44display_error_number / no_display_error_number--noShow/hide error code numbers in output

NVIDIA CUDA-Specific Flags (45--89)

These flags are NVIDIA additions absent from stock EDG. They control CUDA compilation modes, device code generation, and host/device interaction.

IDFlagArgEffect
45gen_c_file_nameyesSet output .int.c file path (qword_106BF20)
46msvc_target_versionyesMSVC version for compatibility (dword_126E1D4)
47host-stub-linkage-explicitnoUse explicit linkage on host stubs
48static-host-stubnoGenerate static host stubs
49device-hidden-visibilitynoApply hidden visibility to device symbols
50no-hidden-visibility-on-unnamed-nsnoExempt unnamed namespaces from hidden visibility
51no-multiline-debugnoDisable multiline debug info
52no-device-int128noDisable __int128 on device
53no-device-float128noDisable __float128 on device
54fe-inliningnoEnable frontend inlining (dword_106C068 = 1)
55modify-stack-limityesControl stack limit modification (dword_106C064)
56fassociative-mathnoEnable associative floating-point math
57orig_src_file_nameyesOriginal source file name (before preprocessing)
58orig_src_path_nameyesOriginal source path name (full path)
59frandom-seedyesRandom seed for reproducible output
60check-template-param-qualnoCheck template parameter qualifications
61check-clock-callnoValidate clock() calls in device code
62check-ffs-callnoValidate ffs() calls in device code
63check-routine-address-takennoCheck when device routine address is taken
64check-memory-clobbernoValidate memory clobber in inline asm
65force-lp64noLP64 data model: pointer=8, long=8
66force-llp64noLLP64 data model: pointer=4, long=4
67pgi_llvmnoPGI/LLVM backend mode
68pgi_arch_ppcnoPGI PowerPC architecture
69pgi_arch_aarch64noPGI AArch64 architecture
70pgi_versionyesPGI compiler version number
71keep-device-functionsnoDo not strip unused device functions
72device-syntax-onlynoDevice-side syntax check without codegen
73device-time-tracenoEnable device compilation time tracing
74force_linkonce_to_weaknoConvert linkonce to weak linkage
75disable_host_implicit_call_checknoSkip implicit call validation on host
76no_strict_cuda_errornoRelax strict CUDA error checking
77device-cnoRelocatable device code (RDC) mode
78no-shadow-functionsnoDisable function shadowing in device code
79disable_ext_lambda_cachenoDisable extended lambda capture cache
80no-constant-variable-inferencingnoDisable constexpr variable inference on device
81zero_init_auto_varsnoZero-initialize automatic variables
82debug_modenoFull debug mode (sets 3 debug globals to 1)
83gen_module_id_filenoGenerate module ID file
84include_file_nameyesForced include file name
85gen_device_file_nameyesDevice-side output file name
86stub_file_nameyesStub file output path
87module_id_file_nameyesModule ID file path
88tile_bc_file_nameyesTile bitcode file path
89tile-onlynoTile-only compilation mode

Architecture and Host Compiler Flags (90--114)

These flags identify the target architecture and host compiler for compatibility emulation.

IDFlagShortArgEffect
90m32--no32-bit mode: pointer=4, long=4, all types sized for ILP32
91m64--no64-bit mode (default on Linux x86-64)
92Version-VnoPrint version with different copyright format, then exit(1)
93compiler_bindir--yesHost compiler binary directory
94sdk_dir--yesSDK directory path
95pgc++--noPGI C++ compiler mode
96icc--noIntel ICC compiler mode
97icc_version--yesIntel ICC version number
98icx--noIntel ICX (oneAPI) compiler mode
99grco--noGRCO compiler mode
100allow_managed--noAllow __managed__ variable declarations
101gen_system_templates_from_text--noGenerate system templates from text
102no_host_device_initializer_list--noDisable HD initializer_list support
103no_host_device_move_forward--noDisable HD std::move/std::forward
104relaxed_constexpr--noRelaxed constexpr rules for device code (--expt-relaxed-constexpr)
105dont_suppress_host_wrappers--noEmit host wrapper functions unconditionally
106arm_cross_compiler--noARM cross-compilation mode
107target_woa--noWindows on ARM target
108gen_div_approx_no_ftz--noGenerate approximate division without flush-to-zero
109gen_div_approx_ftz--noGenerate approximate division with flush-to-zero
110shared_address_immutable--noShared memory addresses are immutable
111uumn--noUnnamed union member naming

C++ Language Feature Toggle Flags (115--275)

The largest group -- approximately 120 paired boolean toggles that control individual C++ language features. Most are inherited from EDG's configuration surface. Each pair shares a case_id and sets a global variable to 1 (--flag) or 0 (--no_flag).

Precompiled Headers (115--121)

IDFlagArgEffect
115unsigned_wchar_tnowchar_t is unsigned
116create_pchyesCreate precompiled header file
117use_pchyesUse existing precompiled header
118pchnoEnable PCH mode
119pch_messages / no_pch_messagesnoShow/hide PCH status messages
120pch_verbose / no_pch_verbosenoVerbose PCH output
121pch_diryesPCH file directory

Core C++ Feature Toggles (122--170)

IDFlagArgDefault
122restrict / no_restrictnoon
123long_lifetime_temps / short_lifetime_tempsno--
124wchar_t_keyword / no_wchar_t_keywordnoon
125pack_alignmentyes--
126alternative_tokens / no_alternative_tokensnoon
127svr4 / no_svr4no--
128brief_diagnostics / no_brief_diagnosticsno--
129nonconst_ref_anachronism / no_nonconst_ref_anachronismno--
130no_preproc_onlyno--
131rtti / no_rttinoon
132building_runtimeno--
133bool / no_boolnoon
134array_new_and_delete / no_array_new_and_deleteno--
135explicit / no_explicitno--
136namespaces / no_namespacesnoon
137using_std / no_using_stdno--
138remove_unneeded_entities / no_remove_unneeded_entitiesnoon
139typename / no_typenameno--
140implicit_typename / no_implicit_typenamenoon
141special_subscript_cost / no_special_subscript_costno--
143old_style_preprocessingno--
144old_for_init / new_for_initno--
145for_init_diff_warning / no_for_init_diff_warningno--
146distinct_template_signatures / no_distinct_template_signaturesno--
147guiding_decls / no_guiding_declsnoon
148old_specializations / no_old_specializationsnoon
149wrap_diagnostics / no_wrap_diagnosticsno--
150implicit_extern_c_type_conversion / no_implicit_extern_c_type_conversionno--
151long_preserving_rules / no_long_preserving_rulesno--
152extern_inline / no_extern_inlineno--
153multibyte_chars / no_multibyte_charsno--
154embedded_c++noEmbedded C++ mode
155vla / no_vlano--
156enum_overloading / no_enum_overloadingno--
157nonstd_qualifier_deduction / no_nonstd_qualifier_deductionno--
158late_tiebreaker / early_tiebreakerno--
159preincludeyes--
160preinclude_macrosyes--
161pending_instantiationsyes--
162const_string_literals / no_const_string_literalsnoon
163class_name_injection / no_class_name_injectionnoon
164arg_dep_lookup / no_arg_dep_lookupnoon
165friend_injection / no_friend_injectionnoon
166nonstd_using_decl / no_nonstd_using_declno--
168designators / no_designatorsno--
169extended_designators / no_extended_designatorsno--
170variadic_macros / no_variadic_macrosno--
171extended_variadic_macros / no_extended_variadic_macrosno--

Include Paths and Module Support (167, 172, 256--265)

Note: These flags use non-contiguous IDs because sys_include and incl_suffixes are registered early, while the C++20 module flags use a separate ID range (256+).

IDFlagArgEffect
167sys_includeyesSystem include directory
172incl_suffixesyesInclude file suffix list (default "::stdh:")
256modules_directoryyesC++20 modules directory
257ms_mod_file_mapyesMSVC module file mapping
258ms_header_unityesMSVC header unit
259ms_header_unit_quoteyesMSVC quoted header unit
260ms_header_unit_angleyesMSVC angle-bracket header unit
261ms_mod_interface / no_ms_mod_interfacenoMSVC module interface mode
262ms_internal_partition / no_ms_internal_partitionnoMSVC internal partition mode
263ms_translate_include / no_ms_translate_includenoMSVC translate #include to import
264modules / no_modulesnoEnable/disable C++20 modules
265module_import_diagnostics / no_module_import_diagnosticsnoModule import diagnostic messages

Host Compiler and Language Feature Toggles (182--239)

Note: All IDs below are verified against the decompiled init_command_line_flags (sub_452010). Flags are registered by sub_451F80 (explicit call) or by inline array population. IDs are not sequential -- gaps exist where flags were removed or repurposed.

IDFlagArgDefault
182gcc / no_gccnoGCC compatibility mode
183g++ / no_g++noG++ mode (alias for GCC C++ mode)
184gnu_versionyesGCC version number (default 80100 = GCC 8.1.0)
185report_gnu_extensionsnoReport use of GNU extensions
186short_enums / no_short_enumsnoUse minimal-size enum representation
187clang / no_clangnoClang compatibility mode
188clang_versionyesClang version number (default 90100 = Clang 9.1.0)
189strict_gnu / no_strict_gnunoStrict GNU mode
190db_nameyesDebug database name
191long_longnoAllow long long type
192context_limityesMaximum template instantiation context depth
193set_flag / clear_flagyesRaw flag manipulation via off_D47CE0 lookup table
194edg_base_diryesEDG base directory (error on invalid path)
195embedded_c / no_embedded_cnoEmbedded C mode (not relevant to CUDA)
196thread_local_storage / no_thread_local_storagenothread_local support
197trigraphs / no_trigraphsnoTrigraph processing (default on)
198nonstd_default_arg_deduction / no_nonstd_default_arg_deductionno--
199stdc_zero_in_system_headers / no_stdc_zero_in_system_headersno--
200template_typedefs_in_diagnostics / no_template_typedefs_in_diagnosticsno--
202uliterals / no_uliteralsnoUnicode literals (u"", U"", u8"")
203type_traits_helpers / no_type_traits_helpersnoIntrinsic type traits
204c++11 / c++0xnoC++11 mode (sets dword_126EF68 to 201103 or 199711)
205list_macrosnoList all defined macros after preprocessing
206dump_configurationnoDump full compiler configuration
207dump_legacy_as_targetyesDump legacy configuration in target format
208signed_bit_fields / unsigned_bit_fieldsnoDefault bit-field signedness
210check_concatenations / no_check_concatenationsnoString literal concatenation checks
211unicode_source_kindyesSource encoding: UTF-8=1, UTF-16LE=2, UTF-16BE=3, none=0
212lambdas / no_lambdasnoC++ lambda expressions
213rvalue_refs / no_rvalue_refsnoRvalue references
214rvalue_ctor_is_copy_ctor / rvalue_ctor_is_not_copy_ctornoRvalue constructor treatment
215gen_move_operations / no_gen_move_operationsnoImplicit move constructor/assignment (default on)
216auto_type / no_auto_typenoC++11 auto type deduction
217auto_storage / no_auto_storagenoauto as storage class (C++03 meaning)
218nonstd_instantiation_lookup / no_nonstd_instantiation_lookupno--
219nullptr / no_nullptrnonullptr keyword
220gcc89_inliningnoGCC 8.9-era inlining behavior
221nonstd_gnu_keywords / no_nonstd_gnu_keywordsnoGNU extension keywords
222default_nocommon_tentative_definitions / default_common_tentative_definitionsnoTentative definition linkage
223no_token_separators_in_pp_outputno--
224c23_typeof / no_c23_typeofnoC23 typeof operator
225c++11_sfinae / no_c++11_sfinaenoC++11 SFINAE rules
226c++11_sfinae_ignore_access / no_c++11_sfinae_ignore_accessnoIgnore access checks in SFINAE
227variadic_templates / no_variadic_templatesnoParameter packs and pack expansion
228c++03noC++03 mode (sets dword_126EF68 to 199711)
229func_prototype_tags / no_func_prototype_tagsno--
230implicit_noexcept / no_implicit_noexceptnoImplicit noexcept on destructors
231unrestricted_unions / no_unrestricted_unionsnoUnrestricted unions (C++11)
232max_depth_constexpr_callyesMaximum constexpr recursion depth (default 200)
233max_cost_constexpr_callyesMaximum constexpr evaluation cost (default 256)
234delegating_constructors / no_delegating_constructorsno--
235lossy_conversion_warning / no_lossy_conversion_warningno--
236deprecated_string_conv / no_deprecated_string_convnoDeprecated string literal to char* conversion
237user_defined_literals / no_user_defined_literalsnoUDL support
238preserve_lvalues_with_same_type_casts / no_...no--
239nonstd_anonymous_unions / no_nonstd_anonymous_unionsno--

Late C++/Architecture/Output Flags (240--258)

IDFlagArgEffect
240c++14noC++14 mode (sets dword_126EF68 to 201402)
241c11noC11 mode (sets dword_126EF68 to 201112)
242c17noC17 mode (sets dword_126EF68 to 201710)
243c23noC23 mode (sets dword_126EF68 to 202311)
244digit_separators / no_digit_separatorsnoC++14 digit separators (1'000'000)
245targetyesSM architecture string, parsed via sub_7525E0 into dword_126E4A8
246c++17noC++17 mode (sets dword_126EF68 to 201703)
247utf8_char_literals / no_utf8_char_literalsnoUTF-8 character literal support
248stricter_template_checkingnoAdditional template constraint checks
249exc_spec_in_func_type / no_exc_spec_in_func_typenoException spec as part of function type (C++17)
250aligned_new / no_aligned_newnoAligned operator new (C++17)
251c++20noC++20 mode (sets dword_126EF68 to 202002)
252c++23noC++23 mode (sets dword_126EF68 to 202302)
253ms_std_preprocessor / no_ms_std_preprocessornoMSVC standard preprocessor mode
268partial-linknoPartial linking mode
273dump_command_optionsnoPrint all registered flag names
274output_modeyesOutput format: text (0) or sarif (1)
275incognito / no_incognitonoIncognito mode

Note: Many IDs in the 240-252 range serve double duty as both C/C++ standard selectors and feature toggles. The standard selection IDs are also cross-referenced in the Language Standard Selection section above.

Inline-Registered Paired Flags

Seven additional paired flags are registered through inline table population rather than calls to register_command_flag. They share the same entry structure but are populated directly into the array:

FlagEffect
relaxed_abstract_checking / no_relaxed_abstract_checkingRelax abstract class checks
concepts / no_conceptsC++20 concepts support
colors / no_colorsColorized diagnostic output
keep_restrict_in_signatures / no_keep_restrict_in_signaturesPreserve restrict in mangled names
check_unicode_security / no_check_unicode_securityUnicode security checks (homoglyph detection)
old_id_chars / no_old_id_charsLegacy identifier character rules
add_match_notes / no_add_match_notesAdd notes about matching overloads

Language Standard Selection

Six language standard flags set dword_126EF68 (the internal __cplusplus / __STDC_VERSION__ value) and trigger corresponding mode changes:

C Standards

IDFlagdword_126EF68 valueC standard
7old_c(K&R)Pre-ANSI C via set_c_mode(1)
179c89198912ANSI C / C89
178c99199901C99
241c11201112C11
242c17201710C17
243c23202311C23

C++ Standards

IDFlagdword_126EF68 valueC++ standard
228c++03199711C++98/03 (also aliased as c++98 via --c++11 flag ID 204 with conditional)
204c++11201103C++11 (sets 199711 if dword_E7FF14 is unset or C mode)
240c++14201402C++14
246c++17201703C++17
251c++20202002C++20
252c++23202302C++23

When a C++ standard is selected, the post-parsing dialect resolution logic automatically enables the corresponding feature flags. For example, selecting --c++11 (value 201103) enables lambdas, rvalue references, auto type deduction, nullptr, variadic templates, and other C++11 features. The resolution logic also interacts with GCC/Clang version thresholds to determine which extensions are available.

Diagnostic Control Flags

The five diag_* flags (IDs 39--43) accept comma-separated lists of diagnostic numbers. The parser strips whitespace, splits on commas, and calls sub_4ED400(number, severity, 1) for each number:

--diag_suppress=1234,5678       # suppress errors 1234 and 5678
--diag_warning=20001            # demote CUDA error 20001 to warning
--diag_error=111                # promote diagnostic 111 to error
--diag_remark=185               # demote diagnostic 185 to remark
--diag_once=175                 # emit diagnostic 175 only once

The error number system is documented in Diagnostic System Overview. Numbers above 3456 in the internal range correspond to the 20000-series CUDA errors via the offset formula display_code = internal_code + 16543.

Post-Parsing Dialect Resolution

After the main parsing loop completes, proc_command_line executes a large block of dialect resolution logic that:

  1. Resolves host compiler mode conflicts -- If both --gcc and --clang are set, or --cfront_2.1 is combined with modern modes, the resolution picks one and adjusts feature flags accordingly
  2. Sets C++ feature flags from __cplusplus version -- Based on the value in dword_126EF68:
    • 199711 (C++98/03): baseline features only
    • 201103 (C++11): enables lambdas, rvalue refs, auto, nullptr, variadic templates, range-based for, delegating constructors, unrestricted unions, user-defined literals
    • 201402 (C++14): adds digit separators, generic lambdas, relaxed constexpr
    • 201703 (C++17): adds aligned new, exception spec in function type, structured bindings
    • 202002 (C++20): adds concepts, modules, coroutines
    • 202302 (C++23): adds latest features
  3. Applies GCC version thresholds -- When in GCC compatibility mode, certain features are gated on the GCC version number stored in qword_126EF98 (default 80100 = GCC 8.1.0). Known thresholds:
    • 40299 (0x9D6B): GCC 4.2
    • 40599 (0x9E97): GCC 4.5
    • 40699 (0x9EFB): GCC 4.6
    • Higher versions enable progressively more features
  4. Opens output files -- Error output, listing file, output file
  5. Processes the input filename -- The remaining non-flag argv entry

Key Globals After Resolution

GlobalTypeContent
dword_126EF68int32__cplusplus / __STDC_VERSION__ value
dword_126EFB4int32Language mode: 0=unset, 1=C, 2=C++
dword_126EFA8int32GCC compatibility enabled
dword_126EFA4int32Clang compatibility enabled
qword_126EF98int64GCC version (default 80100)
qword_126EF90int64Clang version (default 90100)
dword_126EFB0int32GNU extensions enabled
dword_126EFACint32Clang extensions enabled
dword_126E4A8int32SM architecture code (from --target)
dword_126E1D4int32MSVC target version

The set_flag / clear_flag Mechanism

Flag ID 199 (--set_flag / --clear_flag) provides a raw escape hatch. The argument is a flag name looked up in the off_D47CE0 table -- an array of {name, global_address} pairs. If the name is found, the corresponding global variable is set to the provided integer value (--set_flag=name=value) or cleared to 0 (--clear_flag=name). This mechanism allows nvcc to toggle internal EDG configuration flags that do not have dedicated CLI flag registrations.

Default Values

default_init (sub_45EB40) runs before proc_command_line and initializes approximately 350 global configuration variables. Notable non-zero defaults:

GlobalDefaultMeaning
dword_106C2101Exceptions enabled
dword_106C1801RTTI enabled
dword_106C1781bool is keyword
dword_106C1941Namespaces enabled
dword_106C19C1Argument-dependent lookup enabled
dword_106C1A01Class name injection enabled
dword_106C1A41String literals are const
dword_106C1881wchar_t is keyword
dword_106C18C1Alternative tokens enabled
dword_106C1401Compound literals allowed
dword_106C1381Dependent name processing enabled
dword_106C1341Template parsing enabled
dword_106C12C1Friend injection enabled
dword_106BDB81restrict enabled
dword_106BDB01Remove unneeded entities enabled
dword_106BD981Trigraphs enabled
dword_106BD681Guiding declarations allowed
dword_106BD581Old specializations allowed
dword_106BD541Implicit typename enabled
dword_106BE841Generate move operations enabled
dword_106C0641Stack limit modification enabled
qword_106BD10200Max constexpr recursion depth
qword_106BD08256Max constexpr evaluation cost
qword_126EF9880100Default GCC version (8.1.0)
qword_126EF9090100Default Clang version (9.1.0)
qword_126EF781926MSVC version threshold
qword_126EF7099999Some upper bound sentinel

Conflict Detection

Before the main parsing loop, check_conflicting_flags (sub_451E80) verifies that flags 3, 193, 194, and 195 (no_line_commands, set_flag, clear_flag, and related flags) are not used in conflicting combinations. If any conflict is detected, error 1027 is emitted.

Inline-Registered Flags

init_command_line_flags (sub_452010) calls the helper register_command_flag (sub_451F80) 251 times. The remaining 25 flags reach the dispatch switch through a different mechanism: their entries are populated directly into the flag table by inline stores in sub_452010, bypassing the helper. The two registration mechanisms produce identical runtime behavior -- both result in valid 40-byte entries in the table at dword_E80060. The distinction is purely a source-level artifact, and is invisible from the parser's perspective.

There are two reasons a flag is registered inline rather than through the helper:

  1. Non-default field values. A few flags need a visible bit pattern, target compiler gate, or takes_value configuration that the helper's parameter signature cannot express. Inline stores write the exact byte layout required.

  2. Paired-toggle positive forms. EDG paired toggles register both --no_flag and --flag to the same case_id. The --no_flag form is registered through the helper (with the negation-bit flag set). The --flag form is then written inline at the next table slot, sharing the same case_id but with the negation bit clear. The helper is not parameterized for the positive form, so inline stores are the only path.

The 25 inline-registered flag names:

CategoryFlag names
Standalone (no paired toggle)target, version, output_mode, m32, m64, db, icc, icx, grco, uumn, incognito, xref, list
Positive forms of paired toggleslambdas, modules, restrict, bool, rtti, namespaces, trigraphs, nullptr, c++11, c++0x, c89, c99, c11, c17, c23

(The first row has 13 names; the second row has 15 names. The "approximately 25" figure in the wave-58 audit collapses the two c++11/c++0x aliases that share case ID 247, and the c11/c17/c23 C-mode group that share case IDs, into the lower count.)

Sum check. 251 helper-registered flag entries + 25 inline-registered entries = 276 total table entries. This matches the init_command_line_flags exit state, the dword_E80058 counter immediately after registration completes, and the case count of the dispatch switch (IDs 0..275, where 0 is the default sink for unrecognized command-line input).

Version Banners

The CLI exposes two distinct version-printing flags. They share the substring "version" in their long name but differ in case (--version vs. --Version), parsing case ID, format of the printed banner, and whether they terminate the process.

--version (ID 21, -v) -- continues execution

Case 21 in the dispatch switch writes the cudafe++ version banner to stdout via the standard print path, then falls through to the next iteration of the argv loop without calling exit(). The flag is used by build systems that want to query the compiler version while still receiving any subsequent flags on the same command line.

Banner format (lines emitted in order):

cudafe: NVIDIA (R) Cuda Language Front End
Portions Copyright (c) 2005, 2024-YYYY NVIDIA Corporation
Portions Copyright (c) 1988-2018, 2024 Edison Design Group Inc.
Based on Edison Design Group C/C++ Front End, version 6.6
Cuda compilation tools, release 13.0, V13.0.88

The YYYY placeholder is the current compilation year, sourced from a build-time constant. The string "1988-2018, 2024" is the EDG copyright span -- the two date ranges in a single Edison Design Group copyright line are an explicit choice by EDG, not a formatting artifact.

--Version (ID 92, -V) -- emits EDG portion banner then exit(1)

Case 92 in the dispatch switch writes a single EDG-portion copyright line to stdout, then calls exit(1) immediately. The exit code is the conventional "success but no compilation performed" value used by GNU-style version flags. No further argv processing occurs after case 92.

Banner string (single line):

Portions Copyright (c) 1988-2016 Edison Design Group Inc.

The "1988-2016" date range here is intentionally different from the "1988-2018, 2024" range printed by --version: case 92 emits the older, narrower EDG portion-only copyright form. The two banners do not share string data in .rodata; they are two separate string literals, embedded at distinct addresses. Tooling that grep's the binary for either banner should not assume the other will be present in the same form.

Cross-References