我一直在阅读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
实例化?
KR,Iggi
不幸的是,CUDA编译器似乎通常在变量模板方面存在一些问题。如果你look at the assembly,你会发现一切正常。编译器清楚地实例化变量模板并分配相应的设备对象。
.global .align 8 .u64 _Z6d_testIfE;
生成的代码使用此对象就像它应该的那样
ld.global.u64 %rd3, [_Z6d_testIfE];
我认为这个警告是一个编译错误。请注意,我无法在此处重现CUDA 10的问题,所以此问题很可能已经解决。考虑更新你的编译器......
@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>(); }