c++ - 如何编译依赖于 C++ 目标文件的 CUDA 共享库？

Question

我正在编写一个包含多个内核调用的 c++/CUDA 库。

编辑:我觉得原来的帖子有点长，所以我创建了一个更好的例子。原帖如下。

这是简化为最小示例的项目。它不会编译，并给出以下错误:

nvcc -Xcompiler -fPIC -x cu -c -dc -o myclass.o myclass.cpp
nvcc -Xcompiler -fPIC --lib myclass.o kernel.cu -o libhelpme.a -I.
ptxas fatal   : Unresolved extern function '_ZN7myclassC1Ei'
makefile:8: recipe for target 'lib' failed
make: *** [lib] Error 255

关于此主题的所有文档都指向编译可执行文件或目标文件；我不想做这些，而是​​专门做一个静态库。我该怎么做？

代码:

生成文件

program: class lib
    nvcc -o program main.cc -I. -L. -lhelpme

class:
    nvcc -Xcompiler -fPIC -x cu -c -dc -o myclass.o myclass.cpp

lib: class
    nvcc -Xcompiler -fPIC --lib myclass.o kernel.cu -o libhelpme.a -I.

clean:
    rm *.o *.a program

main.cc

#include "stdio.h"
#include <iostream>
#include "kernel.h"
int main()
{
    std::cout << "hello world" << std::endl;
    wrapper();
    return 0;
}

myclass.h

#ifdef __CUDACC__
#define COMMON __host__ __device__
#else
#define COMMON
#endif
#ifndef M
#define M
class myclass
{
    public:
        int x;
        COMMON myclass(int y);
        COMMON void increment();
};
#endif

myclass.cpp

#include "myclass.h"
#ifdef __CUDACC__
#define COMMON __host__ __device__
#else
#define COMMON
#endif
COMMON myclass::myclass(int y)
{
    x = y;
}
COMMON void myclass::increment()
{
    x += 1;
}

内核.h

extern void wrapper();

内核.cu

#include "stdio.h"
#include <iostream>
#include "myclass.h"

class myotherclass
{
    public:
        int x;
        COMMON myotherclass(int y) {x = y;}
        COMMON void decrement() {x -= 1;}
};

__global__ void dokernel()
{
    myotherclass p(8); //This compiles just fine.
    myclass q(7); //This will not compile
}

void wrapper()
{
    std::cout << "hello from wrapper\n";
    myclass q(1);
    myotherclass s(4);
    std::cout << "x = " << s.x << "\n";
    s.decrement();
    std::cout << "x = " << s.x << "\n";
    dokernel<<<1,1>>>();

}

我慢慢地确信这是不可能的......

原始帖子:我有许多 C++ 源/头文件，例如vec.cpp 和 vec.h 被编译成目标文件，例如vec.o

这是一个例子:vec.h

class vec
{
    public:
        realnum x,y,z;
        __host__ __device__ vec(float _x, float _y, float _z);
}

vec.cpp

__host__ __device__ vec::vec(float _x, float _y, float _z) {x = _x; y = _y; z = _z;}

这是我的完整 makefile(仍在制作中):

CC=nvcc
CFLAGS = -Wall -g -O3

HOME_DIR = $(shell pwd)

