nvcc warns about a device variable being a host variable - why?

Question

I've been reading in the CUDA Programming Guide about template functions and is something like this working?

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

I get the correct result but a warning:

"warning: a host variable "d_test [with T=T]" cannot be directly read in a device function" ?

Has the struct in the testfunction to be instantiated with *d_test<float>->val ?

KR, Iggi

Answer 1

Unfortunately, the CUDA compiler seems to generally have some issues with variable templates. If you look at the assembly, you'll see that everything works just fine. The compiler clearly does instantiate the variable template and allocates a corresponding device object.

.global .align 8 .u64 _Z6d_testIfE;

The generated code uses this object just like it's supposed to

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

I'd consider this warning a compiler bug. Note that I cannot reproduce the issue with CUDA 10 here, so this issue has most likely been fixed by now. Consider updating your compiler…