CUDA Driver API全局设备指针内存拷贝报错0001的解决求助

问题场景 #

在CUDA内核中声明了__device__类型的全局设备指针p_GLOBAL_HOLDER_Picture，尝试通过CUDA Driver API从主机端映射该指针并拷贝主机内存到目标地址，但执行cuMemcpyHtoD时触发CUDA API error = 0001（CUDA_ERROR_INVALID_VALUE）错误。

错误原因 #

__device__ unsigned int * p_GLOBAL_HOLDER_Picture是一个设备端的指针变量，它自身仅占用设备全局内存中一个指针大小的空间（64位系统为8字节）。当前代码错误地直接向这个指针变量的地址拷贝8个unsigned int的数据（共32字节），拷贝大小远超过该变量的实际占用空间，导致参数无效错误。

正确逻辑应该是：先在设备端分配一块能存储目标数据的内存，再将这块内存的地址赋值给p_GLOBAL_HOLDER_Picture，最后把主机数据拷贝到这块设备内存中。

解决方案 #

修改主机端test.cpp的关键代码部分：

修改后的核心代码 #

int main(){
  // ... 原有初始化代码 ...

  int count = 8;
  const int SIZE = sizeof(unsigned int);
  unsigned int * ptr = (unsigned int*)malloc(count * SIZE);
  // 初始化主机数据（建议添加，避免读取随机值）
  for(int i=0; i<count; i++) ptr[i] = i;

  // ... 加载模块、获取函数等代码 ...

  // 1. 获取__device__指针变量在设备全局内存中的地址
  CUdeviceptr global_ptr_addr;
  size_t global_ptr_size;
  checkCudaResult(cuModuleGetGlobal(&global_ptr_addr, &global_ptr_size, modId, "p_GLOBAL_HOLDER_Picture"));

  // 2. 在设备端分配存储实际数据的内存
  CUdeviceptr device_data;
  checkCudaResult(cuMemAlloc(&device_data, count * SIZE));

  // 3. 将设备数据内存的地址赋值给__device__指针变量
  checkCudaResult(cuMemcpyHtoD(global_ptr_addr, &device_data, global_ptr_size));

  // 4. 拷贝主机数据到设备数据内存中
  checkCudaResult(cuMemcpyHtoD(device_data, ptr, count * SIZE));

  // ... 启动内核、同步等代码 ...

  // 释放资源：先释放设备数据内存，再释放主机内存
  checkCudaResult(cuMemFree(device_data));
  free(ptr);
  checkCudaResult(cuCtxDestroy(ctx));
}

关键说明 #

1. cuModuleGetGlobal获取的是p_GLOBAL_HOLDER_Picture变量本身的设备地址，我们需要给这个变量赋值一个有效的设备内存地址，而非直接拷贝数据到此处。
2. 先通过cuMemAlloc分配存储实际数据的设备内存，再将该内存地址写入p_GLOBAL_HOLDER_Picture，确保内核中访问p_GLOBAL_HOLDER_Picture[0]时指向正确的内存区域。
3. 错误码0001对应CUDA_ERROR_INVALID_VALUE，原代码因拷贝大小与目标内存空间不匹配触发该错误，修正后参数匹配即可解决。

内容的提问来源于stack exchange，提问作者Sai Krishna Sirimisetty