nvcc 编译器将静态 constexpr 识别为设备代码中未定义

Question

这个问题是的后续问题。
它是关于 nvcc 编译器识别 static constexpr class 变量在设备代码中未定义，如果变量是 odr-used。但是，我找不到为什么它不起作用的原因。
错误信息是：

error: identifier "Tester<int> ::ONE" is undefined in device code

编译为

nvcc -std=c++11 -ccbin=/usr/bin/g++-4.9 -arch=sm_30 main.cu

nvcc编译器版本是release 8.0, V8.0.26。

给出了一个最小的例子（上一个问题中 MWE 的简化版本，专注于这个特定问题）

#include <iostream>
#include <cstdlib>

#ifdef __CUDACC__
    #define HD __host__ __device__
#else
    #define HD
#endif


HD void doSomething(const int& var ) {};

template<typename T> class Tester
{
public:
    static constexpr int ONE = 1;

    HD void test()
    {
        doSomething( ONE );
    }
};
template<typename T> constexpr int Tester<T>::ONE;


int main()
{
    using t = int;

    Tester<t> tester;
    tester.test();

    return EXIT_SUCCESS;
}

问题是不是关于修复这个特定代码（这将通过按值传递 var 而不是 const 引用来完成 - 至少编译器不再抱怨，尽管它是一种 ODR 用途，不是吗？）。
问题是，这是否是 nvcc 编译器的错误，或者是否有充分的理由，为什么这不起作用（我在 NVIDIA 页面上没有找到任何提示...）。

Answer 1

我认为 CUDA 文档 E.2.13. Const-qualified variables 部分的这段摘录解释了：

Let 'V' denote a namespace scope variable or a class static member variable that has const qualified type and does not have execution space annotations (e.g., __device__, __constant__, __shared__). V is considered to be a host code variable.

The value of V may be directly used in device code, if:

V has been initialized with a constant expression before the point of use, and

it has one of the following types:

builtin floating point type except when the Microsoft compiler is used as the host compiler,

builtin integral type.

Device source code cannot contain a reference to V or take the address of V.

我认为您的代码违反了最后一句话 - 您的代码包含引用。

nvcc 编译器将静态 constexpr 识别为设备代码中未定义

nvcc compiler recognizes static constexpr as undefined in device code

c++

cuda

nvcc

constexpr