CUDA图书馆建设链

Question

我的目标是：library2.so使用library1.so，mycode.o使用(libs应该链接) library2.so (可能还有library1.so)。

源代码是(省略一行头文件)：

图书馆1.cu：

__device__ void func1_lib1(void){}

图书馆2.cu：

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

我的代码c：

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

我正在用Makefile构建共享库根据

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'

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

Answer 1

在您的“破坏”方法中，您试图创建一个library1.so (共享库)，它只包含一个__device__函数：

__device__ void func1_lib1(void){}

任何其他希望使用该__device__函数的对象都必须使用可重定位的设备代码/单独编译和链接，这当然是您正在尝试的。

但是，关于库，设备连接库。请注意nvcc手册中的以下语句：

设备链接器能够读取静态主机库格式(.a在Linux上，Mac，.lib在Windows上)。它忽略任何动态(.so或.dll)库。

以及：

注意，设备链接器只支持静态库。

这样你的总体策略就行不通了。一个可能的解决方法是将library 1.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库序列，这在您的“工作”案例中或多或少地得到了演示。