CUDA-1.10

Never use CUDA_ARCH to guard CUDA declarations

Required inputs: IR

CUDA 1.10 [preprocessing.guarded-declarations] Never use CUDA_ARCH to guard CUDA declarations

The declaration of a device variable or global function shall not be in a region of code guarded by a macro dependent on CUDA_ARCH.

Scope: Host, Device.

Audience: CUDA ++.

Category: Required.

Hardware Applicability: All Compute Capabilities.

Rationale

The __CUDA_ARCH__ macro is defined when generating code for the GPU. Any declaration that is guarded by the macro will only be visible on the device. This can lead to missing declarations or objects with different types on the device and host. This can lead to confusion when objects do not behave as expected when data is transferred between the host and the device. This can also cause compilation problems for compilers that perform unified compilation, compiling both the host and device code in a single pass.

Example 1 (Bad)

# ifdef __CUDA_ARCH__

int x; // non-compliant: the declaration is only visible in the device program.
# endif
# ifdef __CUDA_ARCH__
int
# else
float
# endif
  y; // non-compliant: y has type int in device code and float in host code.

Excerpt from NVIDIA CUDA C++ Guidelines for robust and safety-critical programming, Version 3.0.1, Copyright (C) 2018-2023 NVIDIA Corporation.

Possible Messages

Key	Text	Severity	Disabled
decl_depends_on_cuda_arch	This declaration depends on __CUDA_ARCH__	None	False

Options

• This rule shares the following common options: exclude_in_macros, exclude_messages_in_system_headers, excludes, extend_exclude_to_macro_invocations, includes, justification_checker, languages, post_processing, provider, report_at, severity

• The following places define options that affect this rule: Stylechecks, Analysis-GlobalOptions

report_all_nonmember_functions

report_all_nonmember_functions : bool = True

Whether only kernels should be reported, or all non-member functions

 

report_constexpr

report_constexpr : bool = False

Whether constexpr variables should be reported.

 

report_host_vars

report_host_vars : bool = True

Whether only __device__ variables should be reported, or also other non-local variables when their declaration depends on __CUDA_ARCH__