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

Question

我正在尝试在我的项目中使用 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() 设备函数。

这是我的 build/test 脚本：

#!/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 个共享库会出现问题？我读过“设备链接器忽略共享库”，但在这种情况下，它是主机链接器创建共享库，而不是设备链接器，所以我希望这没问题？

谢谢！

Answer 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 命令使其成为每个库中的私有符号，因此不再有冲突。

感谢您的帮助！

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

CUDA separable compilation + shared libraries -> Invalid device function / segfault

c++

cuda

segmentation-fault

nvcc