c - CUDA CUBIN 对象是否向后兼容？

Question

我的应用程序使用 CUDA 内核进行大部分计算。出于充分的理由(超出此问题的范围)，我使用共享对象/链接模型来动态加载对象文件，每个对象文件包含 1 个主机函数和 1 个 CUDA 内核。由于内核不能extern，因此内核的基本结构是:

__global__ kernel() { ...code... }

extern "C" void call_kernel() {
  <<<GRID,BLOCK,SHMEM>>>kernel();
}

我使用一个主机函数，其唯一目的是调用内核。为了构建我使用的共享对象:

nvcc -arch=sm_20 -m64 --compiler-options -fPIC,-shared -link -o kernel0.o kernel0.cu

整个应用程序使用了很多这样的内核，并使用 dlopen() 加载它们。如果一切(构建/加载/执行)都在一台机器 A 上，那么整个事情就可以正常进行。

但是当我在机器 B(cuda 4.1，NVIDIA C2050)上编译/构建共享对象时，稍后在机器 A(cuda 4.0，GTX 480)上 dlopen 它们时，计算不会产生结果与共享对象也是在机器 A 上构建的结果相同。

这对我来说听起来很奇怪。 .o 文件中是否嵌入了一个 CUBIN 对象，其中包含独立于特定 GPU 架构的指令？

我知道建议使用相同的编译器版本进行构建和链接。同样，我有充分的理由为什么不在执行共享对象时在同一台机器上构建共享对象。

Answer 1

我想首先要说明的是，您在应用程序中根本没有使用 CUBIN 文件，您使用的是 CUDA fat 二进制对象。这两件事是不一样的。

但这不是问题的根源。您的问题出在 CUDA 运行时库上。运行时 API 是版本化的，为给定运行时 API 版本编译的任何代码都必须使用该版本运行。此外，运行时 API 库版本具有最低驱动程序版本要求。您不能采用针对 CUDA 4.1 库构建的应用程序并期望在具有 CUDA 4.0 库的机器上运行它。 NVIDIA 推荐的分发运行时 API 代码的方法是将运行时 API 库 (libcudart) 与您的应用程序一起分发，并指定您的代码所需的最低驱动程序版本。这确保应用程序将正确运行(最低驱动程序版本带来最低 CUDA 驱动程序 API 版本，分布式运行时 API 库带来完成要求)。

另一种方法是真正使用 cubin 文件并使用 CUDA 驱动程序 API。它更具可移植性(在最低驱动程序版本要求内)，但在您的主机代码中也需要更多工作。选择权在您。