CUDA nvcc 构建库链

Question

我的目标是：library2.so 使用 library1.so 并且 mycode.o 使用（应链接库）library2.so（也许 library1.so）。

源码为（省略一行头文件）：

library1.cu:

__device__ void func1_lib1(void){}

library2.cu:

#include "library1.h"
__global__ void func1_lib2(void)
{
    func1_lib1();
}
extern "C"
void func2_lib2(void)
{
    func1_lib2<<<1,1>>>();
}

mycode.c:

#include "library2.h"
int main(void)
{
    func2_lib2();
}

我正在使用 Makefile

构建共享库 according to

broken:
    rm -f *.o *.so
    nvcc -arch=sm_30 --compiler-options '-fPIC' -dc library1.cu
    nvcc -arch=sm_30 --compiler-options '-fPIC' -dlink -o cuda-lib1.o library1.o
    gcc  -shared -Wl,-soname,library1.so -o library1.so library1.o uda-lib1.o
    nvcc -arch=sm_30 --compiler-options '-fPIC' -dc library2.cu
    nvcc -arch=sm_30 --compiler-options '-fPIC' -dlink -o cuda-lib2.o library2.o -lrary1
    gcc  -shared -Wl,-soname,library2.so -o library2.so library2.o cuda-lib2.o
    gcc  -c mycode.c
    gcc  -o mycode -L. -lrary2 -lrary1 mycode.o

working:
    rm -f *.o *.so
    nvcc -arch=sm_30 --compiler-options '-fPIC' -dc library1.cu
    nvcc -arch=sm_30 --compiler-options '-fPIC' -dc library2.cu
    nvcc -arch=sm_30 --compiler-options '-fPIC' -dlink -o cuda-lib.o library1.o library2.o
    gcc  -shared -Wl,-soname,library.so -o library.so library1.o library2.o cuda-lib.o
    gcc  -c -fPIC mycode.c                                                      
    gcc  -o mycode -L. -lrary  -L/usr/local/cuda/lib64 -lcuda -lcudart mycode.o

make working 没有任何问题。但它并没有形成一个图书馆链。 library1.cu 和 library2.cu 在同一个 .so 文件中。

make broken 失败

nvcc -arch=sm_30 --compiler-options '-fPIC' -dlink -o cuda-lib2.o library2.o -lrary1
nvlink error   : Undefined reference to '_Z10func1_lib1v' in 'library2.o'

如果我通过 nm 检查 library1.so，则目标 (T) 内有 _Z10func1_lib1v。

Answer 1

在您的 "broken" 方法中，您试图创建一个仅包含 __device__ 函数的 library1.so（共享库）：

__device__ void func1_lib1(void){}

任何其他希望使用该 __device__ 功能的对象都必须使用 relocatable device code/separate compilation and linking，您当然会尝试这样做。

但是，关于库，设备链接 only supports functions contained in static libraries。请注意 nvcc 手册中的这些语句：

The device linker has the ability to read the static host library formats (.a on Linux and Mac OS X, .lib on Windows). It ignores any dynamic (.so or .dll) libraries.

和：

Note that only static libraries are supported by the device linker.

所以你的一般策略是行不通的。一种可能的解决方法是将 library1.cu 代码放在静态库中：

rm -f *.o *.so
nvcc -arch=sm_30 --compiler-options '-fPIC' -dc library1.cu
nvcc -arch=sm_30 --lib -o cuda-lib1.a library1.o
nvcc -arch=sm_30 --compiler-options '-fPIC' -dc library2.cu
nvcc -arch=sm_30 --compiler-options '-fPIC' -dlink -o cuda-lib2.o library2.o cuda-lib1.a
gcc  -shared -Wl,-soname,library2.so -o library2.so -L/usr/local/cuda/lib64 -lcuda -lcudart library2.o cuda-lib2.o cuda-lib1.a
gcc  -c mycode.c
gcc  -o mycode -L. -lrary2  mycode.o

或者创建一系列 .so 库，这些库不需要跨库边界的设备链接，这或多或少在您的 "working" 案例中得到了证明。

CUDA nvcc 构建库链

CUDA nvcc building chain of libraries

cuda

dynamic-linking

nvcc