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 交互的方式，这些 CUDART 通常不能用于创建映射：

$ 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。

CUDA 的 cudaMemcpyFromSymbol 是如何工作的？

How does CUDA's cudaMemcpyFromSymbol work?

cuda