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

Question

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

如何将设备函数作为输入参数传递给主机端函数？

尤其是罗伯特·克罗维拉的回答。我不太明白为什么中间全局符号是必要的。也就是说，为什么这样做：

#include <stdio.h>
__device__ int f1(){ printf("dev f1n"); return 0;}
__device__ void *fptrf1 = (void*) f1;
__global__ void mykernel(int (*fptr)()) {
fptr();
printf("executedn");
}
int main() {
void *hf1;
cudaMemcpyFromSymbol(&hf1, fptrf1, sizeof(int *));
mykernel<<<1,1>>>((int (*)())hf1);
cudaDeviceSynchronize();
}

但这不起作用：

#include <stdio.h>
__device__ int f1(){ printf("dev f1n"); return 0;}
__global__ void mykernel(int (*fptr)()) {
fptr();
printf("executedn");
}
int main() {
void *hf1;
cudaMemcpyFromSymbol(&hf1, f1, sizeof(int *));
mykernel<<<1,1>>>((int (*)())hf1);
cudaDeviceSynchronize();
}

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

Answer 1

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

不，不是。

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

即使使用今天的后 ABI 和后 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 f1n"); return 0; } 
#endif

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

最后是你最初的问题：

为什么这样做：

#include <stdio.h>
__device__ int f1(){ printf("dev f1n"); return 0;}
__device__ void *fptrf1 = (void*) f1;
__global__ void mykernel(int (*fptr)()) {
fptr();
printf("executedn");
}
int main() {
void *hf1;
cudaMemcpyFromSymbol(&hf1, fptrf1, sizeof(int *));
mykernel<<<1,1>>>((int (*)())hf1);
cudaDeviceSynchronize();
}

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