CODEWITHSUNDEEP
cudacompiler-warningsnvcccompiler-bugtemplate-variables
Herzog Igzorn
Herzog Igzorn

Reputation: 33

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

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

Upvotes: 2

Views: 654

Answers (2)

einpoklum
einpoklum

Reputation: 131546

@MichaelKenzel is correct.

This is almost certainly an nvcc bug - which I have now filed (you might need an account to access that.

Also note I've been able to reproduce the issue with less code:

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

and have a look at the result on GodBolt as well.

Upvotes: 1

Michael Kenzel
Michael Kenzel

Reputation: 15941

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…

Upvotes: 3

Related Questions