13

As stated by other questions and according to the link, you can no longer use the symbol name for this function. Now that the feature is gone, when would ever want to use this over cudaMemCpy? When would you ever want to use it at all? What is the tradeoff or benefit?

https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g9bcf02b53644eee2bef9983d807084c7

Edd Inglis
  • 1,067
  • 10
  • 22
rubixibuc
  • 7,111
  • 18
  • 59
  • 98

1 Answers1

46

In short, because cudaMemcpy can't do the same thing as cudaMemcpyToSymbol without an additional API call. Consider a constant memory array:

__constant__ float coeffs[8];

To copy values to this array using cudaMemcpyToSymbol, just do

cudaMemcpyToSymbol(coeffs, hostData, 8*sizeof(float));

To do the same with cudaMemcpy requires this:

float *dcoeffs;
cudaGetSymbolAddress((void **)&dcoeffs, coeffs);
cudaMemcpy(dcoeffs, hostData, 8*sizeof(float), cudaMemcpyHostToDevice);

A direct call to cudaMemcpy is illegal without prior symbol lookup.

[standard disclaimer: all code written in browser without access to documentation or compiler, use at own risk]

talonmies
  • 70,661
  • 34
  • 192
  • 269
  • 1
    ty, what is the relationship between `__constant__` memory and the parameter type pointer to const in the method signature. Does `__constant__` have pointer to const as it's type? Also, why can't you just take &coeffs to get the address? – rubixibuc Apr 13 '13 at 07:15
  • 5
    `__constant__` is a storage class specifier in CUDA that says the symbol is located in a special cached, read-only portion of GPU DRAM. It is completely independent of const. You canoot use `&coeffs` becuase that would imply the address of the symbol *in host memory* not GPU memory, which is what the API call requires – talonmies Apr 13 '13 at 07:31
  • Where is the address in the pointer variable coeffs pointing to, where would a copy to that adress direct memory to? Does this apply to variables declared with `__device__` as well? – rubixibuc Apr 13 '13 at 07:54
  • 1
    A copy to that address via the CUDA API would fail with an invalid argument error because it isn't an address in GPU memory space that the API had previously allocated. And yes, this applies to generic `__device__` pointers and statically declared device symbols as well. – talonmies Apr 13 '13 at 08:35
  • Does the API differentiate between taking the device address of a pointer and a non-pointer device variable with cudaGetSymbolAddress. Does it take the address of the actual pointer variable in memory, or the starting adress of where the memory being pointed is,when you make the call? How does this compare to taking the address of standard int device variable? – rubixibuc Apr 13 '13 at 21:18
  • It is important to mention that cudaGetSymbolAddress(...) should be called from the ".cu" file and not ".cuh" otherwise error number 13 is raise. – ddleon Jul 11 '21 at 12:21