__constant__ and cudaGetSymbolAddress


Documentation reads that

So I expect following code to work… but it doesnt and error is cudaErrorAddressOfConstant or something.

const int MAX_SIZE = 32;

__constant__ uint1 vectors[MAX_SIZE];

int main(){

       int cards;


        cudaSetDevice(cards - 1);

       uint1* adres;

       if (cudaSuccess == cudaGetSymbolAddress((void**)&adres, "vectors")) {

                std::cout << " OK " << std::endl;

        } else {

                std::cout << "not OK" << std::endl;




Got exactly the same problem on CUDA 1.0. However the 1.1 documentation says that the symbol could either be a name or a variable residing in global memory. So it seems we cannot get adresses of constant memory variables. Indeed

const int MAX_SIZE = 32;

__device__ __constant__ uint1 vectors[MAX_SIZE];




  cudaError_t err;

 void* ptr; 

  err = cudaGetSymbolAddress ((void**) &ptr, vectors);

  printf("%p\tError: %s\n", ptr, message_from_cuda_error(err));


  size_t sz;

  err = cudaGetSymbolSize (&sz, vectors);

  printf("%d\tError: %s\n", sz, message_from_cuda_error(err));

  return 0;


Returns a AddressOfConstant error for the first call, but the correct bytes count for the second call (32×4=128).

Seems that we must use cudaMemcpyToSymbol to set a constant variable content and never pass it as reference :-/

Ok, So how does cudaMemcpyToSymbol behave when I use streaming and I have device with compute capability 1.1 (uses stream 0, blocks, something else)?

It would be nice to have cudaMemcpyToSymbolAsync if possible.

Thanks for your help. Using memcpytosymbol seems do have done the trick, gave a 40% speedup, and worked correctly. Still only manages to encrypt 650mbit/s tho.

I have the exact same problem with CUDA 2.3.

Am I right that cudaGetSymbolAddress() cannot fetch the address of a constant variable and thus returns a cudaErrorAddressOfConstant error, and that, in contrast, cudaGetSymbolSize() is able to return the size of a constant variable?