CUDA中CPU与GPU共享constexpr常量的行为可靠性咨询

你好，针对你提到的在CPU和GPU之间共享constexpr编译期常量、遇到编译器警告但运行正常的问题，我来详细解析下相关的行为和可靠性：

首先，先明确你当前的核心场景：你定义了__device__ constexpr fe_ftype vars[2] = {100.0, 300.0};，同时在host端的main函数和device端的foo核函数中直接访问这个数组，编译时出现警告但运行时两端都能正确输出值。

为什么当前代码能正常运行？ #

这完全是constexpr特性带来的结果：C++17中，constexpr变量要求其值在编译期就能被完全确定。对于标记了__device__的constexpr变量，CUDA编译器（nvcc）在处理host端代码时，并不会将其当作设备内存中的变量来访问，而是直接把编译期计算出的常量值嵌入到host代码中——相当于在host端，这个变量的行为和普通的constexpr变量没有区别，自然不会出现读取未初始化内存的问题。

关于编译器警告#20091-D #

这个警告是nvcc针对普通__device__变量设计的：普通的__device__变量只存在于设备内存空间，host端直接访问会导致读取无效的内存地址，触发未定义行为。但你的变量是constexpr，编译器能识别到它是编译期常量，不会去访问设备内存，所以这个警告属于“误报”，对constexpr的__device__变量来说，实际不会引发错误。

用__constant__替代__device__的情况 #

你提到换成__constant__ constexpr也能正常工作，这是因为__constant__变量会被放在设备的常量内存区域，但同样因为constexpr的特性，host端依然会直接使用编译期确定的值，不会去访问设备的常量内存。不过要注意，__constant__变量在设备端的访问特性（比如缓存行为）和__device__变量不同，但这对host端的行为没有影响。

能不能依赖这种行为？ #

在C++17环境下，这种用法是可靠的，理由如下：

• nvcc严格遵守C++标准中constexpr的语义，只要变量确实是编译期可计算的，编译器就会在host端直接嵌入常量值，不会尝试访问设备内存。
• 从CUDA的文档和编译器实现来看，这种跨host/device共享constexpr常量的用法是被支持的，属于合法的C++ CUDA扩展。

不过有几个需要注意的点：

• 确保你的constexpr变量的初始化逻辑完全是编译期可计算的，不能包含任何运行时才能确定的操作（比如调用非constexpr函数、读取运行时变量等），否则constexpr会失效，这时警告就会变成实际的未定义行为。
• 不同版本的nvcc可能存在细微的优化差异，建议在你实际使用的编译环境中做充分测试，确保行为一致。
• 如果觉得警告烦人，可以通过编译器选项-Wno-20091来抑制这个特定警告，但前提是你确认所有被标记__device__/__constant__的constexpr变量都是真正的编译期常量。

内容来源于stack exchange