使用庫傳遞CUDA函數指針

Question

我正在使用CUDA並嘗試使用函數指針將CUDA函數傳遞到稍後在其設備內核中使用此函數的庫，類似於CUDA函數指針示例。

代碼的重要部分是：

/** Type definition for the execution function in #qsched_run. */ 
typedef void (*qsched_funtype)(int , void *); 

__device__ void gpuTest(int type , void *data) 
{ 
    .... 
} 
__device__ qsched_funtype function = gpuTest; 

void main(...) 
{ 
//Various initialization setup. 

if(cudaMemcpyFromSymbol(&func , function , sizeof(qsched_funtype)) != cudaSuccess) 
    error("Failed to copy function pointer from device"); 

qsched_run_CUDA(&s , func); 
} 

的qsched_run_CUDA函數是一個庫函數，它一些初始化，複製函數指針裝置（給一個變量它可以看到），然後運行內核在某些時候使用該函數指針調用gpuTest函數。

代碼編譯正確地提供我使用-G具有以下NVCC呼叫：

nvcc -g -G -m64 -I../src ../src/.libs/libquicksched_cuda.a -L/home/aidan/cuda_6.0/lib -L/home/aidan/cuda_6.0/lib64 -lcudart -lcuda -DWITH_CUDA -gencode arch=compute_30,code=sm_30 -lgomp test_gpu_simple.cu -o out.out 

其中

../src/.libs/libquicksched_cuda.a 

是包含qsched_run_CUDA函數的庫。

的那一刻，我從我的NVCC通話中移除-G標誌後，突然這一切休息，並在qsched_run_CUDA內核運行與無效程序計數器錯誤崩潰，和函數指針（包括我自己。 cu文件）設置爲0x4。

大概我需要使用單獨彙編CUDA（http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#using-separate-compilation-in-cuda）爲Cuda function pointer consistency隱約解釋 - 但是我不知道如何使用庫函數時，要做到這一點，無論是NVCC的指導，也沒有計算器鏈接使它明顯如何做這個。

有沒有人有這方面的經驗？我試圖簡單地嘗試制定nvlink來做到這一點，但我沒有遠遠（我把它傳遞給圖書館似乎並不高興）。

Answer 1

是的，您將需要使用單獨的編譯。根據您目前所顯示的內容以及使用文檔中的nvcc separate compilation library example，我製作了一個簡單的測試用例。下面是代碼：

kernel_lib.cu：

#include <stdio.h> 

#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

/** Type definition for the execution function in #qsched_run. */ 
typedef void (*qsched_funtype)(int , void *); 

__global__ void mykernel(int type, void *data, void *func){ 
    ((qsched_funtype)func)(type, data); 
} 

int qsched_run_CUDA(int val, void *d_data, void *func) 
{ 
    mykernel<<<1,1>>>(val, d_data, func); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("kernel fail"); 
    return 0; 
} 

main.cu：

#include <stdio.h> 
#define DATA_VAL 5 

int qsched_run_CUDA(int, void*, void*); 

#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

/** Type definition for the execution function in #qsched_run. */ 
typedef void (*qsched_funtype)(int , void *); 

__device__ void gpuTest(int type , void *data) 
{ 
    ((int *)data)[0] = type; 
} 
__device__ qsched_funtype function = gpuTest; 


int main() 
{ 
    void *func; 
    cudaMemcpyFromSymbol(&func , function , sizeof(qsched_funtype)); 
    cudaCheckErrors("Failed to copy function pointer from device"); 
    int h_data = 0; 
    int *d_data; 
    cudaMalloc((void **)&d_data, sizeof(int)); 
    cudaCheckErrors("cudaMalloc fail"); 
    cudaMemset(d_data, 0, sizeof(int)); 
    cudaCheckErrors("cudaMemset fail"); 
    int return_val = qsched_run_CUDA(DATA_VAL, (void *)d_data, func); 
    if (return_val != 0) printf("return code error\n"); 
    cudaMemcpy(&h_data, d_data, sizeof(int), cudaMemcpyDeviceToHost); 
    cudaCheckErrors("cudaMemcpy fail"); 
    if (h_data != DATA_VAL) {printf("Fail! %d\n", h_data); return 1;} 
    printf("Success!\n"); 
    return 0; 
} 

編譯命令和結果：

$ nvcc -arch=sm_20 -dc kernel_lib.cu 
$ nvcc -lib kernel_lib.o -o test.a 
$ nvcc -arch=sm_20 -dc main.cu 
$ nvcc -arch=sm_20 main.o test.a -o test 
$ ./test 
Success! 
$ 

我用CUDA 5.0該試驗。