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 构建共享库

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手册中的这些陈述:

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

和:

请注意,设备链接器仅支持静态库.

所以你的总体战略是行不通的.可能的解决方法是将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库,这些库不需要跨库边界的设备链接,这在您的"工作"案例中或多或少地得到了证明.