使用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

我感謝您能給我的任何幫助！ :)

來源

2016-02-12 Christoph

從另一個編譯單元中的設備代碼在一個編譯單元中調用__device__函數需要使用nvcc的separate compilation with device linking。

然而，庫的這種用法only works with static libraries。

因此，如果目標__device__功能是的.so庫，並調用代碼是外.so庫，你的方法不能工作，與當前nvcc工具鏈。

我可以建議的唯一「解決方法」是將所需的目標函數放在靜態庫中，否則將調用者和目標放在同一個庫中。關於cuda標籤有許多問題/答案，它們給出了這些替代方法的例子。

來源

2016-02-12 05:27:07

使用dlopen從共享庫加載設備函數

回答

相關問題