CUDA CUBIN对象是否向后兼容？

Question

我的应用程序使用CUDA内核进行大部分计算。出于很好的原因(超出了这个问题的范围)，我使用共享对象/链接模型来动态加载目标文件，每个目标文件包含一个主机函数和一个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(CUDA4.1，NVIDIA C2050)上编译/构建共享对象，并稍后在机器A (CUDA4.0，GTX480)上dlopen它们时，计算产生的结果并不相同，就好像共享对象也是在机器A上构建的一样。

对我来说这听起来很奇怪。在.o文件中是否嵌入了一个立方体对象，其中包含独立于特定图形处理器架构的指令？

我知道建议在构建和链接时使用相同的编译器版本。同样，我有很好的理由不在执行共享对象时在同一台机器上构建它们。

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。它的可移植性要好得多(在最低驱动程序版本要求范围内)，但在您的主机代码中也需要更多的工作。是你的选择。