NVCC 警告设备变量是主机变量 - 为什么

Question

我一直在阅读有关模板函数的 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 [与 T=T]"无法在设备函数中直接读取" ？

测试函数中的结构是否要用*d_test<float>->val实例化？

氪，伊吉

Answer 1

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

.global .align 8 .u64 _Z6d_testIfE;

生成的代码使用此对象，就像它应该

的那样

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

我认为此警告是编译器错误。请注意，我无法在此处重现 CUDA 10 的问题，因此此问题现在很可能已解决。考虑更新编译器...

Answer 2

>@MichaelKenzel是正确的。

这几乎可以肯定是一个 nvcc 错误 - 我现在已经提交了(您可能需要一个帐户才能访问它。

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

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>(); }

并在GodBolt上查看结果。