__device__ Variable Type Qualifier Strange behavior on Windows

Hi,

I would like to report the following strange behavior on CUDA. The following is a simple code that runs on Linux (openSUSE 11.1), but not run on Windows Vista x64:

[codebox]“main.cu”

#include

#include <cuda_runtime.h>

const float A = 10.0;

host device float sqrtA()

{

  return A * A;

}

int main()

{

  std::cout << sqrtA() << std::endl;

return 0;

}

[/codebox]

When I try to run on Visual Studio, I get the error: identifier “A” is undefined

Then I tried to specify device on variable A:

[codebox]device const float A = 10.0;[/codebox]

It do not give me a compilation error anymore (because nvcc was able to compile function sqrtA() on the device), BUT the value printed is ZERO, which means that the host function compiled by nvcc is not accessing the correct value of variable A.

Please, help me. I have already posted a similar thread and nobody answered.

Best

You tell that this code runs on Linux without any error? To set a constant variable you must use the function cudaMemcpyToSymbol or the device could not be able to set a value for the identifier. When you use the device modifier, the variable is initialised with a default value (ZERO in this case). The error you’ve got that identifier A is undefined was correct. The device wasn’t be able to see the identifier A because it not have the modifier device, so the nvcc set this identifier as a host identifier. When the compiler try to compile the device version of the function it will not have access to identifier declared in host code.

Hope this will help you.

I’m having the same problem. I’m also having this problem on Windows (Win7 in my case).

The previous response is not relevant because it is clearly the host instance of the function that is being called, not the device version. Hence cudaMemcpyToSymbol is not relevant.

It seems that device const does not create the constant for use in the host code (except perhaps in Linux). Unfortunately you can’t say host device const.

I found a workaround: You just declare your constants inside the function body. Unfortunately, this means that you may have to repeatedly define the same constants wherever they are used. Alternatively, you can use the archaic define for your constants.

I’d call this a bug in the Windows version of the cuda compiler, especially since the Linux version works as expected. In any case, it should not be possible to find yourself with a constant arbitrarily changed to zero without a compiler error.

"main.cu"

#include <iostream>

#include <cuda_runtime.h>

__host__ __device__ float sqrtA()

{

      const float A = 10.0;

      return A * A;

}

int main()

{

      std::cout << sqrtA() << std::endl;

return 0;

}