nvcc 警告设备变量是主机变量 - 为什么？答案

【问题标题】：nvcc warns about a device variable being a host variable - why?nvcc 警告设备变量是主机变量 - 为什么？
【发布时间】：2019-08-19 14:14:49
【问题描述】：

我一直在阅读 CUDA 编程指南中关于模板函数的内容，这样的东西是否有效？

#include <cstdio>

/* host struct */
template <typename T>
struct Test {
    T  *val;
    int size;
};

/* struct device */
template <typename T>
__device__ Test<T> *d_test;

/* test function */
template <typename T>
T __device__ testfunc() {
    return *d_test<T>->val;
}

/* test kernel */
__global__ void kernel() {
    printf("funcout = %g \n", testfunc<float>());
}

我得到了正确的结果，但有一个警告：

“警告：不能在设备函数中直接读取主机变量“d_test [with T=T]”？

testfunction 中的结构是否用*d_test<float>->val 实例化？

韩国，伊吉

【问题讨论】：

标签： cuda compiler-warnings nvcc compiler-bug template-variables

【解决方案1】：

不幸的是，CUDA 编译器似乎通常在变量模板方面存在一些问题。如果您look at the assembly，您会发现一切正常。编译器显然会实例化变量模板并分配相应的设备对象。

.global .align 8 .u64 _Z6d_testIfE;

生成的代码使用这个对象，就像它应该做的那样

ld.global.u64   %rd3, [_Z6d_testIfE];

我认为这个警告是编译器错误。请注意，我无法在此处重现 CUDA 10 的问题，因此这个问题很可能现在已经修复。考虑更新你的编译器……

【讨论】：

我能够使用 -std=c++14 在 Fedora 29 上重现 CUDA 10.1 上的问题，所以我认为它可能仍然存在。感谢那些提交错误的人。
好的，谢谢！我正在使用带有 -std=c++14 的 devtoolset-8 中的 CentOS 7.6、CUDA 10.1 和 gcc (g++) 8.2.1 编译器。

【解决方案2】：

@MichaelKenzel 是正确的。

这几乎可以肯定是一个 nvcc 错误 - 我现在有 filed（您可能需要一个帐户才能访问它。

另外请注意，我已经能够用更少的代码重现该问题：

template <typename T>
struct foo { int  val; };

template <typename T>
__device__ foo<T> *x;

template <typename T>
int __device__ f() { return x<T>->val; }

__global__ void kernel() { int y = f<float>(); }

并查看result on GodBolt。

【讨论】：