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 胖二进制对象。这两件事是不一样的。

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

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