problem in reading constant memory variables inside kernels problem in using several __constant__ va

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello to All,

I have 3 kernels (in 3 separate files).

These kernels have to use several constant variables.

These constant variables are declared, defined and allocated on the device in another file.

The problem is that some kernel does not get the right value of constant variables.

I also tried to use cudaGetSymbolAddress in the host function calling the kernel, but I get

“cannot take the address of constant data”

How have I to declare the variables in the files containing kernels?

Here it is how I allocate constant device memory:

[codebox]

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“xdim”, &(ps->xdim), sizeof(int), 0, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“ydim”, &(ps->ydim), sizeof(int), 0, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“kStretch”, &(ps->kStretch), sizeof(REAL), 0, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“kShear”, &(ps->kShear), sizeof(REAL), 0, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“kBend”, &(ps->kBend),sizeof(REAL), 0, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“pitch_position_d”, &ppos, sizeof(size_t), 0, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“pitch_velocity_d”, &pvel, sizeof(size_t), 0, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“pitch_force_d”, &pforce, sizeof(size_t), 0, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(“pitch_mass_d”, &pmass, sizeof(size_t), 0, cudaMemcpyHostToDevice));

[/codebox]

Here it is how I declare variables in each one of the files containing kernels

[codebox]

constant int xdim;

constant int ydim;

constant size_t pitch_position_d;

constant size_t pitch_velocity_d;

constant size_t pitch_force_d;

constant size_t pitch_mass_d;

constant size_t pitch_previousPosition_d;

[/codebox]

Why I can’t take the value of these variables?

I also tried using the device qualifier, but without success.

What’s wrong?

Thanks in advance

Francesco

2

I have just spent a few days trying to work out why my code wasn’t working. Last night I found out that this was the problem - I had all my device variables declared in a .h file, which was then included by the .cu file that interfaced with the rest of the host code (and therefore contained all the cudaMemcpyToSymbol, etc.), and the .cu file that actually contained the device code. Eventually I just merged all of three cuda files (the .h and the two .cu’s) into one file and now it works. It’s very ugly, and I don’t know why nvcc can’t handle this. I also don’t understand why it is not possible to have device functions in a separate file from the device function that calls them - it means I have to have very large and unwieldy files.

3

Hi

thank you for your reply. Yesterday I discovered that only the last .cu file linked (i.e. the last .cu in the makefile) can use constant variables…

The others .cu files don’t read the exact value of the variable. I really don’t know why

Thank you again.

4

This limitation is documented. Check the programming guide: Section B.2.5

So constant variables should only be usable from within the same compilation unit. This limitation has to do with the fact that there is no linker to link the device code and variables.

5

Huge documentation :)

By the way, thank you very much. Now it is really clear!

Francesco