使用 dlopen 从共享库加载设备函数答案

【问题标题】：Load device function from shared library with dlopen使用 dlopen 从共享库加载设备函数
【发布时间】：2016-02-12 13:27:07
【问题描述】：

我对 cuda 编程比较陌生，找不到解决问题的方法。

我正在尝试拥有一个共享库，我们称之为func.so，它定义了一个设备函数

__device__ void hello(){ prinf("hello");}

然后我希望能够通过 dlopen 访问该库，并在我的程序中使用该函数。我尝试了以下几行：

func.cu

#include <stdio.h>
typedef void(*pFCN)();

__device__ void dhello(){
    printf("hello\n")
}

__device__ pFCN ptest = dhello;
pFCN h_pFCN;

extern "C" pFCN getpointer(){
    cudaMemcpyFromSymbol(&h_pFCN, ptest, sizeof(pFCN));
    return h_pFCN;
}

main.cu

#include <dlfcn.h>
#include <stdio.h>

typedef void (*fcn)();
typedef fcn (*retpt)();
retpt hfcnpt;
fcn hfcn;

__device__ fcn dfcn;
__global__ void foo(){
    (*dfcn)();
}
int main() {
    void * m_handle = dlopen("gputest.so", RTLD_NOW);
    hfcnpt = (retpt) dlsym( m_handle, "getpointer");
    hfcn = (*hfcnpt)();
    cudaMemcpyToSymbol(dfcn, &hfcn, sizeof(fcn), 0, cudaMemcpyHostToDevice);
    foo<<<1,1>>>();
    cudaThreadSynchronize();
    return 0;
}

但是这样在使用 cuda-gdb 调试时会出现以下错误：

CUDA Exception: Warp Illegal Instruction

Program received signal CUDA_EXCEPTION_4, Warp Illegal Instruction.
0x0000000000806b30 in dtest () at func.cu:5

感谢大家能给我的任何帮助！ :)

【问题讨论】：

标签： cuda

【解决方案1】：

从另一个编译单元中的设备代码调用一个编译单元中的__device__ 函数需要separate compilation with device linking 使用nvcc。

但是，这种用法与 libraries only works with static libraries.

因此，如果目标__device__ 函数在.so 库中，而调用代码在.so 库之外，则您的方法行不通，使用当前的nvcc 工具链。

我可以建议的唯一“解决方法”是将所需的目标函数放在静态库中，或者将调用者和目标放在同一个 .so 库中。 cuda 标签上有许多问题/答案，它们提供了这些替代方法的示例。

【讨论】：