__constant__ qualifier questions, please help

Accelerated Computing CUDA CUDA Programming and Performance
1

I am having a hard time on using constant memory space.

(1) According to Programming Guide (2.3) B.2.5, constant variables are only allowed at file scope, and they can not be defined by the extern keyword. So the question is: is it possible to make some constant data visible across multiple files?

(2) According to Programming Guide (2.3) B.2.5, the function cudaGetSymbolAddress can be used to take the address of a constant variable. However, the following code complains “cannot take the address of constant data”. What is wrong with it? Or simply we could not take the address of a constant variable, not as stated in the programming guide.

#include
using namespace std;
constant int DATA;
int main(void)
{
int d = 1;
cudaMemcpyToSymbol(“DATA”, &d, sizeof(int));
int data_ptr;
cudaError_t err = cudaGetSymbolAddress((void*)&data_ptr, “DATA”);
cout << cudaGetErrorString(err) << endl;
return 0;
}

Thanks in advance!

WP

2

Hi!

About #1 I’m not entirely sure.

About #2 I also had som minor issues when trying to implement it the first time. I finally didn’t use the cudaGetSymbolAddress stuff… You can simply do:

__constant__ float d_data[length]; // file scope

.....

cudaMemcpyToSymbol(d_data, h_data, data_size);

Hope this helps!

3

No, you cannot. The same is true of device and texture declarations, by the way. They are implicitly declared static for that compilation unit. The root cause of this behavior is that there is no linker for CUDA kernel code.

4

IIRC you can just stick them in a header file included by the relavent .cu files. No exern keyword required. Same with texture declarations.

5

This will not generate any compile errors, no. But if you compile each .cu files in separate compilation units, each gets its own private version of the constant variable or texture reference. If you, say, bind the texture ref in one compilation unit and then try to use it in a kernel in another, you will get a texture not bound error.

6

And this is perhaps one of my biggest remaining sources of irritation with CUDA.

7

Thanks for your answer! This really makes my life a bit hard to maintain a relatively large program.

Do you have any idea about the second question? I feel that the programming guide on this topic is misleading or simply wrong.

8

I haven’t tried it with a constant, but the reference manual and doxygen pages clearly say cudaGetSymbolAddress() works for symbols in global memory. No mention of constant memory.

9

Valkering has suggested a work around here:
[url=“The Official NVIDIA Forums | NVIDIA”]http://forums.nvidia.com/index.php?showtopic=109827[/url]
I haven’t tried that yet.