constant memory declaration How to?

Hi,

I have a question. I want to use constant memory but I cant figure out how to use it.

The situation is like this:

I have two .cu files, one named test.cu the other calculate_ijk_a_kernel.cu

the test.cu is where it needs to be. This is the file where i call the kernel, and where i want to make a constant type.

__constant__ float a_max = 1.0f;

First I had this inside my main function. And got the message that is has to be outside the function. When I then moved it outside and tried to use it in my kernel file I got a strange error stating that a_max is undefined. I tried everything i know. Can someone please give me a hint how I need to do this. I also looked in the examples of the CUDA SDK there was only one example that was making use of constant but what I don’t get is that:

__constant__ float4 atominfo[MAXATOMS];

atominfo never gets a value assigned to it?

this is the cuenergycompsum example.

Thanks,

Jordy

You need to use “cudaMemcpyToSymbol” to copy to constant memory on the device. See the “convolutionSeparable” sample in the SDK for example code.

You will probably also need to #include both .cu files into one .cu and only copmile that one. If you link together two compiled .cu files, textures and (i believe) constant memory aren’t shared among the different compilation units.

Ok I now did this:

test.cu

// define a_max

	float *h_a_max;

	float *d_a_max;

	

	h_a_max = (float*) malloc(sizeof(float));

	cudaMemcpyToSymbol(d_a_max, h_a_max, sizeof(h_a_max));

kernel.cu

__device__ __constant__ float d_a_max;

when I run the application like this I won’t get any errors. The thing is i want to assign a value to the constant in the test.cu how am I going to do this? because when i say h_a_max = 1.0f it won’t work.

and if I do it in the kernel like

__device__ __constant__ float d_a_max = 1.0f;

it works but i have other constants i want to use which are dependent on some vars in test.cu

Something like this, I suppose (note that variable name is enclosed in quotes):

float myConst = 1.0f;

cudaMemcpyToSymbol( "d_a_max", &myConst, sizeof(float) );

Thank you very much it works…

Do you know whethere there is any particular reason for is? I though I could do that with an extern qualifier or something, but I didn’t work?

CUDA Programming guide section 4.2.2.4:

Implied static storage means that the declaration only exists with that compilation unit. If you use the extern keyword with textures (I haven’t tried it with constant memory), the compiler silently accepts it and creates separate textures in each compilation unit. So, if you bind the texture in one .cu file the other gives an error “texture not bound” when executing a kernel.