passing constant memory

Is there a way to pass the address of a constant array from host to kernel? cudaGetSymbolAddress() complains that it “cannot take the address of constant data.”

no, no way, pls check former similar posts.

constant variables cannot be assigned to from the device, only frohost. They are therefore only allowed at file scope.

In my understanding, constant is like global static var. I’ve provided some workable file framework for using constant, in prior posts.

I don’t understand the relevance of the quoted section. I’m not writing to constant memory from the device. I have a kernel which reads from constant memory arrays, and I have several different arrays in constant memory. I would like to be able to tell the kernel which one to read from. My confusion is over why the host function cannot take the address of a constant symbol using cudaGetSymbolAddress() even when it can write to the same symbol using cudaMemcpyToSymbol().

You are absolutely right. yk_cadcg’s comment is misleading. You can get the address. Maybe there is something wrong with your call? Can you post the code?

I did not use selecting from multiple const arrays yet. I always pack all consts into one large block, then have an enum that holds the offsets into the block. I then pass the offset to be used as a kernel argument. As the const array is defined in the CUDA code, the kernel just can do something like array[kernelparam] to get to the correct spot. I put the enum in the .h file that declares the CUDA interface and I can therefore use the same array[enum] on the host side to fill the const array prior to upload.