CUDA可分离编译+共享库->无效的设备功能/段错误

Posted 2023-03-23

【中文标题】CUDA可分离编译+共享库->无效的设备功能/段错误

【英文标题】：CUDA separable compilation + shared libraries -> Invalid device function / segfault

【发布时间】：2021-11-02 19:16:33

【问题描述】：

我正在尝试在我的项目中使用 CUDA 可分离编译。该项目由依赖于几个共享库（全部构建在同一个构建系统中）的二进制文件组成。这些共享库依次使用通用 CUDA 代码。运行二进制文件时，我得到一个类似于here 的段错误。当我创建一个最小示例时，我得到“无效的设备功能”错误。 如果我将共享库转换为静态库，错误就会消失。不幸的是，我无法控制它，需要让它与共享库一起使用。

我在 SO 中看到过一些类似的帖子，但他们使用 CMake，并且解决方案通常涉及将库从共享更改为静态，这在我的项目中是做不到的。我已经仔细检查了我在正确的 GPU 上运行代码（如果我进行一些更改，它确实可以工作，见下文），所以这不是问题。

我相信我在进行 CUDA 可分离编译、设备链接或创建共享库时遗漏了一些东西。

以下是该问题的完全可重现的最小示例：

// common.h
#ifndef COMMON_H
#define COMMON_H

__device__ int common();

#endif

// common.cu
#include "common.h"

__device__ int common()

    return 123;

// a.h
#ifndef A_H
#define A_H

__attribute__((__visibility__("default")))
void runA();

#endif

// a.cu
#include "a.h"

#include <cstdio>
#include <iostream>

#include "common.h"

__global__ void kernelA()

    printf("Running A: %d\n", 456 + common());


void runA()

    kernelA<<<1,1>>>();
    std::cout << cudaGetErrorString(cudaPeekAtLastError()) << std::endl;
    cudaDeviceSynchronize();

// b.h
#ifndef B_H
#define B_H

__attribute__((__visibility__("default")))
void runB();

#endif

// b.cu
#include "b.h"

#include <cstdio>
#include <iostream>

#include "common.h"

__global__ void kernelB()

    printf("Running B: %d\n", 321 + common());


void runB()

    kernelB<<<1,1>>>();
    std::cout << cudaGetErrorString(cudaPeekAtLastError()) << std::endl;
    cudaDeviceSynchronize();

// main.cpp
#include "a.h"
#include "b.h"

int main()

    runA();
    runB();

所以基本上是一个依赖于 2 个共享库 A 和 B 的二进制文件，它们都利用了common() 设备函数。

这是我的构建/测试脚本：

#!/usr/bin/env bash
set -euxo pipefail

CUDA_ROOT=/usr/local/cuda-10.2
NVCC=$CUDA_ROOT/bin/nvcc
CC=/usr/bin/g++
GENCODE="arch=compute_75,code=sm_75"

# Clean previous build
rm -f *.o *.so main

# Compile relocatable CUDA code
$NVCC -gencode=$GENCODE -dc -Xcompiler -fPIC,-fvisibility=hidden common.cu -o common.cu.o
$NVCC -gencode=$GENCODE -dc -Xcompiler -fPIC,-fvisibility=hidden      a.cu -o      a.cu.o
$NVCC -gencode=$GENCODE -dc -Xcompiler -fPIC,-fvisibility=hidden      b.cu -o      b.cu.o

# Build shared library A
$NVCC -gencode=$GENCODE -dlink common.cu.o a.cu.o -o a.dlink.o
$CC -shared common.cu.o a.cu.o a.dlink.o -L$CUDA_ROOT/lib64 -lcudart -o liba.so

# Build shared library B
$NVCC -gencode=$GENCODE -dlink common.cu.o b.cu.o -o b.dlink.o
$CC -shared common.cu.o b.cu.o b.dlink.o -L$CUDA_ROOT/lib64 -lcudart -o libb.so

# Build final executable
$CC main.cpp -L. -la -lb -o main

# Run it
LD_LIBRARY_PATH=. ./main

运行它我得到：

invalid device function
invalid device function

经过反复试验，我注意到问题已通过以下方式解决：

设备链接时不链接任一库中的common.cu.o（显然我需要使任一库不再使用common() 函数。 制作 A 和 B 静态库。 将 A 和 B 合并到一个共享库中。

很遗憾，我无法在我的项目中应用这些解决方案。为什么拥有 2 个共享库会成为问题？我读过“忽略共享库的设备链接器”，但在这种情况下，创建共享库的是主机链接器，而不是设备链接器，所以我希望没关系？

谢谢！

【问题讨论】：

没有共享 device 库之类的东西。您必须使用静态设备端链接，无论您是在构建时执行还是在运行时执行。这是没有商量余地的。

我没有创建共享设备库，库的入口点是调用 CUDA 代码的常规 C++ 函数。否则，如何实现 CuDNN、TensoRT 等？它们是内部调用 CUDA 代码的共享库。共享设备库将是一个 .so 文件，它提供了一些 __device__ 函数，可以在其他地方导入。这不是我在这里做的。我想在每个共享库中静态链接所有 CUDA 函数。

@user1011113：但实际上是这样，因为您在两个动态链接的主机库之间存在设备代码依赖性。那永远行不通。而这正是 CUDNN 和 CUBLAS 等库不做的事情

@talonmies 两个共享库都使用相同的公共代码，是的，但是它们彼此独立地静态链接。另外，我确保使用“-fvisibility=hidden”来防止导出符号——这不会导致每个库都拥有自己的“common ()”函数的静态副本吗？也许正如罗伯特所说，无论出于何种原因，这样的重复符号都是不可能的，就像运行时链接器不知道如何处理那样。我在Nvidia论坛上问过，让我们看看他们怎么说！谢谢！

玩了一会儿，我现在尝试只创建 1 个共享库，但仍然执行 2 个设备链接步骤（a.dlink.o，b.dlink.o）。如果我尝试将这两个链接到一个共享库中，我会得到： `__cudaRegisterLinkedBinary_41_tmpxft_00001022_00000000_6_common_cpp1_ii_e5ca4d49' 的多个定义 所以我想如果我分别链接这些库，这不会被注意到。所以问题一定是运行时链接器不知道如何处理这些重复的符号？在常规的 C++ 中，它只会选择一个，如果两者相同（就像这里一样），则无关紧要。

【参考方案1】：

与Nvidia接触后，终于找到了解决问题的办法：

将-Xcompiler -fvisibility=hidden 添加到每个dlink 命令

$NVCC -gencode=$GENCODE -dlink -Xcompiler -fvisibility=hidden common.cu.o a.cu.o -o a.dlink.o
$NVCC -gencode=$GENCODE -dlink -Xcompiler -fvisibility=hidden common.cu.o b.cu.o -o b.dlink.o

原来的问题是，事实上，liba.so 和 libb.so 都在公开导出相同的符号 (__cudaRegisterLinkedBinary_41_tmpxft_00000c46_00000000_6_common_cpp1_ii_e5ca4d49)，这当然会导致冲突。将-fvisibility=hidden 添加到dlink 命令使其成为每个库中的私有符号，因此不再有冲突。

感谢您的帮助！