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

Question

这个问题是this one的后续问题.
它是关于 nvcc 编译器将 static constexpr 类变量识别为设备代码中未定义的，如果变量是 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。

#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

我认为这段摘自 E.2.13. Const-qualified variables CUDA 文档部分解释说:

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.

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