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

问题描述

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

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

#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();
}

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

标签： pointerscuda

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

不，不是。

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

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

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

看上面。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 公开，并且设备端链接器可以（现在）在链接期间静态分配正确的值，以便在运行时初始化时，该值是正确的。但请注意，它是静态分配，在运行时不会发生任何事情。

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

问题描述

解决方案

推荐阅读