CUDASafety-1.6ΒΆ

Constexpr host variables shall only be used in valid contexts

Required inputs: IR

CUDA SAFETY 1.6 Constexpr host variables shall only be used in valid contexts.

Let 'V' denote a namespace scope variable or a class static member variable that has been marked constexpr and that does not have execution space annotations (e.g., __device__, __constant__, __shared__). V is considered to be a host code variable. If V is of scalar type other than long double and the type is not volatile-qualified, the value of V can be directly used in device code. In addition, if V is of a non-scalar type then scalar elements of V can be used inside a constexpr __device__ or __host__ __device__ function, if the call to the function is a constant expression. Device source code cannot contain a reference to V or take the address of V.

Scope: Device.
Audience: CUDA C++.
Category: Mandatory.
Hardware Applicability: All Compute Capabilities.
Rationale

Accessing host variables may exhibit unexpected behavior at runtime.

Example 1 (Bad/Good)
constexpr int xxx = 10;
constexpr int yyy = xxx + 4;

struct S1_t { static constexpr int qqq = 100; };
constexpr int host_arr[] = { 1, 2, 3};
constexpr __device__ int get(int idx) { return host_arr[idx]; }

__device__ int foo(int idx) {
  int v1 = xxx + yyy + S1_t::qqq; // OK
  const int &v2 = xxx; // error: reference to host constexpr
// variable
  const int *v3 = &xxx; // error: address of host constexpr
// variable
  const int &v4 = S1_t::qqq; // error: reference to host constexpr
// variable
  const int *v5 = &S1_t::qqq; // error: address of host constexpr
// variable
  v1 += get(2); // OK: 'get(2)' is a constant
// expression.
  v1 += get(idx); // error: 'get(idx)' is not a constant
// expression
  v1 += host_arr[2]; // error: 'host_arr' does not have
// scalar type.
  return v1;
}
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

device_addr_constexpr_host_var

Device source code cannot contain a reference to a constexpr host variable or take the address of it

None

False

invalid_constexpr_host_var_access

Invalid use of constexpr host variable in device code

None

False

Options