Nvlink、可重定位设备代码和静态设备库

Question

在调查可重定位设备代码的一些问题时，我遇到了一些我不太理解的东西。

这是幻灯片6上所示内容的一个用例。我使用Robert Crovella的答案作为复制代码的基础。想法是我们有一些编译成静态库（例如某些数学/工具箱库）的可重定位设备代码，并且我们想要在程序的另一个设备库中使用该预编译库的某些函数：

libutil.a ---> libtest.so ---> test_pgm

假设这个外部库包含以下函数：

__device__ int my_square (int a);

libutil.a 例如是在另一个项目中按照以下方式生成的：

nvcc ${NVCC_FLAGS} -dc util.cu
nvcc ${NVCC_FLAGS} -dlink util.o -o util_dlink.o
nvcc ${NVCC_FLAGS} -lib util_dlink.o util.o -o libutil.a

然后，在我们的项目中，生成 libtest.so：

nvcc ${NVCC_FLAGS} -dc test.cu
nvcc ${NVCC_FLAGS} -dlink test.o libutil.a -o test_dlink.o
g++ -shared -Wl,-soname,libtest.so -o libtest.so test.o test_dlink.o libutil.a -L${CUDA_LIBDIR} -lcudart

但是当我生成test_dlink.o时，会出现以下错误：

nvlink error   : Undefined reference to '_Z9my_squarei' in 'test.o'

链接器找不到我们的虚拟函数my_square(int)。如果我们使用（假设我们可以访问util.o）：

nvcc ${NVCC_FLAGS} -dlink test.o util.o -o test_dlink.o

链接器成功了，之后一切都正常工作。

进一步调查：

$ nm -C libutil.a

util_dlink.o:
                 U atexit
                 U __cudaRegisterFatBinary
0000000000000015 T __cudaRegisterLinkedBinary_39_tmpxft_0000106a_00000000_6_util_cpp1_ii_843d693d
  ...

util.o:
                 U __cudaInitModule
                 U __cudaRegisterLinkedBinary_39_tmpxft_0000106a_00000000_6_util_cpp1_ii_843d693d
  ...
0000000000000015 T my_square(int)
  ...

符号在存档的util.o中存在，但是nvlink（由nvcc调用）似乎找不到它。为什么会这样？根据官方文档的说法：

设备链接器有能力读取静态主机库格式（Linux和Mac上的.a，Windows上的.lib）。

当然，我们可以提取对象文件并与其链接：

ar x libutil.a `ar t libutil.a | grep -v "dlink"`
nvcc ${NVCC_FLAGS} -dlink test.o util.o -o test_dlink.o

但是这不像是期望的解决方案...那么我错过了什么？还有其他nvcc选项可以解决这个问题吗？在生成libutil.a和/或libtest.so时是否存在错误？

请注意，此测试是在Arch Linux上使用CUDA 6.5进行的。

编辑：使用带注释的代码修复

Makefile

NVCC_FLAGS=-m64 -arch=sm_20 -Xcompiler '-fPIC'
CUDA_LIBDIR=${CUDA_HOME}/lib64

testmain : main.cpp libtest.so
    g++ -c main.cpp
    g++ -o testmain -L. -ldl -Wl,-rpath,. -ltest -L${CUDA_LIBDIR} -lcudart main.o

libutil.a : util.cu util.cuh
    nvcc ${NVCC_FLAGS} -dc util.cu
    # ---> FOLLOWING LINES THAT WERE WRONG <---
    # nvcc ${NVCC_FLAGS} -dlink util.o -o util_dlink.o
    # nvcc ${NVCC_FLAGS} -lib util.o util_dlink.o -o libutil.a
    # INSTEAD:
    nvcc ${NVCC_FLAGS} -lib util.o -o libutil.a
    # Assuming util is an external library, so util.o is not available
    rm util.o

libtest.so : test.cu test.h libutil.a util.cuh
    nvcc ${NVCC_FLAGS} -dc test.cu
    # Use NVCC for device linking + G++
    nvcc -v ${NVCC_FLAGS} -dlink test.o libutil.a -o test_dlink.o
    g++ -shared -o libtest.so test.o test_dlink.o libutil.a -L${CUDA_LIBDIR} -lcudart
    # Or let NVCC generate the shared library
    #nvcc -v ${NVCC_FLAGS} -shared -L. -lutil test.o -o libtest.so

clean :
    rm -f testmain *.o *.a *.so

test.h

#ifndef TEST_H
# define TEST_H

int my_test_func();

#endif //! TEST_H

test.cu

#include <stdio.h>

#include "test.h"
#include "util.cuh"

#define DSIZE 1024
#define DVAL 10
#define SQVAL 3
#define nTPB 256

#define cudaCheckErrors(msg)                             \
  do {                                                   \
    cudaError_t __err = cudaGetLastError();              \
    if (__err != cudaSuccess) {                          \
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
              msg, cudaGetErrorString(__err),            \
              __FILE__, __LINE__);                       \
      fprintf(stderr, "*** FAILED - ABORTING\n");        \
      exit(1);                                           \
    }                                                    \
  } while (0)

__global__ void my_kernel(int *data){
  int idx = threadIdx.x + (blockDim.x *blockIdx.x);
  if (idx < DSIZE) data[idx] =+ DVAL + my_square (SQVAL);
}

