CUDA的cudaMemcpyFromSymbol如何运作？

Question

我理解传递符号的概念,但想知道幕后到底发生了什么.如果它不是变量的地址,那么它是什么？

Answer 1

我相信细节是每个__device__变量,cudafe创建一个普通的全局变量,如C和CUDA特定的PTX变量.使用全局C变量,以便主机程序可以通过其地址引用变量,并且PTX变量用于变量的实际存储.主变量的存在还允许主机编译器成功解析程序.当设备程序执行时,它在按名称操作变量时对PTX变量进行操作.

如果您编写了一个程序来打印__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

由于两个处理器都不同意变量的地址,这使得复制到它有问题,并且实际上__host__不允许函数__device__直接访问变量:

__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允许从__device__变量复制数据,前提是程序员恰好知道源程序中变量的(受损)名称.

cudafe通过在程序初始化时创建从错位名称到变量的设备地址的映射来促进这一点.程序通过查询CUDA驱动程序获取驱动程序令牌的错误名称来发现每个变量的设备地址.

所以cudaMemcpyFromSymbol在伪代码中看起来像这样的实现:

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);
}

如果查看输出nvcc --keep,您可以自己查看程序与通常无法创建映射的特殊CUDART API交互的方式:

$ 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); }

如果检查输出,则可以看到cudafe已插入调用__cudaRegisterVariable以创建映射device_variable.用户不应尝试自己使用此API.