为什么我需要一个中间全局来获取 CUDA 设备端函数地址？

Question

我一直在思考这个问题的答案：

How to pass device function as an input argument to host-side function?

尤其是 Robert Crovella 的 answer。我不太明白为什么需要中间全局符号。也就是说，为什么会这样：

#include <stdio.h>

__device__ int f1(){ printf("dev f1\n"); return 0;}

__device__ void *fptrf1 = (void*) f1;

__global__ void mykernel(int (*fptr)()) {
  fptr();
  printf("executed\n");
}

int main() {
  void *hf1;
  cudaMemcpyFromSymbol(&hf1, fptrf1, sizeof(int *));
  mykernel<<<1,1>>>((int (*)())hf1);
  cudaDeviceSynchronize();
}

但这不起作用：

#include <stdio.h>

__device__ int f1(){ printf("dev f1\n"); return 0;}

__global__ void mykernel(int (*fptr)()) {
  fptr();
  printf("executed\n");
}

int main() {
  void *hf1;
  cudaMemcpyFromSymbol(&hf1, f1, sizeof(int *));
  mykernel<<<1,1>>>((int (*)())hf1);
  cudaDeviceSynchronize();
}

？我的意思是，函数不是符号吗？如果全局指针的设备端地址可以被我的主机端代码“知道”，为什么函数本身不能呢？ 如果它不起作用 - 为什么它会编译而不是抱怨？

Answer 1

我的意思是，函数不是符号吗？

不，不是。

我在这里没有什么特别的见解，但毫无疑问，部分原因是历史原因：当 CUDA API 被发明时，__device__ 函数只是一种编程辅助。没有 ABI，没有函数指针支持，所有设备函数都被编译器内联扩展。发出的唯一静态设备符号是__global__ 函数、纹理引用和__device__ 变量。因此，在 15 年前将语言和 API 放在一起时，绝对不可能设想或可能使用这种用法。

即使使用今天的 post ABI 和 post ELF 格式的设备工具链（最初所有内容都是带有嵌入字符串的纯文本），您也不会发现设备对象文件 ELF 接口公开的 __device__ 函数。与 __global__ 函数和其他设备符号不同，无法通过任何主机 API 检索任意 __device__ 函数。

如果我的主机端代码可以“知道”全局指针的设备端地址，为什么函数本身不能呢？

见上文。 API 从未公开过这一点。

如果它不起作用 - 为什么它会编译而不是抱怨？

因为编译轨迹。 CUDA 前端对主机代码中的 __device__ 函数执行此操作（这里没有区别，它对每个 __device__ 函数执行此操作，包括内部工具链函数和设备库）：

# 3 "unobtainium.cu"
__attribute__((unused)) int f1() {int volatile ___ = 1;::exit(___);}
#if 0
# 3
{ printf("dev f1\n"); return 0; } 
#endif

即它创建一个虚拟主机存根，以便一切都编译。内核和设备符号也有存根，但具有不同的样板。这些样板存根与内部运行时函数用来使主机端运行时 API 工作的标签相匹配。但设备函数不会，因为它们没有被 CUDA 设备代码 API 公开。

最后是你原来的问题：

为什么会这样：

#include <stdio.h>

__device__ int f1(){ printf("dev f1\n"); return 0;}

__device__ void *fptrf1 = (void*) f1;

__global__ void mykernel(int (*fptr)()) {
  fptr();
  printf("executed\n");
}

int main() {
  void *hf1;
  cudaMemcpyFromSymbol(&hf1, fptrf1, sizeof(int *));
  mykernel<<<1,1>>>((int (*)())hf1);
  cudaDeviceSynchronize();
}

这里有趣的是它并不总是有效。曾几何时，您必须运行设置内核来初始化设备端函数指针。在 CUDA 5 附近的某个地方，它开始以这种方式工作。为什么相对简单 - 编译单元范围 __device__ 变量是有效的设备符号，因此由主机 API 公开，并且设备端链接器可以（现在）静态在链接期间分配正确的值以便在运行时初始化时，该值是正确的。但请注意，它是静态赋值，在运行时不会发生任何事情。