Is it possible to write to kernel constants ?


I’ve been wondering now and then how to setup kernel constants from the host side.

Since these just constants it should be possible to set these up before the kernel actually runs.

CUDA C Programming Guide section B.2.2. mentions how to retrieve a “handle” or “cuda pointer” if you will… to these constants by “quering/retrieving” them by name via api’s.

It mentions nice api’s for the runtime api which can be used to copy to and from symbols, for the driver api it only mentions a “get” function.

So my question is pretty simple really: Can the copy functions from the driver api be used to write and read values to and from these constants via the retrieved cuda pointer ?

My guess would be yes… but this would be nice to know… if so then guide could also be updated a little bit and mention that “cuda driver api memory copy functions” can be used to read/write the constants in the kernel/cuda device. ;)

Perhaps it should even have a special section which is simply called: “reading and writing to constants” that’ll be nice ;) :)


While I haven’t used the driver api much, I don’t see why it should work. I’ve always had the impression that the Runtime api is just a wrapper of the driver api anyway. Structures between driver api and runtime api are often interchangeable which I guess indicates some sort of compatibility.

You can write a simple test program.

Are you talking about variables with a constant qualifier in front of them?

If so, there is no way that you can write to them from within device code.

If you want to write to a variable that is cached in a similar manner to constant variables, then you could try ‘surface’ references. I think the cache for surface/texture fetches is about 8kb, maybe 4kb I don’t quite remember. I do know that you can only have up to 8 surface references bound at a given time.

What I ment is this kernel code:

constant int a;



host code:

int InitConstant = 5000;

CudaPointer := cuGetSymbolNameOrSomething( Device, Module, “a” );

MemCopyHostToDevice( CudaPointer, InitConstant, 4 );

From the looks of it this should be possible, because it returns a cuda pointer…

Variables declared with the constant qualifier are strictly read only as far as the device code is concerned.

You have to remember that just because something is a device pointer, doesn’t mean that it behaves like a normal pointer. There is probably some hardware specific reason why device code can’t write to constant memory, and probably has something to do with the performance of reads from constant memory.

I can understand if device code can’t write to constants.

But what about host code writing to device code via memory copy functions ? ;)

Somehow constants should be initialized anyway… so I see no reason why host code couldn’t do that ;)

I guess I will have to try it out sometime ;) :)

Suppose it wasn’t possible, then it would still be possible to modify the ptx strings if they are embedded or change the files when they are loaded… and then ptx assembly instructions could be changed so their constant values or changed so one way or the other it’s probably possible ;)

Oh, you weren’t talking about only the device code. In that case, yes it is possible to write to constant memory using cudaMemcpyToSymbol(). I’m not sure that you can do it with cudaMemcpy() even after you get the device pointer with cudaGetSymbolAddress().

Constant memory variables may behave like texture references in some ways, and if you have ever dealt with texture references you’d know how much of a pain those are. But I’m not entirely certain on this point.

If you have the book CUDA By Example, there is a very good section on using constant memory.

If you’re using the Driver API, and you have some constant variable declared in your kernel, you can get a reference to it on the host with cuModuleGetGlobal(), then memcpy to it for initialization (via cuMemcpyHtoD)


in the kernel:

__constant__ float constVal[LENGTH];

in the host:

CUdeviceptr constDevPtr;

cuModuleGetGlobal(&constDevPtr, NULL, cumodule, "constVal");

cuMemcpyHtoD(constDevPtr, h_data, LENGTH * sizeof(float));

where the cumodule is the loaded PTX file corresponding to the .cu file w/ the constant variable, and h_data is the host-based memory to copy to the constant