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

Question

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

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

我认为这是来自E.2.13 的摘录。CUDA 文档的Const 限定变量部分解释说：

让'V' 表示命名空间范围变量或具有const 限定类型且没有执行空间注释（例如，__device__、__constant__、__shared__）的类静态成员变量。V 被认为是主机代码变量。

V 的值可以直接用于设备代码，如果：

• V 在使用点之前已经用常量表达式初始化，并且
• 它具有以下类型之一：
  • 内置浮点类型，除非使用 Microsoft 编译器作为主机编译器，
  • 内置整数类型。

设备源代码不能包含对 V 的引用或采用 V 的地址。

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