How does CUDA's cudaMemcpyFromSymbol work?

Question

I understand the concept of passing a symbol, but was wondering what exactly is going on behind the scenes. If it's not the address of the variable, then what is it?

Answer 1

I believe the details are that for each __device__ variable, cudafe creates a normal global variable as in C and also a CUDA-specific PTX variable. The global C variable is used so that the host program can refer to the variable by its address, and the PTX variable is used for the actual storage of the variable. The presence of the host variable also allows the host compiler to successfully parse the program. When the device program executes, it operates on the PTX variable when it manipulates the variable by name.

If you wrote a program to print the address of a __device__ variable, the address would differ depending on whether you printed it out from the host or device:

#include <cstdio>

__device__ int device_variable = 13;

__global__ void kernel()
{
  printf("device_variable address from device: %p\n", &device_variable);
}

int main()
{
  printf("device_variable address from host: %p\n", &device_variable);

  kernel<<<1,1>>>();
  cudaDeviceSynchronize();

  return 0;
}

$ nvcc test_device.cu -run
device_variable address from host: 0x65f3e8
device_variable address from device: 0x403ee0000

Since neither processor agrees on the address of the variable, that makes copying to it problematic, and indeed __host__ functions are not allowed to access __device__ variables directly:

__device__ int device_variable;

int main()
{
  device_variable = 13;

  return 0;
}

$ nvcc warning.cu
error.cu(5): warning: a __device__ variable "device_variable" cannot be directly written in a host function

cudaMemcpyFromSymbol allows copying data back from a __device__ variable, provided the programmer happens to know the (mangled) name of the variable in the source program.

cudafe facilitates this by creating a mapping from mangled names to the device addresses of variables at program initialization time. The program discovers the device address of each variable by querying the CUDA driver for a driver token given its mangled name.

So the implementation of cudaMemcpyFromSymbol would look something like this in pseudocode:

std::map<const char*, void*> names_to_addresses;

cudaError_t cudaMemcpyFromSymbol(void* dst, const char* symbol, size_t count, size_t offset, cudaMemcpyKind kind)
{
  void* ptr = names_to_addresses[symbol];

  return cudaMemcpy(dst, ptr + offset, count, kind);
}

If you look at the output of nvcc --keep, you can see for yourself the way that the program interacts with special CUDART APIs that are not normally available to create the mapping:

$ nvcc --keep test_device.cu
$ grep device_variable test_device.cudafe1.stub.c
static void __nv_cudaEntityRegisterCallback( void **__T22) {  __nv_dummy_param_ref(__T22); __nv_save_fatbinhandle_for_managed_rt(__T22); __cudaRegisterEntry(__T22, ((void ( *)(void))kernel), _Z6kernelv, (-1)); __cudaRegisterVariable(__T22, __shadow_var(device_variable,::device_variable), 0, 4, 0, 0); }

If you inspect the output, you can see that cudafe has inserted a call to __cudaRegisterVariable to create the mapping for device_variable. Users should not attempt to use this API themselves.