int my_test_func()
{
  int *d_data, *h_data;
  h_data = (int *) malloc(DSIZE * sizeof(int));
  if (h_data == 0) {printf("malloc fail\n"); exit(1);}
  cudaMalloc((void **)&d_data, DSIZE * sizeof(int));
  cudaCheckErrors("cudaMalloc fail");
  for (int i = 0; i < DSIZE; i++) h_data[i] = 0;
  cudaMemcpy(d_data, h_data, DSIZE * sizeof(int), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy fail");
  my_kernel<<<((DSIZE+nTPB-1)/nTPB), nTPB>>>(d_data);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel");
  cudaMemcpy(h_data, d_data, DSIZE * sizeof(int), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy 2");
  for (int i = 0; i < DSIZE; i++)
    if (h_data[i] != DVAL + SQVAL*SQVAL)
    {
      printf("Results check failed at offset %d, data was: %d, should be %d\n",
             i, h_data[i], DVAL);
      exit(1);
    }
  printf("Results check passed!\n");
  return 0;
}

util.cuh

#ifndef UTIL_CUH
# define UTIL_CUH

__device__ int my_square (int a);

#endif //! UTIL_CUH

util.cu

#include "util.cuh"

__device__ int my_square (int a)
{
  return a * a;
}

main.cpp

#include "test.h"

int main()
{
  my_test_func();
  return 0;
}

- BenC

---

这似乎是一个简单的名称混淆问题。请注意，链接器正在查找一个混淆的C++符号名称，而对象包含一个未混淆的C风格名称。这里没有足够的信息告诉你如何/什么去修复。 - talonmies

Answer 1

我建议在问题中放置一个完整的简单示例，就像我下面所做的那样。不提倡使用外部代码链接。当它们变得陈旧时，问题的价值就会降低。

是的，你在生成libutil.a时出现了错误。创建具有公开设备链接的静态库与创建没有公开设备链接（按定义）的共享库不同。请注意我在你链接的先前问题中提到“无CUDA包装器”的提及。本问题中的示例具有公开的设备链接，因为my_square在库中但被用于库外的代码中。

请查看nvcc relocatable device code compiling examples，您将找到一个生成可链接设备的静态库的示例。在静态库创建过程中没有设备链接步骤。设备链接步骤在最终可执行文件创建时进行（或在本例中，在so的创建时，即“CUDA边界”）。静态库创建中的“额外”设备链接操作是您观察到的错误的直接原因。

这里是一个完整的示例：

$ cat util.h

__device__ float my_square(float);

$ cat util.cu

__device__ float my_square(float val){ return val*val;}

$ cat test.h

float dbl_sq(float val);

$ cat test.cu
#include "util.h"

__global__ void my_dbl_sq(float *val){
  *val = 2*my_square(*val);
}

float dbl_sq(float val){
  float *d_val, h_val;
  cudaMalloc(&d_val, sizeof(float));
  h_val = val;
  cudaMemcpy(d_val, &h_val, sizeof(float), cudaMemcpyHostToDevice);
  my_dbl_sq<<<1,1>>>(d_val);
  cudaMemcpy(&h_val, d_val, sizeof(float), cudaMemcpyDeviceToHost);
  return h_val;
}
$ cat main.cpp
#include <stdio.h>
#include "test.h"

int main(){

  printf("%f\n", dbl_sq(2.0f));
  return 0;
}
$ nvcc -arch=sm_35 -Xcompiler -fPIC -dc util.cu
$ nvcc -arch=sm_35 -Xcompiler -fPIC -lib util.o -o libutil.a
$ nvcc -arch=sm_35 -Xcompiler -fPIC -dc test.cu
$ nvcc -arch=sm_35 -shared -Xcompiler -fPIC -L. -lutil test.o -o libtest.so
$ g++ -o main main.cpp libtest.so
$ cuda-memcheck ./main
========= CUDA-MEMCHECK
8.000000
========= ERROR SUMMARY: 0 errors
$

在这个例子中，设备链接会自动发生在用于创建 .so 库的 nvcc 调用中。在我的这个例子中，我已经将我的 LD_LIBRARY_PATH 环境变量设置为包含我的工作目录。在 CentOS 6.2 上使用 CUDA 6.5 进行测试。（请注意，在创建可执行文件期间可以执行多个设备链接操作，但这些设备链接操作必须位于不同的链接域内，即用户代码或用户代码入口点不能在域之间共享。在这里不是这种情况。）

- Robert Crovella

---

感谢指出在生成 libutil.a 时设备链接是无用的，需要在与其链接时进行。简单地删除该步骤即可解决问题。此后，在那个虚拟示例中我不需要依赖于 nvcc，但是 nvcc 是否做了一些使得总是使用它来生成类似的 libtest.so 的操作呢？nvcc 生成了一个调用 g++ 的命令，并带有一些额外的链接选项（-lcudadevrt -lcudart_static -lrt -lpthread -ldl），这些选项可能并不总是必需的。请注意，我正在调查一个基于 CMake 的复杂 CUDA 构建链中的奇怪错误，因此这实际上很重要。 - BenC

我们可能需要深入研究您的情况（或者我的情况）。在我的情况下，我遇到了使用g++创建的libtest.so的问题，当我稍后尝试链接它时，仍然存在一些未解析的符号（与您问题中的符号无关）。当我使用nvcc -shared来创建so时，这些问题消失了。正如您指出的那样，如果跳过静态库，则可以很好地创建so。因此，普通对象+静态库+设备链接成为so的组合意味着我必须使用nvcc来创建so，但我没有进一步研究它。 - Robert Crovella

“-Xcompiler -fPIC” 不是多余的吗？也就是说，即使使用了“-dc”或“-rdc”，我们仍然需要它吗？ - einpoklum

"-rdc" 指导设备代码生成行为。 "-Xcompiler" 是对主机编译器的指示。它们不是冗余的。 - Robert Crovella