为啥即使使用 -cudart static 编译,库用户仍然需要链接到 cuda 运行时
Posted
技术标签:
【中文标题】为啥即使使用 -cudart static 编译,库用户仍然需要链接到 cuda 运行时【英文标题】:Why is linking to cuda runtime still necessary for library user even when compiling with -cudart static为什么即使使用 -cudart static 编译,库用户仍然需要链接到 cuda 运行时 【发布时间】:2020-07-12 11:16:10 【问题描述】:我有一些简单的 cuda 代码,我正在使用 nvcc
编译到静态库,以及一些我正在使用 g++
编译并链接到先前编译的静态库的用户代码。尝试链接时,即使我在nvcc
编译命令行中使用了-cudart static
选项,我也会收到cudaMalloc
之类的链接器错误。
这是我的代码:
//kern.hpp
#include <cstddef>
class Kern
private:
float* d_data;
size_t size;
public:
Kern(size_t s);
~Kern();
void set_data(float *d);
;
//kern.cu
#include <iostream>
#include <kern.hpp>
__global__ void kern(float* data, size_t size)
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if(idx < size)
data[idx] = 0;
Kern::Kern(size_t s) : size(s)
cudaMalloc((void**)&d_data, size*sizeof(float));
Kern::~Kern()
cudaFree(d_data);
void Kern::set_data(float* d)
size_t grid_size = size;
std::cout << "Starting kernel with grid size " << grid_size << " and block size " << 1 <<
std::endl;
kern<<<grid_size, 1>>>(d_data, size);
cudaError_t err = cudaGetLastError();
if(err != cudaSuccess)
std::cout << "ERROR: " << cudaGetErrorString(err) << std::endl;
cudaDeviceSynchronize();
cudaMemcpy((void*)d, (void*)d_data, size*sizeof(float), cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
//main.cpp
#include <iostream>
#include <kern.hpp>
int main(int argc, char** argv)
std::cout << "starting" << std::endl;
Kern k(256);
float arr[256];
k.set_data(arr);
bool ok = true;
for(int i = 0; i < 256; ++i) ok &= arr[i] == 0;
std::cout << (ok ? "done" : "wrong") << std::endl;
我正在用nvcc
编译紧缩如下:
nvcc -I ./ -lib --compiler-options '-fPIC' -o libkern.a kern.cu -cudart static
然后用g++
main 如下:
g++ -o main main.cpp -I ./ -L. -L/opt/cuda/lib64 -lkern
产生错误的原因:
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `Kern::Kern(unsigned long)':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x4d): undefined reference to `cudaMalloc'
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `Kern::~Kern()':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x6b): undefined reference to `cudaFree'
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `Kern::set_data(float*)':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x152): undefined reference to `__cudaPushCallConfiguration'
/usr/bin/ld: tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x175): undefined reference to `cudaGetLastError'
/usr/bin/ld: tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x1a1): undefined reference to `cudaGetErrorString'
/usr/bin/ld: tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x1c6): undefined reference to `cudaDeviceSynchronize'
/usr/bin/ld: tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x1ee): undefined reference to `cudaMemcpy'
/usr/bin/ld: tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x1f3): undefined reference to `cudaDeviceSynchronize'
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `__cudaUnregisterBinaryUtil()':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x24e): undefined reference to `__cudaUnregisterFatBinary'
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `__nv_init_managed_rt_with_module(void**)':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x269): undefined reference to `__cudaInitModule'
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `__device_stub__Z4kernPfm(float*, unsigned long)':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x305): undefined reference to `__cudaPopCallConfiguration'
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `__nv_cudaEntityRegisterCallback(void**)':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x430): undefined reference to `__cudaRegisterFunction'
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `__sti____cudaRegisterAll()':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x44b): undefined reference to `__cudaRegisterFatBinary'
/usr/bin/ld: tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x47c): undefined reference to `__cudaRegisterFatBinaryEnd'
/usr/bin/ld: ./libkern.a(tmpxft_00001d30_00000000-8_kern.o): in function `cudaError cudaLaunchKernel<char>(char const*, dim3, dim3, void**, unsigned long, CUstream_st*)':
tmpxft_00001d30_00000000-5_kern.cudafe1.cpp:(.text+0x4d9): undefined reference to `cudaLaunchKernel'
collect2: error: ld returned 1 exit status
但如果我执行以下操作:
g++ -o main main.cpp -I ./ -L. -L/opt/cuda/lib64 -lkern -lcudart
一切正常。
我的问题是,既然我在nvcc
编译行中有一个-cudart static
,那么libkern.a
不应该已经解决了cuda 运行时的符号吗?为什么-lcudart
仍然需要在g++
行中?
另外,如果我将 libkern.a
更改为共享对象,则在 g++
行中不链接到 cuda 运行时有效。也就是说,以下工作:
nvcc -I ./ -shared --compiler-options '-fPIC' -o libkern.so kern.cu -cudart static
g++ -o main main.cpp -I ./ -L. -L/opt/cuda/lib64 -lkern
为什么静态库版本失败,但共享对象版本有效?
请注意,在 nvcc
行中将 -cudart static
替换为 -lcudart_static
后,我已经尝试了上述方案,并且进行该替换后行为没有任何变化。这是意料之中的,因为这两个选项基本上做同样的事情,对吗?
我在linux上。
nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2019 NVIDIA Corporation
Built on Wed_Oct_23_19:24:38_PDT_2019
Cuda compilation tools, release 10.2, V10.2.89
g++ --version
g++ (GCC) 10.1.0
Copyright (C) 2020 Free Software Foundation, Inc.
This is free software; see the source for copying conditions. There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
非常感谢任何帮助和/或澄清。
【问题讨论】:
你从来没有在任何地方链接过运行时库,为什么会这样呢? 是的,我的错,我完全误解了文档。 【参考方案1】:如果您研究the nvcc documentation,很明显-lib
选项创建一个静态库(并指定不链接),而-shared
选项创建一个共享库并指定链接。例如,摘录:
4.2.2.1。 --链接(-链接) 指定默认行为:编译并链接所有输入文件。
4.2.2.2。 --lib (-lib) 如有必要,将所有输入文件编译为目标文件,并将结果添加到指定的库输出文件中。
4.2.3.11。 --共享(-共享) 在链接期间生成一个共享库。 当需要其他链接器选项进行更多控制时,请使用选项 --linker-options。
我相信这或多或少与典型的 gcc/g++ 用法一致。如果你在“g++ create static library”上进行谷歌搜索,你会得到任意数量的references,这表明你基本上应该这样做:
g++ -c my_source_file.cpp ...
ar ...
换句话说,指定了源到对象的编译,但没有指定链接。举一个例子,cudaMalloc
是 CUDA 运行时库的一部分,与它的连接将在链接阶段完成。
nvcc
是一个相当复杂的动物,但我们应该记住,对于某些功能,它主要使用已安装的主机工具链。这包括主机代码的编译,也包括最后的链接阶段。
结合这一点,我相信您在这里想要做的是“部分”链接或增量链接。在最终链接阶段之前执行一些最终链接阶段。
GNU 链接器(同样,nvcc
将在后台使用,默认情况下在 linux 上使用)supports that,所以如果我们不考虑编译可重定位设备代码的任何问题,应该可以这样做你想要的如下:
$ nvcc -Xcompiler '-fPIC' -I. -c kern.cu
$ ld -o kern.ro -r kern.o -L/usr/local/cuda/lib64 -lcudart_static -lculibos
$ ar rs libkern.a kern.ro
ar: creating libkern.a
$ g++ -o main main.cpp -I ./ -L. -lkern -lpthread -lrt -ldl
$ cuda-memcheck ./main
========= CUDA-MEMCHECK
starting
Starting kernel with grid size 256 and block size 1
done
========= ERROR SUMMARY: 0 errors
$
注意事项:
-lpthread -lrt -ldl
是 cudart/culibos 的标准库依赖项,因此需要在最终链接阶段提供这些依赖项,但它们不依赖于任何 CUDA 工具包项。如果您希望这些依赖项也从增量链接对象中删除,我认为这是一个单独的问题,与 CUDA 无关。
归档步骤(创建库)对于这个简单的案例来说不是必需的。我们可以直接将增量链接的 (-r
) 对象 kern.ro
传递到最终的编译/链接步骤。
请注意,您的 CUDA 安装显然位于不同的位置,因此可能需要更改上述某些库路径 (-L
)。
【讨论】:
我想我完全误解了文档,非常抱歉。它现在可以正常编译和链接。非常感谢您,您的回答为我澄清了很多事情。以上是关于为啥即使使用 -cudart static 编译,库用户仍然需要链接到 cuda 运行时的主要内容,如果未能解决你的问题,请参考以下文章
为啥我不能在编译时将整数添加到泛型集合中,即使使用引用类型作为数字创建的泛型? [复制]
std::any 用于仅移动模板,其中 copy-ctor 内的 static_assert 等于编译错误,但为啥呢?
为啥即使模板访问私有属性,Angular AOT 也会编译?
为啥即使实现了 Iterable,我也会收到 foreach 编译器错误?