CUDA:如何从单独的编译中 link 特定的 obj、ptx、cubin?

CUDA: How to link a specific obj, ptx, cubin from a separate compilation?

我有一个相当大的 CUDA/C++ 编译成静态库的项目。工具链是CUDA Toolkit 9.0/9.2和VS 2017。我无法更改公司工具链。我们最昂贵的内核受到 9.0 工具包中引入的 nvcc 编译器回归的影响。我已将此提交给 Nvidia 开发人员的网站,并收到了回归确认。那是大约一年前的事了,门票仍然开放。也许 10.0 工具包会修复它。

但我等不及了。所以我的计划是使用 8.0 nvcc 编译器和 v140 (VS 2015) 编译器只编译这个特定的内核。它是一个单独的 .hpp 文件,带有用于内核声明的 __device__ 装饰器,以及一个带有定义的 .cu 文件。内核不调用其他内核;这是一个相当简单的内核。

在 v140 Native Tools 命令提示符中,我执行了:

nvcc -x cu -arch=sm_61 -dc kernel.cu

并获得了一个kernel.obj文件。我已经阅读了NVCC documentation on CUDA Compiler Driver NVCC。我承认不完全理解。有几个编译阶段,我看不出哪个是适合我的情况。

我的问题是如何link这个目标文件到我更大的静态库中?如果有人能指出正确的命令系列,或者更好的是,如何将其包含到 VS 项目中,大概使用 kernel.hpp 和 kernel.obj,我将不胜感激。

根据 Njuffa 的上述评论,最简单的解决方案是使用该内核的早期高性能工具链(VS 2015 和 CUDA 8.0 Tookit)创建静态库。然后 link 将该库与后来的工具链一起放入更大的项目中。我成功了。

我在 VS 2015 中创建了一个 CUDA 8.0 模板项目,只有内核源代码和 header。编译目标设置为静态库。这创建了一个 .lib 文件。然后使用 VS 2017 和 CUDA 9.0 将 .lib 文件和 header 添加到更大项目的 C++ linker 设置中。使用此静态库的所有测试可执行文件均通过。这是一个比尝试使用中间编译格式(ptx、cubin 等)重新编译更简单的解决方案

尽管最终,真正的解决方案是重构内核以更有效地使用共享内存,从而不再需要旧的 nvcc 版本。