CODEWITHSUNDEEP
cuda
Bwawhwmopoloin
Bwawhwmopoloin

Reputation: 91

Scope of __device__ qualifier

On http://docs.nvidia.com/cuda/cuda-c-programming-guide/#device-variable-qualifier it says that a __device_ qualifier variable has the "lifetime of an application". Does this mean the kernel? If there are several kernels, how can CUDA know which variable belongs to which kernel?

If i declare a __device_ variable like so: 

void someHOSTfunction() {
   __device__ int var;
   // Launch kernel etc...
}

Is "var" still global in the sense that it is still accessible from a kernel launched from a different function, even though it is "local" on the stack of someHOSTfunction() and gets scoped (?) when someHOSTfunction() returns? Does it make any difference to write it like this: 

__device__ int var;
void someHOSTfunction() {
   // Launch kernel etc...
}

Now var is a global variable. But that means it is also accessible from other translation units. This probably won't work to prevent that: 

static __device__ int var;
void someHOSTfunction() {
   // Launch kernel etc...
}

What would be the appropriate way of doing it?

Upvotes: 4

Views: 2434

Answers (1)

talonmies
talonmies

Reputation: 72349

This:

void someHOSTfunction() {
   __device__ int var;
   // Launch kernel etc...
}

is illegal in CUDA. A __device__ variable declaration is not allowed inside a function body, and the compiler will emit an error if you try to do so. You must declare them at translation unit scope. The restriction applies to any function, be it __host__ or __device__.

If you need different statically declared __device__ variables for different kernels, then use a different variable name for each. Or use a runtime allocated variable and pass it as an argument to your kernel, or use template parameter variables, or something else. But what you describe is impossible in CUDA as it exists today.

Upvotes: 3

Related Questions