CUDA 构建共享库

Question

我需要为 cuda 创建一个共享库。库的编译工作正常，但是当我尝试在我的程序中使用它时 nvcc 返回链接器或 ptxas 错误。

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

文件A.h

#ifndef FILEA_H_
#define FILEA_H_

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

文件A.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;
}

临时类.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

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

主程序

#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++ 正确编译。你有什么建议如何解决这个问题。

Answer 1

您似乎正在尝试跨库边界进行设备代码链接。目前，这仅适用于静态库。

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