分离库的主机端和 CUDA 设备端版本

Question

我有一个包含一些 __host__ __device__ 函数的库。我还有一个 #ifdef __CUDACC__ 小工具，它可以确保常规 C++ 编译器看不到 __host__ __device__，从而可以编译这些函数。

现在，我想在普通 C++ 静态库文件（Linux 上的.a）中使用库函数的编译主机端版本——我什至希望在 CUDA 不可用时可以编译该库;我希望在单独的静态库中编译设备端版本。

我快到了（我想），但遇到链接错误。以下是此类库的玩具源代码、测试程序（调用设备端和主机端版本的函数）和我使用的构建命令。

我做错了什么？

---

• my_lib.hpp（库头）：

#ifdef __CUDACC__
__host__ __device__
#endif
void foo(int*x, int* y);
int bar();

• my_lib.cu（库来源）：

#include "my_lib.hpp"

#ifdef __CUDACC__
__host__ __device__
#endif
void foo(int*x, int* y)  { *x = *y; }

int bar() { return 5; }

• main.cu（测试程序）：

#include "my_lib.hpp"

__global__ void my_kernel() {
  int z { 78 };
  int w { 90 };
  foo(&z,&w);
}

int main() {
  int z { 123 };
  int w { 456 };
  foo(&z,&w);
  my_kernel<<<1,1>>>();
  cudaDeviceSynchronize();
  cudaDeviceReset();
}

我的构建命令：

c++ -c -x c++ -o my_lib-noncuda.o my_lib.cu
ar qc my_lib-noncuda.a my_lib-noncuda.o
ranlib my_lib-noncuda.a
nvcc -dc -o my_lib-cuda.o my_lib.cu
ar qc my_lib-cuda.a my_lib-cuda.o
ranlib my_lib-cuda.a
nvcc -dc -o main.rdc.o main.cu
nvcc -dlink -o main.o main.rdc.o my_lib-cuda.a
c++ -o main main.o my_lib-noncuda.a -lcudart

以及我得到的错误 - 最后一个，链接，命令：

/usr/bin/ld: main.o: in function `__cudaRegisterLinkedBinary_39_tmpxft_00003f88_00000000_6_main_cpp1_ii_e7ab3416':
link.stub:(.text+0x5a): undefined reference to `__fatbinwrap_39_tmpxft_00003f88_00000000_6_main_cpp1_ii_e7ab3416'
/usr/bin/ld: main.o: in function `__cudaRegisterLinkedBinary_41_tmpxft_00003f69_00000000_6_my_lib_cpp1_ii_ab44b3f6':
link.stub:(.text+0xaa): undefined reference to `__fatbinwrap_41_tmpxft_00003f69_00000000_6_my_lib_cpp1_ii_ab44b3f6'
collect2: error: ld returned 1 exit status

注意事项：

• 我在 Devuan GNU/Linux 上使用 CUDA 10.1 和 g++ 9.2.1。
• 这是对已删除问题的“跟进”； @talonmies 评论说我最好准确地展示我的所作所为；这在一定程度上改变了问题。
• 有点相关的问题：this one。

Answer 1

以下是创建两个库的方法，一个仅包含 CUDA 设备函数，另一个仅包含主机函数。 您可以省略“复杂的”#if 和 #ifndef 守卫。但是，您的库中也会有“非 CUDA 代码”my_lib-cuda.a。

对于其他问题，请参阅@talonmies 社区 wiki 答案或参考我已经在 cmets 中发布的链接：https://devblogs.nvidia.com/separate-compilation-linking-cuda-device-code/ - “高级用法：使用不同的链接器”部分。

my_lib.cu

#include "my_lib.hpp"

#ifdef __CUDA_ARCH__
__device__
#endif
#if (defined __CUDA_ARCH__) || (not defined __CUDACC__)
void foo(int*x, int* y)  { *x = *y; }
#endif

#ifndef __CUDACC__
int bar() { return 5; }
#endif

库的构建过程保持不变：（仅将 ar qc 更改为 ar rc 以替换现有文件，因此在不事先删除库的情况下重新构建时不会出错）

c++ -c -x c++ -o my_lib-noncuda.o my_lib.cu
ar rc my_lib-noncuda.a my_lib-noncuda.o
ranlib my_lib-noncuda.a
nvcc -dc -o my_lib-cuda.o my_lib.cu
ar rc my_lib-cuda.a my_lib-cuda.o 
ranlib my_lib-cuda.a 

构建 CUDA 程序：（通过仅使用 nvcc 而不是 c++ 进行简化，或者查看 @talonmies 社区 wiki 答案）

nvcc -dc main.cu -o main.o
nvcc main.o my_lib-cuda.a my_lib-noncuda.a -o main

如上所述，如果您还省略了my_lib.cu 中的#if 和#ifndef，则可以省略指向my_lib-noncuda.a 的链接。

构建一个 C++ 程序：（假设在 main.cu 的 CUDA 代码周围有 #ifdef __CUDACC__ 守卫）

c++ -x c++ -c main.cu -o main.o
c++ main.o my_lib-noncuda.a -o main

Answer 2

让我们将您的示例修改为我认为您的实际用例。修改将main()放入.cpp文件，由g++编译，将CUDA代码放入单独的.cu文件，由nvcc编译。这对于使您的两个库设置工作很重要；并且合理，因为“main 包含需要单独编译和链接的 CUDA 内核”是 nvcc 编译模型的特殊极端情况。

重组后的代码：

main.cu：

include "my_lib.hpp"

__global__ void my_kernel() {
  int z { 78 };
  int w { 90 };
  foo(&z,&w);
}

int cudamain()
{
  my_kernel<<<1,1>>>();
  return 0;
}

main.cpp：

#include <cuda_runtime_api.h>
#include "my_lib.hpp"

extern int cudamain();

int main() {
  int z { 123 };
  int w { 456 };
  foo(&z,&w);
  cudamain();
  cudaDeviceSynchronize();
  cudaDeviceReset();
}

所有其他文件都保留在问题中。

现在构建程序所需的命令是：

c++ -c -x c++ -o my_lib-noncuda.o my_lib.cu
ar qc my_lib-noncuda.a my_lib-noncuda.o
ranlib my_lib-noncuda.a

nvcc -std=c++11 -dc -o my_lib-cuda.rdc.o my_lib.cu
ar qc my_lib-cuda.a my_lib-cuda.rdc.o
ranlib my_lib-cuda.a

# Until this line - identical to what you have tried in your question

nvcc -std=c++11 -c -rdc=true main.cu -o main.cu.o 
nvcc -dlink -o main.o main.cu.o my_lib-cuda.a

c++ -std=c++11 -o main main.cpp main.o main.cu.o -I/path/to/cuda/include \
    -L/path/to/cuda/lib64 my_lib-cuda.a my_lib-noncuda.a -lcudart -lcudadevrt

需要牢记的重要一点是主机端组件需要在构建中进行。因此，您必须将 CUDA 主机代码的nvcc 输出传递给主链接，并且您还必须将您的 CUDA 端库添加到主链接。否则，您的代码的主机端运行时 API 支持将丢失。另请注意，您必须链接设备运行时库才能使其工作。