CUDA构建共享库

Question

我需要创建一个针对cuda的共享库。编译库时没有问题，但是当我尝试在我的程序中使用它时，nvcc会返回链接器或ptxas错误。

我将问题简化为以下代码。该库必须替换不同的C函数（这里是memset）。 该库由三个C++文件组成：

FileA.h

#ifndef FILEA_H_
#define FILEA_H_

namespace A {
    __device__ 
    void* memset(void* _in, int _val, int _size);
};
#endif

FileA.cpp

#include "FileA.h"

__device__ 
void* A::memset(void* _in, int _val, int _size) {
    char* tmp = (char*)_in;
    for(int i = 0; i < _size; i++) tmp[i] = _val;
    return _in;
}

TempClass.h

#ifndef TEMPCLASS_H_
#define TEMPCLASS_H_

#include "FileA.h"

namespace A {
    template <typename T>
    class TC {
    public:
        __device__ 
        TC() {
            data = new T[10];
        }

        __device__ 
        ~TC(){
            delete [] data;
        }

        __device__ 
        void clear(){
            A::memset(data, 0, 10*sizeof(T));
        }

        T* data;
    };
};
#endif

以下是创建共享库的命令：

nvcc -Xcompiler -fPIC -x cu -rdc=true -c FileA.cpp -o FileA.o
nvcc -Xcompiler -fPIC --shared -o libTestA.so FileA.o -lcudart

这个库应该在主程序中使用：

main.cpp

#include <cuda.h>
#include <TempClass.h>
#include <iostream>

__device__
int doSomthing() {
    A::TC<int>* tc = new A::TC<int>();
    tc->clear();
    for (int i = 0; i < 5; i++) tc->data[i] = i;

    int sum = 0;
    for (int i = 0; i < 5; i++)  sum += tc->data[i];
    delete tc;
    return sum;
}

__global__
void kernel(int* _res) {
    _res[0] = doSomthing();
}

int main(int argc, char** argv) {
    int* devVar;
    int* hostVar;
    hostVar = new int[1];
    hostVar[0] = -1;
    cudaMalloc(&devVar, sizeof(int));
    cudaMemcpy(devVar, hostVar, sizeof(int), cudaMemcpyHostToDevice);

    kernel<<< 1, 1>>> (devVar);

    cudaMemcpy(hostVar, devVar, sizeof(int), cudaMemcpyDeviceToHost);

    std::cout << "kernel done. sum " << *hostVar << std::endl;

    return 0;
}

如果我尝试使用以下命令编译程序：

nvcc -Xcompiler -fPIC -I. -L. -rdc=true -x cu -c main.cpp -o main.o 
nvcc -Xcompiler -fPIC -I. -L. main.o -o main -lTestA

我收到了错误信息：

nvlink error   : Undefined reference to '_ZN1A6memsetEPvii' in 'main.o'

如果我试图直接编译文件，我会收到相同的错误提示：

nvcc -Xcompiler -fPIC -I. -L. -rdc=true -x cu main.cpp -o main -lTestA

命令nm libTestA.so显示该库包含函数符号_ZN1A6memsetEPvii。

当我在链接时删除-rdc=true选项时，我收到一个ptxas错误：

ptxas fatal   : Unresolved extern function '_ZN1A6memsetEPvii'

在我的情况下，静态链接不可行，我需要一个共享库。我也尝试将memset设为extern "C"函数，但这会与原始的C函数发生冲突。使用g++编译代码是正确的。你有什么建议来解决这个问题吗？

- Tim

---

1. FileA.cpp 应该改为 FileA.cu。
2. 你没有编译 FileA.cu，你怎么能指望链接成功呢？
3. 如果你“需要”一个共享库，为什么要静态链接呢？我不明白。

- einpoklum

1. 选项 -x cu 可让 nvcc 处理类似 cu 文件的 cpp 文件。
2. 我将其编译为共享库（请参见 main.cpp 前面的命令）。
3. 你在哪里看到任何静态链接？也许我漏掉了什么。

- Tim

1. 你说得对，nvcc会将其视为CUDA文件，但你仍然应该按其实际类型命名。
2. 啊，所以前面的两行出现在后面的两行之前...现在我明白了。

- einpoklum

@RobertCrovella：我想这是一个答案......但是-也太差了吧！为什么nvcc不支持ths？ - einpoklum

Answer 1

看起来你正在尝试跨库边界进行设备代码链接。目前，这只有使用静态库才可能。

我知道的选项是切换到静态库/链接安排，或者重构您的代码，以便您不需要在动态库边界上链接设备代码。

- Robert Crovella

---

非常感谢您的回答。我们/我不得不处理这个问题，直到Nvidia解决它。希望尽快解决，因为这真的很不方便。 - Tim

这个限制已经存在多年了，自从rdc被引入以来。如果你想看到改变，你可能需要在developer.nvidia.com上提交一个请求。 - Robert Crovella