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

问题描述：

我想有一个共享库，让叫它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

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

答

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

然而，库的这种用法only works with static libraries。

因此，如果目标__device__功能是的.so库，并调用代码是外.so库，你的方法不能工作，与当前nvcc工具链。

我可以建议的唯一“解决方法”是将所需的目标函数放在静态库中，否则将调用者和目标放在同一个库中。关于cuda标签有许多问题/答案，它们给出了这些替代方法的例子。

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

相关推荐