CUDA nvcc构建图书馆链

Question

我的目标是：library2.so正在使用library1.so和mycode。o正在使用（应该链接库）library2.so（可能还有library1.so

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

library1.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

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

使破坏失败

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

在您的“损坏”方法中，您正在尝试创建一个仅包含__device__函数的 library1.so（共享库）：

__device__ void func1_lib1(void){}

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

然而，对于库，设备链接仅支持静态库中包含的函数。请注意nvcc手册中的以下语句：

设备链接器能够读取静态主机库格式（Linux和Mac OS X上的.a，Windows上的.lib）。它忽略任何动态（.so或.dll）库。

和：

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

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

rm -f *.o *.so
nvcc -arch=sm_30 --compiler-html" target="_blank">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

或者创建一个< code >序列。所以库不需要跨库边界的设备链接，这或多或少在您的“工作”案例中得到了证明。