CUDA构建共享库

Posted

tags:

篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了CUDA构建共享库相关的知识,希望对你有一定的参考价值。

我需要为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 ++正确编译。您有如何解决此问题的建议吗?

答案

您似乎正在尝试跨库边界进行设备代码链接。目前,那是only possible with a static library

我所知道的选项是切换到静态库/链接安排,或者重构代码,这样就不需要跨动态库边界链接设备代码。

以上是关于CUDA构建共享库的主要内容,如果未能解决你的问题,请参考以下文章

银行冲突CUDA共享内存?

ndk的系统构建

CUDA内核包装器的共享库未定义引用

CUDA:重载共享内存以实现多个数组的归约方法

使用共享内存时不执行 CUDA 内核代码

CUDA 内存库冲突