__constant__ and cudaGetSymbolAddress

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

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;

        cudaGetDeviceCount(&cards);

        cudaSetDevice(cards - 1);

       uint1* adres;

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

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

        } else {

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

        };

       cudaThreadExit();

};
2

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];

int

main()

{

  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 :-/

3

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.

4

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.

5

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?

Thx