cudafe++ reserves internal error indices 3457--3794 (338 slots) for CUDA-specific diagnostics. These are displayed to users as numbers 20000--20337 using the formula display = internal + 16543. Of the 338 slots, approximately 210 carry unique message templates; the remainder are reserved or share templates with parametric fill-ins. Every CUDA error can be controlled by its numeric code or diagnostic tag name via --diag_suppress, --diag_warning, --diag_error, or the #pragma nv_diagnostic system.
This page is a flat lookup table. For the diagnostic pipeline architecture (severity stack, pragma scoping, SARIF output), see Diagnostic Overview. For narrative discussion of each category with implementation details, see CUDA Errors.
Scope note. This catalog enumerates the CUDA-extension namespace (~248 documented tags spanning internal codes 3457--3794, plus 6 pragma-action tags listed for completeness). The full cudafe++ binary defines roughly 859 snake_case diagnostic tags across all of EDG 6.6's C++ frontend; the remaining ~605 tags govern standard C++ diagnostics (templates, constexpr evaluation, modules, attributes, type mismatches, overload ambiguity, etc.) inherited from EDG and shared with the wider EDG ecosystem. Those tags use the original error-code range 0--3456 and are accepted by --diag_suppress / #pragma nv_diag_suppress identically to CUDA tags but are not enumerated here. See EDG Diagnostic Namespace at the bottom of this page for a category-prefix breakdown. Confidence: HIGH for the 859 total (counted directly from cudafe_strings.json), MED for the 248 / 605 split (prefix-based bucketing).
User-visible: file.cu(42): error #20042-D: calling a __device__ function from ...
^^^^^
display code = internal + 16543
| Direction | Formula | Example |
|---|
| Display to internal | internal = display - 16543 | 20042 maps to internal 3499 |
| Internal to display | display = internal + 16543 | 3457 maps to display 20000 |
The -D suffix appears when severity is 7 or below (note, remark, warning, soft error). Hard errors (severity 8+) omit the suffix.
| Code | Level | Suppressible |
|---|
| 2 | note | yes |
| 4 | remark | yes |
| 5 | warning | yes |
| 6 | command-line warning | no |
| 7 | error (soft) | yes |
| 8 | error (hard, from pragma) | no |
| 9 | catastrophic error | no |
| 10 | command-line error | no |
| 11 | internal error | no |
Suppress by numeric code:
nvcc --diag_suppress=20042
Suppress by tag name:
nvcc --diag_suppress=unsafe_device_call
In source code:
#pragma nv_diag_suppress unsafe_device_call
#pragma nv_diag_suppress 20042
Checks performed by the call-graph walker comparing the execution-space byte at entity offset +182 of caller vs. callee.
| Tag | Sev | Message |
|---|
unsafe_device_call | W | calling a __device__ function(%sq1) from a __host__ function(%sq2) is not allowed |
unsafe_device_call | W | calling a __device__ function(%sq1) from a __host__ __device__ function(%sq2) is not allowed |
unsafe_device_call | W | calling a __host__ function(%sq1) from a __device__ function(%sq2) is not allowed |
unsafe_device_call | W | calling a __host__ function(%sq1) from a __global__ function(%sq2) is not allowed |
unsafe_device_call | W | calling a __host__ function(%sq1) from a __host__ __device__ function(%sq2) is not allowed |
unsafe_device_call | W | calling a __host__ function from a __host__ __device__ function is not allowed |
These fire when --expt-relaxed-constexpr is not enabled.
| Tag | Sev | Message |
|---|
unsafe_device_call | W | calling a constexpr __device__ function(%sq1) from a __host__ function(%sq2) is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this. |
unsafe_device_call | W | calling a constexpr __device__ function(%sq1) from a __host__ __device__ function(%sq2) is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this. |
unsafe_device_call | W | calling a constexpr __host__ function(%sq1) from a __device__ function(%sq2) is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this. |
unsafe_device_call | W | calling a constexpr __host__ function(%sq1) from a __global__ function(%sq2) is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this. |
unsafe_device_call | W | calling a constexpr __host__ function(%sq1) from a __host__ __device__ function(%sq2) is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this. |
unsafe_device_call | W | calling a constexpr __host__ function from a __host__ __device__ function is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this. |
Override checker (sub_432280) extracts the 0x30 mask from the execution-space byte. __global__ is excluded because kernels cannot be virtual.
| Tag | Sev | Message |
|---|
| -- | E | execution space mismatch: overridden entity (%n1) is a __device__ function, but overriding entity (%n2) is a __host__ function |
| -- | E | execution space mismatch: overridden entity (%n1) is a __device__ function, but overriding entity (%n2) is a __host__ __device__ function |
| -- | E | execution space mismatch: overridden entity (%n1) is a __host__ function, but overriding entity (%n2) is a __device__ function |
| -- | E | execution space mismatch: overridden entity (%n1) is a __host__ function, but overriding entity (%n2) is a __host__ __device__ function |
| -- | E | execution space mismatch: overridden entity (%n1) is a __host__ __device__ function, but overriding entity (%n2) is a __device__ function |
| -- | E | execution space mismatch: overridden entity (%n1) is a __host__ __device__ function, but overriding entity (%n2) is a __host__ function |
Checked in decl_routine (sub_4CE420) and check_cuda_attribute_consistency (sub_4C6D50).
| Tag | Sev | Message |
|---|
device_function_redeclared_with_global | E | a __device__ function(%no1) redeclared with __global__ |
global_function_redeclared_with_device | E | a __global__ function(%no1) redeclared with __device__ |
global_function_redeclared_with_host | E | a __global__ function(%no1) redeclared with __host__ |
global_function_redeclared_with_host_device | E | a __global__ function(%no1) redeclared with __host__ __device__ |
global_function_redeclared_without_global | E | a __global__ function(%no1) redeclared without __global__ |
host_function_redeclared_with_global | E | a __host__ function(%no1) redeclared with __global__ |
host_device_function_redeclared_with_global | E | a __host__ __device__ function(%no1) redeclared with __global__ |
| Tag | Sev | Message |
|---|
device_function_redeclared_with_host | W | a __device__ function(%no1) redeclared with __host__, hence treated as a __host__ __device__ function |
device_function_redeclared_with_host_device | W | a __device__ function(%no1) redeclared with __host__ __device__, hence treated as a __host__ __device__ function |
device_function_redeclared_without_device | W | a __device__ function(%no1) redeclared without __device__, hence treated as a __host__ __device__ function |
host_function_redeclared_with_device | W | a __host__ function(%no1) redeclared with __device__, hence treated as a __host__ __device__ function |
host_function_redeclared_with_host_device | W | a __host__ function(%no1) redeclared with __host__ __device__, hence treated as a __host__ __device__ function |
| Tag | Sev | Message |
|---|
global_function_return_type | E | a __global__ function must have a void return type |
global_function_deduced_return_type | E | a __global__ function must not have a deduced return type |
global_function_has_ellipsis | E | a __global__ function cannot have ellipsis |
global_rvalue_ref_type | E | a __global__ function cannot have a parameter with rvalue reference type |
global_ref_param_restrict | E | a __global__ function cannot have a parameter with __restrict__ qualified reference type |
global_va_list_type | E | A __global__ function or function template cannot have a parameter with va_list type |
global_function_with_initializer_list | E | a __global__ function or function template cannot have a parameter with type std::initializer_list |
global_param_align_too_big | E | cannot pass a parameter with a too large explicit alignment to a __global__ function on win32 platforms |
| Tag | Sev | Message |
|---|
global_class_decl | E | A __global__ function or function template cannot be a member function |
global_friend_definition | E | A __global__ function or function template cannot be defined in a friend declaration |
global_function_in_unnamed_inline_ns | E | A __global__ function or function template cannot be declared within an inline unnamed namespace |
global_operator_function | E | An operator function cannot be a __global__ function |
global_new_or_delete | E | (__global__ on operator new/delete) |
| -- | E | function main cannot be marked __device__ or __global__ |
| Tag | Sev | Message |
|---|
global_function_constexpr | E | A __global__ function or function template cannot be marked constexpr |
global_function_consteval | E | A __global__ function or function template cannot be marked consteval |
global_function_inline | E | (__global__ with inline) |
global_exception_spec | E | An exception specification is not allowed for a __global__ function or function template |
| Tag | Sev | Message |
|---|
global_private_type_arg | E | A type that is defined inside a class and has private or protected access (%t) cannot be used in the template argument type of a __global__ function template instantiation, unless the class is local to a __device__ or __global__ function |
global_private_template_arg | E | A template that is defined inside a class and has private or protected access cannot be used in the template template argument of a __global__ function template instantiation |
global_unnamed_type_arg | E | An unnamed type (%t) cannot be used in the template argument type of a __global__ function template instantiation, unless the type is local to a __device__ or __global__ function |
global_func_local_template_arg | E | A type defined inside a __host__ function (%t) cannot be used in the template argument type of a __global__ function template instantiation |
global_lambda_template_arg | E | The closure type for a lambda (%t%s) cannot be used in the template argument type of a __global__ function template instantiation, unless the lambda is defined within a __device__ or __global__ function, or the flag '-extended-lambda' is specified and the lambda is an extended lambda |
local_type_used_in_global_function | W | a local type %t (defined in %sq1) used in global function %sq2 template argument, the global function cannot be launched from host code. |
| Tag | Sev | Message |
|---|
variable_template_private_type_arg | E | (private/protected type in variable template instantiation) |
variable_template_private_template_arg | E | (private template template arg in variable template) |
variable_template_unnamed_type_template_arg | E | An unnamed type (%t) cannot be used in the template argument type of a variable template instantiation, unless the type is local to a __device__ or __global__ function |
variable_template_func_local_template_arg | E | A type defined inside a __host__ function (%t) cannot be used in the template argument type of a variable template instantiation |
variable_template_lambda_template_arg | E | The closure type for a lambda (%t%s) cannot be used in the template argument type of a variable template instantiation, unless the lambda is defined within a __device__ or __global__ function, or the lambda is an 'extended lambda' and the flag --extended-lambda is specified |
| Tag | Sev | Message |
|---|
global_function_multiple_packs | E | Multiple pack parameters are not allowed for a variadic __global__ function template |
global_function_pack_not_last | E | Pack template parameter must be the last template parameter for a variadic __global__ function template |
| Tag | Sev | Message |
|---|
bounds_attr_only_on_global_func | E | %s is only allowed on a __global__ function |
maxnreg_attr_only_on_global_func | E | (__maxnreg__ only on __global__) |
missing_launch_bounds | W | no __launch_bounds__ specified for __global__ function |
cuda_specifier_twice_in_group | E | (duplicate CUDA specifier on same declaration) |
bounds_maxnreg_incompatible_qualifiers | E | (__launch_bounds__ and __maxnreg__ conflict) |
| -- | E | The %s qualifiers cannot be applied to the same kernel |
| -- | E | Multiple %s specifiers are not allowed |
| -- | E | incorrect value for launch bounds |
Extended lambdas (__device__ or __host__ __device__ lambdas in host code, enabled by --extended-lambda) must have closure types serializable for device transfer.
| Tag | Sev | Message |
|---|
extended_lambda_reference_capture | E | An extended %s lambda cannot capture variables by reference |
extended_lambda_pack_capture | E | An extended %s lambda cannot capture an element of a parameter pack |
extended_lambda_too_many_captures | E | An extended %s lambda can only capture up to 1023 variables |
extended_lambda_array_capture_rank | E | An extended %s lambda cannot capture an array variable (type: %t) with more than 7 dimensions |
extended_lambda_array_capture_assignable | E | An extended %s lambda cannot capture an array variable whose element type (%t) is not assignable on the host |
extended_lambda_array_capture_default_constructible | E | An extended %s lambda cannot capture an array variable whose element type (%t) is not default constructible on the host |
extended_lambda_init_capture_array | E | An extended %s lambda cannot init-capture variables with array type |
extended_lambda_init_capture_initlist | E | An extended %s lambda cannot have init-captures with type std::initializer_list |
extended_lambda_capture_in_constexpr_if | E | An extended %s lambda cannot first-capture variable in constexpr-if context |
this_addr_capture_ext_lambda | W | Implicit capture of 'this' in extended lambda expression |
extended_lambda_hd_init_capture | E | init-captures are not allowed for extended __host__ __device__ lambdas |
| -- | E | Unless enabled by language dialect, *this capture is only supported when the lambda is either __device__ only, or is defined within a __device__ or __global__ function |
| Tag | Sev | Message |
|---|
extended_lambda_capture_local_type | E | A type local to a function (%t) cannot be used in the type of a variable captured by an extended __device__ or __host__ __device__ lambda |
extended_lambda_capture_private_type | E | A type that is a private or protected class member (%t) cannot be used in the type of a variable captured by an extended __device__ or __host__ __device__ lambda |
extended_lambda_call_operator_local_type | E | A type local to a function (%t) cannot be used in the return or parameter types of the operator() of an extended __device__ or __host__ __device__ lambda |
extended_lambda_call_operator_private_type | E | A type that is a private or protected class member (%t) cannot be used in the return or parameter types of the operator() of an extended __device__ or __host__ __device__ lambda |
extended_lambda_parent_local_type | E | A type local to a function (%t) cannot be used in the template argument of the enclosing parent function (and any parent classes) of an extended __device__ or __host__ __device__ lambda |
extended_lambda_parent_private_type | E | A type that is a private or protected class member (%t) cannot be used in the template argument of the enclosing parent function (and any parent classes) of an extended __device__ or __host__ __device__ lambda |
extended_lambda_parent_private_template_arg | E | A template that is a private or protected class member cannot be used in the template argument of the enclosing parent function (and any parent classes) of an extended %s lambda |
| Tag | Sev | Message |
|---|
extended_lambda_enclosing_function_local | E | The enclosing parent function (%sq2) for an extended %s1 lambda must not be defined inside another function |
extended_lambda_enclosing_function_not_found | E | (no enclosing function found for extended lambda) |
extended_lambda_inaccessible_parent | E | The enclosing parent function (%sq2) for an extended %s1 lambda cannot have private or protected access within its class |
extended_lambda_enclosing_function_deducible | E | The enclosing parent function (%sq2) for an extended %s1 lambda must not have deduced return type |
extended_lambda_cant_take_function_address | E | The enclosing parent function (%sq2) for an extended %s1 lambda must allow its address to be taken |
extended_lambda_parent_non_extern | E | On Windows, the enclosing parent function (%sq2) for an extended %s1 lambda cannot have internal or no linkage |
extended_lambda_parent_class_unnamed | E | The enclosing parent function (%sq2) for an extended %s1 lambda cannot be a member function of a class that is unnamed |
extended_lambda_parent_template_param_unnamed | E | The enclosing parent function (%sq2) for an extended %s1 lambda cannot be in a template which has a unnamed parameter: %nd |
extended_lambda_nest_parent_template_param_unnamed | E | The enclosing parent %n for an extended %s lambda cannot be a template which has a unnamed parameter |
extended_lambda_multiple_parameter_packs | E | The enclosing parent template function (%sq2) for an extended %s1 lambda cannot have more than one variadic parameter, or it is not listed last in the template parameter list. |
extended_lambda_no_parent_func | E | (extended lambda has no parent function) |
extended_lambda_illegal_parent | E | (extended lambda in illegal parent context) |
| Tag | Sev | Message |
|---|
extended_lambda_enclosing_function_generic_lambda | E | An extended %s1 lambda cannot be defined inside a generic lambda expression(%sq2). |
extended_lambda_enclosing_function_hd_lambda | E | An extended %s1 lambda cannot be defined inside an extended __host__ __device__ lambda expression(%sq2). |
extended_lambda_inaccessible_ancestor | E | An extended %s1 lambda cannot be defined inside a class (%sq2) with private or protected access within another class |
extended_lambda_inside_constexpr_if | E | For this host platform/dialect, an extended lambda cannot be defined inside the 'if' or 'else' block of a constexpr if statement |
extended_lambda_multiple_parent | E | Cannot specify multiple __nv_parent directives in a lambda declaration |
extended_host_device_generic_lambda | E | __host__ __device__ extended lambdas cannot be generic lambdas |
| -- | E | If an extended %s lambda is defined within the body of one or more nested lambda expressions, each of these enclosing lambda expressions must be defined within the immediate or nested block scope of a function. |
| Tag | Sev | Message |
|---|
extended_lambda_disallowed | E | __host__ or __device__ annotation on lambda requires --extended-lambda nvcc flag |
extended_lambda_constexpr | E | The %s1 specifier is not allowed for an extended %s2 lambda |
lambda_operator_annotated | E | The operator() function for a lambda cannot be explicitly annotated with execution space annotations (__host__/__device__/__global__), the annotations are derived from its closure class |
extended_lambda_discriminator | E | (extended lambda discriminator collision) |
General restrictions that apply to all GPU-side code (__device__ and __global__ function bodies).
| Tag | Sev | Message |
|---|
cuda_device_code_unsupported_operator | E | The operator '%s' is not allowed in device code |
unsupported_type_in_device_code | E | %t %s1 a %s2, which is not supported in device code |
| -- | E | device code does not support exception handling |
no_coroutine_on_device | E | device code does not support coroutines |
| -- | E | operations on vector types are not supported in device code |
undefined_device_entity | E | cannot use an entity undefined in device code |
undefined_device_identifier | E | identifier %sq is undefined in device code |
thread_local_in_device_code | E | cannot use thread_local specifier for variable declarations in device code |
unrecognized_pragma_device_code | W | unrecognized #pragma in device code |
| -- | E | zero-sized parameter type %t is not allowed in device code |
| -- | E | zero-sized variable %sq is not allowed in device code |
| -- | E | dynamic initialization is not supported for a function-scope static %s variable within a __device__/__global__ function |
| -- | E | function-scope static variable within a __device__/__global__ function requires a memory space specifier |
use_of_virtual_base_on_compute_1x | E | Use of a virtual base (%t) requires the compute_20 or higher architecture |
| -- | E | alloca() is not supported for architectures lower than compute_52 |
| Tag | Sev | Message |
|---|
device_launch_no_sepcomp | E | kernel launch from __device__ or __global__ functions requires separate compilation mode |
missing_api_for_device_side_launch | E | device-side kernel launch could not be processed as the required runtime APIs are not declared |
| -- | W | explicit stream argument not provided in kernel launch |
| -- | E | kernel launches from templates are not allowed in system files |
device_side_launch_arg_with_user_provided_cctor | E | cannot pass an argument with a user-provided copy-constructor to a device-side kernel launch |
device_side_launch_arg_with_user_provided_dtor | E | cannot pass an argument with a user-provided destructor to a device-side kernel launch |
| Tag | Sev | Message |
|---|
device_var_read_in_host | E | a %s1 %n1 cannot be directly read in a host function |
device_var_written_in_host | E | a %s1 %n1 cannot be directly written in a host function |
device_var_address_taken_in_host | E | address of a %s1 %n1 cannot be directly taken in a host function |
host_var_read_in_device | E | a host %n1 cannot be directly read in a device function |
host_var_written_in_device | E | a host %n1 cannot be directly written in a device function |
host_var_address_taken_in_device | E | address of a host %n1 cannot be directly taken in a device function |
| Tag | Sev | Message |
|---|
illegal_local_to_device_function | E | %s1 %sq2 variable declaration is not allowed inside a device function body |
illegal_local_to_host_function | E | %s1 %sq2 variable declaration is not allowed inside a host function body |
shared_specifier_in_range_for | E | the __shared__ memory space specifier is not allowed for a variable declared by the for-range-declaration |
bad_shared_storage_class | E | __shared__ variables cannot have external linkage |
device_variable_in_unnamed_inline_ns | E | A %s variable cannot be declared within an inline unnamed namespace |
| -- | E | member variables of an anonymous union at global or namespace scope cannot be directly accessed in __device__ and __global__ functions |
shared_inside_struct | E | shared type inside a struct or union is not allowed |
shared_parameter | E | (__shared__ as function parameter) |
| Tag | Sev | Message |
|---|
auto_device_fn_ref | E | A non-constexpr __device__ function (%sq1) with "auto" deduced return type cannot be directly referenced %s2, except if the reference is absent when __CUDA_ARCH__ is undefined |
device_var_constexpr | E | (constexpr rules for __device__ variables) |
device_var_structured_binding | E | (structured bindings on __device__ variables) |
The __grid_constant__ annotation (compute_70+) marks a kernel parameter as read-only grid-wide.
| Tag | Sev | Message |
|---|
grid_constant_non_kernel | E | __grid_constant__ annotation is only allowed on a parameter of a __global__ function |
grid_constant_not_const | E | a parameter annotated with __grid_constant__ must have const-qualified type |
grid_constant_reference_type | E | a parameter annotated with __grid_constant__ must not have reference type |
grid_constant_unsupported_arch | E | __grid_constant__ annotation is only allowed for architecture compute_70 or later |
grid_constant_incompat_redecl | E | incompatible __grid_constant__ annotation for parameter %s in function redeclaration (see previous declaration %p) |
grid_constant_incompat_templ_redecl | E | incompatible __grid_constant__ annotation for parameter %s in function template redeclaration (see previous declaration %p) |
grid_constant_incompat_specialization | E | incompatible __grid_constant__ annotation for parameter %s in function specialization (see previous declaration %p) |
grid_constant_incompat_instantiation_directive | E | incompatible __grid_constant__ annotation for parameter %s in instantiation directive (see previous declaration %p) |
JIT mode (-dc for device-only compilation) restricts host constructs.
| Tag | Sev | Message |
|---|
no_host_in_jit | E | A function explicitly marked as a __host__ function is not allowed in JIT mode |
unannotated_function_in_jit | E | A function without execution space annotations (__host__/__device__/__global__) is considered a host function, and host functions are not allowed in JIT mode. Consider using -default-device flag to process unannotated functions as __device__ functions in JIT mode |
unannotated_variable_in_jit | E | A namespace scope variable without memory space annotations (__device__/__constant__/__shared__/__managed__) is considered a host variable, and host variables are not allowed in JIT mode. Consider using -default-device flag to process unannotated namespace scope variables as __device__ variables in JIT mode |
unannotated_static_data_member_in_jit | E | A class static data member with non-const type is considered a host variable, and host variables are not allowed in JIT mode. Consider using -default-device flag to process such data members as __device__ variables in JIT mode |
host_closure_class_in_jit | E | The execution space for the lambda closure class members was inferred to be __host__ (based on context). This is not allowed in JIT mode. Consider using -default-device to infer __device__ execution space for namespace scope lambda closure classes. |
| Tag | Sev | Message |
|---|
| -- | E | An inline __device__/__constant__/__managed__ variable must have internal linkage when the program is compiled in whole program mode (-rdc=false) |
template_global_no_def | E | when "-static-global-template-stub=true" in whole program compilation mode ("-rdc=false"), a __global__ function template instantiation or specialization (%sq) must have a definition in the current translation unit |
extern_kernel_template | E | when "-static-global-template-stub=true", extern __global__ function template is not supported in whole program compilation mode ("-rdc=false") |
| -- | W | address of internal linkage device function (%sq) was taken (nv bug 2001144). mitigation: no mitigation required if the address is not used for comparison, or if the target function is not a CUDA C++ builtin |
CUDA atomics lowered to PTX instructions with size, type, scope, and memory-order constraints.
| Tag | Sev | Message |
|---|
nv_atomic_functions_not_supported_below_sm60 | E | __nv_atomic_* functions are not supported on arch < sm_60. |
nv_atomic_operation_not_in_device_function | E | atomic operations are not in a device function. |
nv_atomic_function_no_args | E | atomic function requires at least one argument. |
nv_atomic_function_address_taken | E | nv atomic function must be called directly. |
invalid_nv_atomic_operation_size | E | atomic operations and, or, xor, add, sub, min and max are valid only on objects of size 4, or 8. |
invalid_nv_atomic_cas_size | E | atomic CAS is valid only on objects of size 2, 4, 8 or 16 bytes. |
invalid_nv_atomic_exch_size | E | atomic exchange is valid only on objects of size 4, 8 or 16 bytes. |
invalid_data_size_for_nv_atomic_generic_function | E | generic nv atomic functions are valid only on objects of size 1, 2, 4, 8 and 16 bytes. |
non_integral_type_for_non_generic_nv_atomic_function | E | non-generic nv atomic load, store, cas and exchange are valid only on integral types. |
invalid_nv_atomic_operation_add_sub_size | E | atomic operations add and sub are not valid on signed integer of size 8. |
nv_atomic_add_sub_f64_not_supported | W | atomic add and sub for 64-bit float is supported on architecture sm_60 or above. |
invalid_nv_atomic_operation_max_min_float | E | atomic operations min and max are not supported on any floating-point types. |
floating_type_for_logical_atomic_operation | E | For a logical atomic operation, the first argument cannot be any floating-point types. |
nv_atomic_cas_b16_not_supported | E | 16-bit atomic compare-and-exchange is supported on architecture sm_70 or above. |
nv_atomic_exch_cas_b128_not_supported | E | 128-bit atomic exchange or compare-and-exchange is supported on architecture sm_90 or above. |
nv_atomic_load_store_b128_version_too_low | E | 128-bit atomic load and store are supported on architecture sm_70 or above. |
| Tag | Sev | Message |
|---|
nv_atomic_load_order_error | E | atomic load's memory order cannot be release or acq_rel. |
nv_atomic_store_order_error | E | atomic store's memory order cannot be consume, acquire or acq_rel. |
nv_atomic_operation_order_not_constant_int | E | atomic operation's memory order argument is not an integer literal. |
nv_atomic_operation_scope_not_constant_int | E | atomic operation's scope argument is not an integer literal. |
invalid_nv_atomic_memory_order_value | E | (invalid memory order enum value) |
invalid_nv_atomic_thread_scope_value | E | (invalid thread scope enum value) |
| Tag | Sev | Message |
|---|
nv_atomic_operations_scope_fallback_to_membar | W | atomic operations' scope argument is supported on architecture sm_60 or above. Fall back to use membar. |
nv_atomic_operations_memory_order_fallback_to_membar | W | atomic operations' argument of memory order is supported on architecture sm_70 or above. Fall back to use membar. |
nv_atomic_operations_scope_cluster_change_to_device | W | atomic operations' scope of cluster is supported on architecture sm_90 or above. Using device scope instead. |
nv_atomic_load_store_scope_cluster_change_to_device | W | atomic load and store's scope of cluster is supported on architecture sm_90 or above. Using device scope instead. |
NVPTX backend supports fewer inline-assembly constraint letters than x86.
| Tag | Sev | Message |
|---|
asm_constraint_letter_not_allowed_in_device | E | asm constraint letter '%s' is not allowed inside a __device__/__global__ function |
asm_constraint_must_have_single_letter | E | an asm operand may specify only one constraint letter in a __device__/__global__ function |
| -- | E | The 'C' constraint can only be used for asm statements in device code |
cc_clobber_in_device | E | The cc clobber constraint is not supported in device code |
cuda_xasm_strict_placeholder_format | E | (strict placeholder format in CUDA asm) |
addr_of_label_in_device_func | E | address of label extension is not supported in __device__/__global__ functions |
Controls calling convention for device functions, adjusting parameter passing to match PTX ABI.
| Tag | Sev | Message |
|---|
nv_abi_pragma_bad_format | E | (malformed #pragma nv_abi) |
nv_abi_pragma_invalid_option | E | #pragma nv_abi contains an invalid option |
nv_abi_pragma_missing_arg | E | #pragma nv_abi requires an argument |
nv_abi_pragma_duplicate_arg | E | #pragma nv_abi contains a duplicate argument |
nv_abi_pragma_not_constant | E | #pragma nv_abi argument must evaluate to an integral constant expression |
nv_abi_pragma_not_positive_value | E | #pragma nv_abi argument value must be a positive value |
nv_abi_pragma_overflow_value | E | #pragma nv_abi argument value exceeds the range of an integer |
nv_abi_pragma_device_function | E | #pragma nv_abi must be applied to device functions |
nv_abi_pragma_device_function_context | E | #pragma nv_abi is not supported inside a host function |
nv_abi_pragma_next_construct | E | #pragma nv_abi must appear immediately before a function declaration, function definition, or an expression statement |
Forces all parameters to be passed in registers (compute_80+).
| Tag | Sev | Message |
|---|
register_params_not_enabled | E | __nv_register_params__ support is not enabled |
register_params_unsupported_arch | E | __nv_register_params__ is only supported for compute_80 or later architecture |
register_params_unsupported_function | E | __nv_register_params__ is not allowed on a %s function |
register_params_ellipsis_function | E | __nv_register_params__ is not allowed on a function with ellipsis |
__CUDACC_RTC__name_expr forms the mangled name of a __global__ function or __device__/__constant__ variable at compile time.
| Tag | Sev | Message |
|---|
name_expr_parsing | E | Error in parsing name expression for lowered name lookup. Input name expression was: %sq |
name_expr_non_global_routine | E | Name expression cannot form address of a non-__global__ function. Input name expression was: %sq |
name_expr_non_device_variable | E | Name expression cannot form address of a variable that is not a __device__/__constant__ variable. Input name expression was: %sq |
name_expr_not_routine_or_variable | E | Name expression must form address of a __global__ function or the address of a __device__/__constant__ variable. Input name expression was: %sq |
name_expr_extra_tokens | E | Extra tokens found after parsing name expression for lowered name lookup. Input name expression was: %sq |
name_expr_internal_error | E | Internal error in parsing name expression for lowered name lookup. Input name expression was: %sq |
| Tag | Sev | Message |
|---|
texture_surface_variable_in_unnamed_inline_ns | E | A texture or surface variable cannot be declared within an inline unnamed namespace |
| -- | E | A texture or surface variable cannot be used in the non-type template argument of a __device__, __host__ __device__ or __global__ function template instantiation |
reference_to_text_surf_type_in_device_func | E | a reference to texture/surface type cannot be used in __device__/__global__ functions |
reference_to_text_surf_var_in_device_func | E | taking reference of texture/surface variable not allowed in __device__/__global__ functions |
addr_of_text_surf_var_in_device_func | E | cannot take address of texture/surface variable %sq in __device__/__global__ functions |
addr_of_text_surf_expr_in_device_func | E | cannot take address of texture/surface expression in __device__/__global__ functions |
indir_into_text_surf_var_in_device_func | E | indirection not allowed for accessing texture/surface through variable %sq in __device__/__global__ functions |
indir_into_text_surf_expr_in_device_func | E | indirection not allowed for accessing texture/surface through expression in __device__/__global__ functions |
| Tag | Sev | Message |
|---|
managed_const_type_not_allowed | E | a __managed__ variable cannot have a const qualified type |
managed_reference_type_not_allowed | E | a __managed__ variable cannot have a reference type |
managed_cant_be_shared_constant | E | __managed__ variables cannot be marked __shared__ or __constant__ |
unsupported_arch_for_managed_capability | E | __managed__ variables require architecture compute_30 or higher |
unsupported_configuration_for_managed_capability | E | __managed__ variables are not yet supported for this configuration (compilation mode (32/64 bit) and/or target operating system) |
decltype_of_managed_variable | E | A __managed__ variable cannot be used as an unparenthesized id-expression argument for decltype() |
| Tag | Sev | Message |
|---|
device_function_has_ellipsis | E | __device__ or __host__ __device__ function with ellipsis requires compute_30 or higher architecture |
device_func_tex_arg | E | (device function with texture argument restriction) |
no_host_device_initializer_list | E | (std::initializer_list in __host__ __device__ context) |
no_host_device_move_forward | E | (std::move/forward in __host__ __device__ context) |
no_strict_cuda_error | W | (relaxed error checking mode) |
Warp Group Matrix Multiply-Accumulate builtins (sm_90a+).
| Tag | Sev | Message |
|---|
wgmma_mma_async_not_enabled | E | __wgmma_mma_async builtins are only available for sm_90a |
wgmma_mma_async_nonconstant_arg | E | Non-constant argument to __wgmma_mma_async call |
wgmma_mma_async_missing_args | E | The 'A' or 'B' argument to __wgmma_mma_async call is missing |
wgmma_mma_async_bad_shape | E | The shape %s is not supported for __wgmma_mma_async builtin |
wgmma_mma_async_bad_A_type | E | (invalid type for operand A) |
wgmma_mma_async_bad_B_type | E | (invalid type for operand B) |
Architecture-dependent launch configuration attributes.
| Tag | Sev | Message |
|---|
block_size_unsupported | E | __block_size__ is not supported for this GPU architecture |
block_size_must_be_positive | E | (block size values must be positive) |
cluster_dims_unsupported | E | __cluster_dims__ is not supported for this GPU architecture |
cluster_dims_must_be_positive | E | (__cluster_dims__ values must be positive) |
cluster_dims_too_large | E | cluster dimension value is too large |
conflict_between_cluster_dim_and_block_size | E | cannot specify the second tuple in __block_size__ while __cluster_dims__ is present |
max_blocks_per_cluster_unsupported | E | cannot specify max blocks per cluster for this GPU architecture |
max_blocks_per_cluster_negative | E | max blocks per cluster must not be negative |
max_blocks_per_cluster_too_large | E | max blocks per cluster is too large |
too_many_blocks_in_cluster | E | total number of blocks in cluster computed from %s exceeds __launch_bounds__ specified limit for max blocks in cluster |
shared_block_size_must_be_positive | E | the block size of a shared array must be greater than zero |
shared_block_size_too_large | E | (shared block size exceeds maximum) |
mismatched_shared_block_size | E | shared block size does not match one previously specified |
ambiguous_block_size_spec | E | (ambiguous block size specification) |
multiple_block_sizes | E | multiple block sizes not allowed |
threads_dimension_requires_definite_block_size | E | a dynamic THREADS dimension requires a definite block size |
shared_nonthreads_dim | E | (shared array dimension is not THREADS-based) |
shared_affinity_type | E | (shared affinity type mismatch) |
| Tag | Sev | Message |
|---|
inline_hint_forceinline_conflict | E | "__inline_hint__" and "__forceinline__" may not be used on the same declaration |
inline_hint_noinline_conflict | E | "__inline_hint__" and "__noinline__" may not be used on the same declaration |
| Tag | Sev | Message |
|---|
local_maxnreg | E | (__local_maxnreg__ attribute applied) |
local_maxnreg_attr_only_nonmember_func | E | (__local_maxnreg__ only on non-member functions) |
local_maxnreg_attribute_conflict | E | (__local_maxnreg__ conflicts with existing attribute) |
local_maxnreg_negative | E | (__local_maxnreg__ value is negative) |
local_maxnreg_too_large | E | (__local_maxnreg__ value exceeds maximum) |
maxnreg_attr_only_nonmember_func | E | (__maxnreg__ only on non-member functions) |
bounds_attr_only_nonmember_func | E | (launch bounds only on non-member functions) |
| Tag | Sev | Message |
|---|
cuda_displaced_new_or_delete_operator | E | (displaced new/delete in CUDA context) |
cuda_demote_unsupported_floating_point | W | (unsupported floating-point type demoted) |
illegal_ucn_in_device_identifer | E | Universal character is not allowed in device entity name (%sq) |
thread_local_for_device_vars | E | cannot use thread_local specifier for a %s variable |
global_qualifier_not_allowed | E | (execution space qualifier not allowed here) |
unsupported_nv_attribute | W | (unrecognized NVIDIA attribute) |
addr_of_nv_builtin_var | E | (address-of applied to NVIDIA builtin variable) |
shared_address_immutable | E | (__shared__ variable address is immutable) |
nonshared_blocksizeof | E | (BLOCKSIZEOF applied to non-__shared__ variable) |
nonshared_strict_relaxed | E | (strict/relaxed qualifier on non-__shared__ variable) |
extern_shared | W | (extern __shared__ variable) |
invalid_nvvm_builtin_intrinsic | E | (invalid NVVM builtin intrinsic) |
unannotated_static_not_allowed_in_device | E | (unannotated static not allowed in device code) |
missing_pushcallconfig | E | (cudaConfigureCall not found for kernel launch lowering) |
All 286 CUDA-specific diagnostic tag names extracted from the cudafe++ binary, organized alphabetically within functional groups. Every tag can be used with --diag_suppress, --diag_warning, --diag_error, or #pragma nv_diag_suppress / nv_diag_warning / nv_diag_error.
| # | Tag Name |
|---|
| 1 | unsafe_device_call |
| # | Tag Name |
|---|
| 2 | device_function_redeclared_with_global |
| 3 | device_function_redeclared_with_host |
| 4 | device_function_redeclared_with_host_device |
| 5 | device_function_redeclared_without_device |
| 6 | global_function_redeclared_with_device |
| 7 | global_function_redeclared_with_host |
| 8 | global_function_redeclared_with_host_device |
| 9 | global_function_redeclared_without_global |
| 10 | host_device_function_redeclared_with_global |
| 11 | host_function_redeclared_with_device |
| 12 | host_function_redeclared_with_global |
| 13 | host_function_redeclared_with_host_device |
| # | Tag Name |
|---|
| 14 | bounds_attr_only_on_global_func |
| 15 | bounds_maxnreg_incompatible_qualifiers |
| 16 | cuda_specifier_twice_in_group |
| 17 | global_class_decl |
| 18 | global_exception_spec |
| 19 | global_friend_definition |
| 20 | global_func_local_template_arg |
| 21 | global_function_consteval |
| 22 | global_function_constexpr |
| 23 | global_function_deduced_return_type |
| 24 | global_function_has_ellipsis |
| 25 | global_function_in_unnamed_inline_ns |
| 26 | global_function_inline |
| 27 | global_function_multiple_packs |
| 28 | global_function_pack_not_last |
| 29 | global_function_return_type |
| 30 | global_function_with_initializer_list |
| 31 | global_lambda_template_arg |
| 32 | global_new_or_delete |
| 33 | global_operator_function |
| 34 | global_param_align_too_big |
| 35 | global_private_template_arg |
| 36 | global_private_type_arg |
| 37 | global_qualifier_not_allowed |
| 38 | global_ref_param_restrict |
| 39 | global_rvalue_ref_type |
| 40 | global_unnamed_type_arg |
| 41 | global_va_list_type |
| 42 | local_type_used_in_global_function |
| 43 | maxnreg_attr_only_on_global_func |
| 44 | missing_launch_bounds |
| 45 | template_global_no_def |
| # | Tag Name |
|---|
| 46 | extended_host_device_generic_lambda |
| 47 | extended_lambda_array_capture_assignable |
| 48 | extended_lambda_array_capture_default_constructible |
| 49 | extended_lambda_array_capture_rank |
| 50 | extended_lambda_call_operator_local_type |
| 51 | extended_lambda_call_operator_private_type |
| 52 | extended_lambda_cant_take_function_address |
| 53 | extended_lambda_capture_in_constexpr_if |
| 54 | extended_lambda_capture_local_type |
| 55 | extended_lambda_capture_private_type |
| 56 | extended_lambda_constexpr |
| 57 | extended_lambda_disallowed |
| 58 | extended_lambda_discriminator |
| 59 | extended_lambda_enclosing_function_deducible |
| 60 | extended_lambda_enclosing_function_generic_lambda |
| 61 | extended_lambda_enclosing_function_hd_lambda |
| 62 | extended_lambda_enclosing_function_local |
| 63 | extended_lambda_enclosing_function_not_found |
| 64 | extended_lambda_hd_init_capture |
| 65 | extended_lambda_illegal_parent |
| 66 | extended_lambda_inaccessible_ancestor |
| 67 | extended_lambda_inaccessible_parent |
| 68 | extended_lambda_init_capture_array |
| 69 | extended_lambda_init_capture_initlist |
| 70 | extended_lambda_inside_constexpr_if |
| 71 | extended_lambda_multiple_parameter_packs |
| 72 | extended_lambda_multiple_parent |
| 73 | extended_lambda_nest_parent_template_param_unnamed |
| 74 | extended_lambda_no_parent_func |
| 75 | extended_lambda_pack_capture |
| 76 | extended_lambda_parent_class_unnamed |
| 77 | extended_lambda_parent_local_type |
| 78 | extended_lambda_parent_non_extern |
| 79 | extended_lambda_parent_private_template_arg |
| 80 | extended_lambda_parent_private_type |
| 81 | extended_lambda_parent_template_param_unnamed |
| 82 | extended_lambda_reference_capture |
| 83 | extended_lambda_too_many_captures |
| 84 | this_addr_capture_ext_lambda |
| # | Tag Name |
|---|
| 85 | addr_of_label_in_device_func |
| 86 | asm_constraint_letter_not_allowed_in_device |
| 87 | asm_constraint_must_have_single_letter |
| 88 | auto_device_fn_ref |
| 89 | cc_clobber_in_device |
| 90 | cuda_device_code_unsupported_operator |
| 91 | cuda_xasm_strict_placeholder_format |
| 92 | illegal_ucn_in_device_identifer |
| 93 | no_coroutine_on_device |
| 94 | no_strict_cuda_error |
| 95 | thread_local_in_device_code |
| 96 | undefined_device_entity |
| 97 | undefined_device_identifier |
| 98 | unrecognized_pragma_device_code |
| 99 | unsupported_type_in_device_code |
| 100 | use_of_virtual_base_on_compute_1x |
| # | Tag Name |
|---|
| 101 | device_func_tex_arg |
| 102 | device_function_has_ellipsis |
| 103 | no_host_device_initializer_list |
| 104 | no_host_device_move_forward |
| # | Tag Name |
|---|
| 105 | device_launch_no_sepcomp |
| 106 | device_side_launch_arg_with_user_provided_cctor |
| 107 | device_side_launch_arg_with_user_provided_dtor |
| 108 | missing_api_for_device_side_launch |
| # | Tag Name |
|---|
| 109 | device_var_address_taken_in_host |
| 110 | device_var_constexpr |
| 111 | device_var_read_in_host |
| 112 | device_var_structured_binding |
| 113 | device_var_written_in_host |
| 114 | device_variable_in_unnamed_inline_ns |
| 115 | host_var_address_taken_in_device |
| 116 | host_var_read_in_device |
| 117 | host_var_written_in_device |
| 118 | illegal_local_to_device_function |
| 119 | illegal_local_to_host_function |
| # | Tag Name |
|---|
| 120 | variable_template_func_local_template_arg |
| 121 | variable_template_lambda_template_arg |
| 122 | variable_template_private_template_arg |
| 123 | variable_template_private_type_arg |
| 124 | variable_template_unnamed_type_template_arg |
| # | Tag Name |
|---|
| 125 | decltype_of_managed_variable |
| 126 | managed_cant_be_shared_constant |
| 127 | managed_const_type_not_allowed |
| 128 | managed_reference_type_not_allowed |
| 129 | unsupported_arch_for_managed_capability |
| 130 | unsupported_configuration_for_managed_capability |
| # | Tag Name |
|---|
| 131 | grid_constant_incompat_instantiation_directive |
| 132 | grid_constant_incompat_redecl |
| 133 | grid_constant_incompat_specialization |
| 134 | grid_constant_incompat_templ_redecl |
| 135 | grid_constant_non_kernel |
| 136 | grid_constant_not_const |
| 137 | grid_constant_reference_type |
| 138 | grid_constant_unsupported_arch |
| # | Tag Name |
|---|
| 139 | floating_type_for_logical_atomic_operation |
| 140 | invalid_data_size_for_nv_atomic_generic_function |
| 141 | invalid_nv_atomic_cas_size |
| 142 | invalid_nv_atomic_exch_size |
| 143 | invalid_nv_atomic_memory_order_value |
| 144 | invalid_nv_atomic_operation_add_sub_size |
| 145 | invalid_nv_atomic_operation_max_min_float |
| 146 | invalid_nv_atomic_operation_size |
| 147 | invalid_nv_atomic_thread_scope_value |
| 148 | non_integral_type_for_non_generic_nv_atomic_function |
| 149 | nv_atomic_add_sub_f64_not_supported |
| 150 | nv_atomic_cas_b16_not_supported |
| 151 | nv_atomic_exch_cas_b128_not_supported |
| 152 | nv_atomic_function_address_taken |
| 153 | nv_atomic_function_no_args |
| 154 | nv_atomic_functions_not_supported_below_sm60 |
| 155 | nv_atomic_load_order_error |
| 156 | nv_atomic_load_store_b128_version_too_low |
| 157 | nv_atomic_load_store_scope_cluster_change_to_device |
| 158 | nv_atomic_operation_not_in_device_function |
| 159 | nv_atomic_operation_order_not_constant_int |
| 160 | nv_atomic_operation_scope_not_constant_int |
| 161 | nv_atomic_operations_memory_order_fallback_to_membar |
| 162 | nv_atomic_operations_scope_cluster_change_to_device |
| 163 | nv_atomic_operations_scope_fallback_to_membar |
| 164 | nv_atomic_store_order_error |
| # | Tag Name |
|---|
| 165 | host_closure_class_in_jit |
| 166 | no_host_in_jit |
| 167 | unannotated_function_in_jit |
| 168 | unannotated_static_data_member_in_jit |
| 169 | unannotated_variable_in_jit |
| # | Tag Name |
|---|
| 170 | extern_kernel_template |
| 171 | template_global_no_def |
| # | Tag Name |
|---|
| 172 | nv_abi_pragma_bad_format |
| 173 | nv_abi_pragma_device_function |
| 174 | nv_abi_pragma_device_function_context |
| 175 | nv_abi_pragma_duplicate_arg |
| 176 | nv_abi_pragma_invalid_option |
| 177 | nv_abi_pragma_missing_arg |
| 178 | nv_abi_pragma_next_construct |
| 179 | nv_abi_pragma_not_constant |
| 180 | nv_abi_pragma_not_positive_value |
| 181 | nv_abi_pragma_overflow_value |
| # | Tag Name |
|---|
| 182 | register_params_ellipsis_function |
| 183 | register_params_not_enabled |
| 184 | register_params_unsupported_arch |
| 185 | register_params_unsupported_function |
| # | Tag Name |
|---|
| 186 | name_expr_extra_tokens |
| 187 | name_expr_internal_error |
| 188 | name_expr_non_device_variable |
| 189 | name_expr_non_global_routine |
| 190 | name_expr_not_routine_or_variable |
| 191 | name_expr_parsing |
| # | Tag Name |
|---|
| 192 | addr_of_text_surf_expr_in_device_func |
| 193 | addr_of_text_surf_var_in_device_func |
| 194 | indir_into_text_surf_expr_in_device_func |
| 195 | indir_into_text_surf_var_in_device_func |
| 196 | reference_to_text_surf_type_in_device_func |
| 197 | reference_to_text_surf_var_in_device_func |
| 198 | texture_surface_variable_in_unnamed_inline_ns |
| # | Tag Name |
|---|
| 199 | wgmma_mma_async_bad_A_type |
| 200 | wgmma_mma_async_bad_B_type |
| 201 | wgmma_mma_async_bad_shape |
| 202 | wgmma_mma_async_missing_args |
| 203 | wgmma_mma_async_nonconstant_arg |
| 204 | wgmma_mma_async_not_enabled |
| # | Tag Name |
|---|
| 205 | ambiguous_block_size_spec |
| 206 | block_size_must_be_positive |
| 207 | block_size_unsupported |
| 208 | cluster_dims_must_be_positive |
| 209 | cluster_dims_too_large |
| 210 | cluster_dims_unsupported |
| 211 | conflict_between_cluster_dim_and_block_size |
| 212 | max_blocks_per_cluster_negative |
| 213 | max_blocks_per_cluster_too_large |
| 214 | max_blocks_per_cluster_unsupported |
| 215 | mismatched_shared_block_size |
| 216 | multiple_block_sizes |
| 217 | shared_affinity_type |
| 218 | shared_block_size_must_be_positive |
| 219 | shared_block_size_too_large |
| 220 | shared_nonthreads_dim |
| 221 | threads_dimension_requires_definite_block_size |
| 222 | too_many_blocks_in_cluster |
| # | Tag Name |
|---|
| 223 | inline_hint_forceinline_conflict |
| 224 | inline_hint_noinline_conflict |
| # | Tag Name |
|---|
| 225 | bounds_attr_only_nonmember_func |
| 226 | local_maxnreg |
| 227 | local_maxnreg_attr_only_nonmember_func |
| 228 | local_maxnreg_attribute_conflict |
| 229 | local_maxnreg_negative |
| 230 | local_maxnreg_too_large |
| 231 | maxnreg_attr_only_nonmember_func |
| # | Tag Name |
|---|
| 232 | lambda_operator_annotated |
| # | Tag Name |
|---|
| 233 | addr_of_nv_builtin_var |
| 234 | bad_shared_storage_class |
| 235 | cuda_demote_unsupported_floating_point |
| 236 | cuda_displaced_new_or_delete_operator |
| 237 | extern_shared |
| 238 | invalid_nvvm_builtin_intrinsic |
| 239 | missing_pushcallconfig |
| 240 | nonshared_blocksizeof |
| 241 | nonshared_strict_relaxed |
| 242 | shared_address_immutable |
| 243 | shared_inside_struct |
| 244 | shared_parameter |
| 245 | shared_specifier_in_range_for |
| 246 | thread_local_for_device_vars |
| 247 | unannotated_static_not_allowed_in_device |
| 248 | unsupported_nv_attribute |
| # | Tag Name |
|---|
| 249 | nv_diag_default |
| 250 | nv_diag_error |
| 251 | nv_diag_once |
| 252 | nv_diag_remark |
| 253 | nv_diag_suppress |
| 254 | nv_diag_warning |
The following standard EDG error codes (0--3456) are repurposed or frequently triggered by CUDA-specific validation. These display with their original number (not the 20000-D series).
| Internal # | Display # | Context |
|---|
| 21 | 21 | CUDA auto type with template deduction |
| 147 | 147 | redeclaration mismatch |
| 149 | 149 | illegal CUDA storage class at namespace scope |
| 246 | 246 | static member of non-class type |
| 298 | 298 | typedef/using with template name |
| 325 | 325 | thread_local in CUDA |
| 337 | 337 | calling convention mismatch |
| 453 | 453 | in template instantiation context |
| 551 | 551 | not a member function |
| 795 | 795 | definition in class scope with external linkage (CUDA) |
| 799 | 799 | definition in class scope (C++20 CUDA) |
| 891 | 891 | anonymous type in variable declaration |
| 892 | 892 | auto with __constant__ variable |
| 893 | 893 | auto with CUDA variable |
| 948 | 948 | calling convention mismatch on redeclaration |
| 992 | 992 | fatal error (suppress-all sentinel) |
| 1034 | 1034 | explicit instantiation with conflicting attributes |
| 1063 | 1063 | in include file context |
| 1118 | 1118 | CUDA attribute on namespace-scope variable |
| 1150 | 1150 | context lines truncated |
| 1158 | 1158 | auto return type with __global__ |
| 1306 | 1306 | CUDA memory space mismatch on redeclaration |
| 1418 | 1418 | incomplete type in definition |
| 1430 | 1430 | function attribute mismatch in template |
| 1560 | 1560 | CUDA constexpr class with non-trivial destructor |
| 1580 | 1580 | redeclaration with different template parameters |
| 1655 | 1655 | tentative definition of constexpr |
| 2384 | 2384 | constexpr mismatch on redeclaration (CUDA) |
| 2442 | 2442 | extern variable at block scope with CUDA attribute |
| 2443 | 2443 | extern variable at block scope with CUDA attribute (variant) |
| 2502 | 2502 | no_unique_address mismatch |
| 2503 | 2503 | no_unique_address mismatch (variant) |
| 2656 | 2656 | internal error (assertion failure) |
| 2885 | 2885 | CUDA attribute on deduction guide |
| 2937 | 2937 | structured binding with CUDA attribute |
| 3033 | 3033 | incompatible constexpr CUDA target |
| 3116 | 3116 | restrict qualifier on definition |
| 3414 | 3414 | auto with volatile/atomic qualifier |
| 3510 | 3510 | __shared__ variable with VLA |
| 3566 | 3566 | __constant__ with constexpr with auto |
| 3567 | 3567 | CUDA variable with VLA type |
| 3568 | 3568 | __constant__ with constexpr |
| 3578 | 3578 | CUDA attribute in discarded constexpr-if branch |
| 3579 | 3579 | CUDA attribute at namespace scope with structured binding |
| 3580 | 3580 | CUDA attribute on variable-length array |
| 3648 | 3648 | CUDA __constant__ with external linkage |
| 3698 | 3698 | parameter type mismatch |
| 3709 | 3709 | warnings treated as errors |
CUDA error messages use the same fill-in system as EDG base errors, expanded by process_fill_in (sub_4EDCD0).
| Specifier | Kind | Meaning | Example |
|---|
%sq | 3 | Quoted entity name | Function name in cross-space call |
%sq1, %sq2 | 3 | Indexed quoted names | Caller and callee |
%no1 | 4 | Entity name (omit kind prefix) | Function in redeclaration |
%n1, %n2 | 4 | Entity names | Override base/derived pair |
%nd | 4 | Entity with declaration location | Template parameter |
%s, %s1, %s2 | 3 | String fill-in | Execution space keyword |
%t | 6 | Type fill-in | Type in template arg errors |
%p | 2 | Source position | Previous declaration location |
Quick reference for minimum architecture required by various CUDA features.
| Feature | Minimum Architecture |
|---|
| Virtual bases | compute_20 |
__device__/__host__ __device__ with ellipsis | compute_30 |
__managed__ variables | compute_30 |
| alloca() | compute_52 |
__nv_atomic_* functions | sm_60 |
| Atomic scope argument | sm_60 |
| Atomic add/sub for f64 | sm_60 |
__grid_constant__ | compute_70 |
| Atomic memory order argument | sm_70 |
| 16-bit atomic CAS | sm_70 |
| 128-bit atomic load/store | sm_70 |
__nv_register_params__ | compute_80 |
| Cluster scope for atomics | sm_90 |
| 128-bit atomic exchange/CAS | sm_90 |
__wgmma_mma_async builtins | sm_90a |
Beyond the 286 CUDA-extension tags above, the cudafe++ binary embeds approximately 573 additional snake_case diagnostic tags belonging to EDG 6.6's standard C++ frontend. They share the same suppression machinery -- --diag_suppress=<tag>, #pragma nv_diag_suppress <tag>, numeric #pragma nv_diag_suppress <N> for codes 0--3456 -- but live in the original error-code range and are emitted without the 20000-D renumbering.
The table below summarizes the top category prefixes harvested from the binary's string table. Counts are tag totals per prefix; the role column describes what the prefix governs.
| Category prefix | Tags | Role |
|---|
invalid_ | 150 | Structural validity (invalid declarator, invalid attribute location, invalid asm qualifier, ...) |
constexpr_ | 112 | Constexpr interpreter errors (access past object, allocation mismatch, call not interpretable, undefined behaviour, ...) |
type_ | 56 | Type-system inconsistencies (conflicts-with-tag, after-integral-promotion, constraint-failed, ...) |
template_ | 48 | Template engine diagnostics (instance linkage conflict, arg involves error entity, depth mismatch, ...) |
function_ | 34 | Function declaration/definition issues (body processing, defaulted in friend decl, does-not-match-arguments, ...) |
ambiguous_ | 28 | Overload and lookup ambiguity (constructor, conversion function, partial spec, virtual override, ...) |
class_ | 26 | Class-layout and inheritance constraints |
module_ | 24 | C++20 modules support (import conflict, file-not-found, declaration-position, primary-name, ...) |
attribute_ | 24 | GCC/MSVC/standard attribute application rules (cleanup-storage, condition-satisfied, ignored-on-incomplete-class, ...) |
static_ | 21 | Static specifier and linkage |
virtual_ | 19 | Virtual function rules |
symbol_ / return_ | 19 / 19 | Symbol resolution; return-type mismatches |
member_ / implicit_ | 18 / 18 | Member access; implicit conversion / construction |
struct_ / pragma_ / lambda_ / enum_ | 15 each | Struct layout; #pragma parsing; lambda restrictions; enum constraints |
⚡ QUIRK — diagnostic tags survive renumbering
A CUDA error displayed as #20042-D and an EDG error displayed as #147 are both suppressible by tag name through the same global lookup. Internally cudafe++ resolves the tag string against all 859 entries before the per-error-code severity arrays (byte_1067920 / byte_1067921 / byte_1067922) are consulted, so users never need to know whether a tag belongs to the CUDA extension range or the EDG base range. The numeric #pragma nv_diag_suppress 20042 form silently subtracts 16543 before indexing the severity arrays.
The EDG-base namespace is intentionally not enumerated tag-by-tag in this wiki: the upstream catalog is maintained by EDG and tags are largely stable across EDG 6.x point releases. For a deep dive into specific subsystems consult:
Confidence: MED -- prefix bucketing is mechanical, but a handful of tags straddle categories (e.g. module_id_for_source_corresp could be either modules or RDC).