Problem setting __constant__ memory

Hi all. I am having a problem with the setting of constant memory using the cudaCpymemToSymbol function. The guide is not very clear on right way to go about doing this and so I do not know if I am the one making the mistake, or if it is a possible bug.

The code is as follows:

__constant__ float min;

float tmp = 1.0f;

CUDA_SAFE_CALL( cudaMemcpyToSymbol( min , &tmp , sizeof( float ) , 0 ) );

__constant__ float3 pos;

float3 cPos;

... ( cPos is set to some value )

CUDA_SAFE_CALL( cudaMemcpyToSymbol( pos , &cPos , sizeof( float3 ) , 0 ) );

This code is executed on the host. Getting the constant data values in the kernel, resulted in min always having the value 0, but pos had the correct values. I assumed that it might not be possible to allocate a single float value in constant memory, so I constructed a single struct to hold all the constant data I need ( which is just 32 bytes in size ).

typedef struct


	int voxelGridSize;

	float voxelSize;

	float originOffset;

	float minPos;

	float maxPos;

	float3 beamPosition;

} ConstantData;

__constant__ ConstantData __constantData;

ConstantData tmp;

... ( values of tmp are set )

CUDA_SAFE_CALL( cudaMemcpyToSymbol( __constantData , &tmp , sizeof( ConstantData ) ) );

This struct shows the exact same behaviour: the float3 pos value is correctly set to the right values, but al single float values in the struct have bogus values assigned to them.

I am using CUDA version 0.8 on a 8800GTX, 8 processor Xeon E5320 machine running Windows XP SP2.

Some help will be sincerely appreciated! Thanks.

__constant__ float min;

float tmp = 1.0f;

CUDA_SAFE_CALL( cudaMemcpyToSymbol( min , &tmp , sizeof( float ) , 0 ) );

From your system spec I see that you are using the VC compiler. It defines min as a macro that you need to avoid by #define NOMINMAX . The CUDA SDK projects do that in the vcproj settings. Make sure this is not an issue in your project.

For the struct - the VC compiler sometimes has trouble figuring the sizeof structs. Try the following: define the struct rather than declaring a type


Thanks for the tip. I renamed the variable, but I do not think this was a problem.

I was actually using the Intel compiler v9.0 in VS7, which I forgot to mention. There was obviously something going wrong with the struct alignment. I tried some compiler switches, but it did not help. Switching to the VC++ compiler solved the issue, but this reduces the performance of the needed C host code considerably. I could not find much about Intel compiler support in Windows. Will it be supported sometime in the future, or should I stick with the VC compiler?

Also puzzling is the fact that single floats were not being set correctly in constant memory. Is this not supported? I haven’t tested it yet with VC++.

Thanks for the help.

Setting single floats is supported. The Intel compiler might not understand the alignment statements correctly if they are part of a struct declaration (the float3 declaration has one). Try adding a CUDA align(16) to your struct declaration.


I added the aligns and all seems to be working perfectly right now using CUDA 0.8 and the Intel C++ compiler v9.0 in VS7.0. Thanks for the help!