CUDA跨编译单元下cudaMemcpyFromSymbol获取函数指针异常咨询

问题背景 #

尝试获取CUDA设备函数指针并传递给核函数，实现了辅助函数GetDeviceSymbolAddress——通过const引用接收设备符号，借助cudaMemcpyFromSymbol获取其值（设备函数指针）。但发现当核函数定义在不同编译单元（不同.cu文件）时，GetDeviceSymbolAddress返回的地址无效，导致核函数调用失败。可通过定义/取消定义SWITCH宏切换核函数所在的编译单元。

示例代码 #

test.cuh #

#ifndef TEST_H
#define TEST_H

#include <iostream>

#define SWITCH 1

using pfloat__fn_pvoid = float* (*)(const void*);

// Test and log Cuda errors
#define CUDA_CHECK(err)                                     
  do {                                                      
    auto e = err; /* err might be a function call */        
    if (e != cudaSuccess) {                                 
      std::cout << "CUDA ERROR: " << cudaGetErrorString(e); 
      abort();                                              
    }                                                       
  } while (0)

template <typename T>
T GetDeviceSymbolAddress(const T& dvc_symbol) {
  T hst_symbol{};
  CUDA_CHECK(cudaMemcpyFromSymbol(&hst_symbol, dvc_symbol, sizeof(T)));
  std::cout << (void*)hst_symbol << std::endl;
  return hst_symbol;
}

void RunTest(pfloat__fn_pvoid result_fn);

#endif /* TEST_H */

test.cu #

#include "test.cuh"

#ifdef SWITCH

__global__ void Test(pfloat__fn_pvoid fn) {
  fn(nullptr);
}

void RunTest(pfloat__fn_pvoid fn) {
  Test<<<1, 1>>>(fn);
  CUDA_CHECK(cudaDeviceSynchronize());
}

#endif

test1.cu #

#include "test.cuh"

__device__ float* Foo(const void* data) {
  return nullptr;
}

__device__ pfloat__fn_pvoid pFoo = Foo;

#ifndef SWITCH

__global__ void Test(pfloat__fn_pvoid fn) {
  fn(nullptr);
}

void RunTest(pfloat__fn_pvoid fn) {
  Test<<<1, 1>>>(fn);
  CUDA_CHECK(cudaDeviceSynchronize());
}
#endif

int main() {
  RunTest(GetDeviceSymbolAddress(pFoo));
  return 0;
}

问题解答 #

1. 核函数的编译单元位置为何会影响调用结果？ #

本质是跨编译单元的设备符号解析逻辑问题：

• 当核函数与__device__变量pFoo在同一编译单元（SWITCH未定义时），编译器能直接关联主机端代码与设备端符号，cudaMemcpyFromSymbol可以正确获取到设备函数指针的有效地址，传递给核函数后能正常调用。
• 当核函数在另一编译单元（SWITCH定义时），CUDA默认将__device__符号视为内部链接（类似C++的static），不同编译单元会把pFoo当成独立符号。此时GetDeviceSymbolAddress中传递的dvc_symbol是主机端生成的无效代理对象，导致cudaMemcpyFromSymbol无法定位到设备端的实际符号，返回的地址自然无效，核函数调用失败。

2. cudaMemcpyFromSymbol的调用编译单元是否会对结果产生影响？ #

会产生影响，核心在于设备符号的可见性与编译时解析逻辑：

• 当cudaMemcpyFromSymbol的调用代码与目标设备符号在同一编译单元时，编译器能直接将符号的设备端地址信息嵌入调用逻辑，确保函数能正确获取到符号值。
• 跨编译单元调用时，如果没有显式声明设备符号的外部可见性（比如用extern __device__），主机端的符号代理无法正确映射到设备端的实际符号，cudaMemcpyFromSymbol会因无法找到有效符号而返回错误地址。此外，若通过变量引用传递符号而非字符串名称，跨编译单元的符号解析差异会进一步放大这个问题。

修复建议 #

1. 显式声明跨编译单元的设备符号：在test.cuh中添加extern __device__ pfloat__fn_pvoid pFoo;，确保所有编译单元都能正确识别该符号为同一设备端实体。
2. 改用符号名称调用cudaMemcpyFromSymbol：修改GetDeviceSymbolAddress，让它接收符号名字符串而非变量引用，避免跨编译单元的代理对象问题：

template <typename T>
T GetDeviceSymbolAddress(const char* symbol_name) {
  T hst_symbol{};
  CUDA_CHECK(cudaMemcpyFromSymbol(&hst_symbol, symbol_name, sizeof(T)));
  std::cout << (void*)hst_symbol << std::endl;
  return hst_symbol;
}

调用时改为RunTest(GetDeviceSymbolAddress<pfloat__fn_pvoid>("pFoo"));

内容的提问来源于stack exchange，提问作者Gils