SRC_DIR := ${HOME_DIR}/../src
OBJ_DIR := ${HOME_DIR}/../lib
LIB_DIR := ${HOME_DIR}/../lib
KER_DIR := ${HOME_DIR}/../kernel
SRC_FILES := $(wildcard $(SRC_DIR)/*.cpp)
OBJ_FILES := $(patsubst $(SRC_DIR)/%.cpp,$(OBJ_DIR)/%.o,$(SRC_FILES))

LPROPS := -L${LIB_DIR} -lcuprops
LMAIN := -L${LIB_DIR} -lsharc
LRDR := -L${LIB_DIR} -lcurdr

INCL_PROPS := -I${SRC_DIR} -I${KER_DIR}

program: $(LIB_DIR)/libcurdr.so ${LIB_DIR}/libsharc.so $(LIB_DIR)/libcuprops.so $(OBJ_FILES)
    ${CC} -o $@ main.cc -I${SRC_DIR} ${LPROPS} ${LMAIN} ${LRDR}

${LIB_DIR}/libsharc.so: $(OBJ_FILES) $(LIB_DIR)/libcuprops.so
    ${CC} -Xcompiler -fPIC --shared  ${OBJ_FILES} -o $(LIB_DIR)/libsharc.so ${INCL_PROPS}

$(OBJ_DIR)/%.o: $(SRC_DIR)/%.cpp
    ${CC} -Xcompiler -fPIC -dc -o $@ $< ${INCL_PROPS}

$(LIB_DIR)/libcuprops.so:
    ${CC} -Xcompiler -fPIC --shared -o $(LIB_DIR)/libcuprops.so ${KER_DIR}/nvidia_properties.cu ${INCL_PROPS}

$(LIB_DIR)/libcurdr.so: $(OBJ_FILES)
    ${CC} -Xcompiler -fPIC --shared ${OBJ_FILES} ${KER_DIR}/gpu_rdr.cu -o $(LIB_DIR)/libcurdr.so ${INCL_PROPS}

clean:
    rm ${LIB_DIR}/*

当我make时，我得到以下信息:

ptxas fatal   : Unresolved extern function '_ZN3vecC1Eddd'

我有一个内核，我在其中尝试初始化一个 vector :

__global__ void SOME_KERNEL()
{
    int row = blockIdx.y*blockDim.y + threadIdx.y;
    int col = blockIdx.x*blockDim.x + threadIdx.x;
    if (row < dev_height && col < dev_width)
    {
        vec t(0,0,0); //Compiles nicely when I comment out this line!
    }
}

我读过 separate compiling and linking据称典型的项目架构(我相信我正在使用)通过以下方式与单独的编译和链接兼容:

objects = main.o particle.o v3.o

all: $(objects)
    nvcc -arch=sm_20 $(objects) -o app

%.o: %.cpp
    nvcc -x cu -arch=sm_20 -I. -dc $< -o $@

clean:
    rm -f *.o app

注意“-dc”标志的使用，它与this一致回答。

在这一点上，我已经尝试了很多东西，以至于我完全迷失了。那么，我该如何编译这个项目呢？

在有用的情况下，这里是 make 的完整输出:

nvcc -Xcompiler -fPIC -dc -o /home/wvn/dirs/projects/sharc/build/../lib/mat33.o /home/wvn/dirs/projects/sharc/build/../src/mat33.cpp -I/home/wvn/dirs/projects/sharc/build/../src -I/home/wvn/dirs/projects/sharc/build/../kernel
nvcc -Xcompiler -fPIC -dc -o /home/wvn/dirs/projects/sharc/build/../lib/vec.o /home/wvn/dirs/projects/sharc/build/../src/vec.cpp -I/home/wvn/dirs/projects/sharc/build/../src -I/home/wvn/dirs/projects/sharc/build/../kernel
nvcc -Xcompiler -fPIC -dc -o /home/wvn/dirs/projects/sharc/build/../lib/sharc.o /home/wvn/dirs/projects/sharc/build/../src/sharc.cpp -I/home/wvn/dirs/projects/sharc/build/../src -I/home/wvn/dirs/projects/sharc/build/../kernel
nvcc -Xcompiler -fPIC -dc -o /home/wvn/dirs/projects/sharc/build/../lib/boundingbox.o /home/wvn/dirs/projects/sharc/build/../src/boundingbox.cpp -I/home/wvn/dirs/projects/sharc/build/../src -I/home/wvn/dirs/projects/sharc/build/../kernel
nvcc -Xcompiler -fPIC --shared /home/wvn/dirs/projects/sharc/build/../lib/mat33.o /home/wvn/dirs/projects/sharc/build/../lib/vec.o /home/wvn/dirs/projects/sharc/build/../lib/sharc.o /home/wvn/dirs/projects/sharc/build/../lib/boundingbox.o /home/wvn/dirs/projects/sharc/build/../kernel/gpu_rdr.cu -o /home/wvn/dirs/projects/sharc/build/../lib/libcurdr.so -I/home/wvn/dirs/projects/sharc/build/../src -I/home/wvn/dirs/projects/sharc/build/../kernel
ptxas fatal   : Unresolved extern function '_ZN3vecC1Eddd'
makefile:32: recipe for target '/home/wvn/dirs/projects/sharc/build/../lib/libcurdr.so' failed
make: *** [/home/wvn/dirs/projects/sharc/build/../lib/libcurdr.so] Error 255

Answer 1

指的是更新后的帖子，而不是原始帖子。

如评论中所述，当您编译以 .cpp 结尾但包含 CUDA 构造或设备的文件时，需要做的事情之一是添加 -x cu代码。您已经将其添加到适当的位置。

您围绕 COMMON 宏的讨论与此无关。它们有不同的用途。该宏不能替代 -x cu。

另一件遗漏的事情是您需要在静态库生成点指示 nvcc 需要可重定位设备代码和设备链接。您可以通过将 -rdc=true 添加到编译命令行来执行此操作。您的 Makefile 可以这样修改:

lib: class
    nvcc -Xcompiler -fPIC -rdc=true --lib myclass.o kernel.cu -o libhelpme.a -I.

这是必需的，因为您在一个编译单元 (kernel.cu) 中有设备代码，它正在调用另一个编译单元 (myclass.cpp) 中的设备代码。

有了这个改变，你更新后的帖子/项目对我来说编译没有问题，运行也没有错误。

$ make clean
rm *.o *.a program
$ make
nvcc -Xcompiler -fPIC -x cu -c -dc -o myclass.o myclass.cpp
nvcc -Xcompiler -fPIC -rdc=true --lib myclass.o kernel.cu -o libhelpme.a -I.
nvcc -o program main.cc -I. -L. -lhelpme
$ cuda-memcheck ./program
========= CUDA-MEMCHECK
hello world
hello from wrapper
x = 4
x = 3
========= ERROR SUMMARY: 0 errors
$