CUDA內核與函數指針和可變參數模板

Question

我想設計一個cuda框架，它將接受用戶函數並通過設備函數指針將它們轉發給內核。 CUDA可以使用可變參數模板（-stc = C++ 11），並且非常好。

但是，當內核調用設備函數指針時出現問題。顯然內核運行沒有問題，但GPU使用率爲0％。如果我簡單地用實際函數替換回調指針，那麼GPU使用率爲99％。這裏的代碼非常簡單，大的循環範圍僅僅是爲了使事情可以測量。我測與GPU狀態：

nvidia-smi --query-gpu=utilization.gpu,utilization.mory,memory.used --format=csv -lms 100 -f out.txt 

IIRC，用戶功能需要在同一個文件爲單位的內核（可能是執行#included）以NVCC成功。 func_d就在源代碼中，並且它編譯並運行良好，除了不使用函數指針（這是本設計中的全部要點）。

我的問題是： 爲什麼帶回調設備函數指針的內核不工作？

需要注意的是，當我printf的諾斯回調和func_d地址，它們是相同的，因爲在此示例輸出：

size of Args = 1 
callback() address = 4024b0 
func_d() address = 4024b0 

另一個奇怪的是，如果取消註釋在kernel()的callback()呼叫，則GPU使用率回到0％，即使func_d()調用仍然在那裏... func_d版本需要大約4秒鐘的運行時間，而回調版本不需要任何（約0.1秒）。

系統規格和編譯命令位於以下代碼的頭部。

謝謝！

// compiled with: 
// nvcc -g -G -O0 -std=c++11 -arch=sm_20 -x cu sample.cpp 
// 
// Nvidia Quadro 6000 (compute capability 2.0) 
// CUDA 6.5 (V6.5.12), 
// Arch Linux, Nvidia driver 343.22-4, gcc 4.9.1 
// Nov, 2014 


#include <stdio.h> 

__device__ 
void func_d(double* vol) 
{ 
*vol += 5.4321f; 
} 


// CUDA kernel function 
template <typename... Types> 
__global__ void kernel(void (*callback)(Types*...)) 
{ 
double val0 = 1.2345f; 

// // does not use gpu (0% gpu utilization) 
// for (int i = 0; i < 1000000; i++) { 
// callback(&val0); 
// } 

// uses gpu (99% gpu utilization) 
for (int i = 0; i < 10000000; i++) { 
func_d(&val0); 
} 
} 


// host function 
template <typename... Types> 
void host_func(void (*callback)(Types*...)) 
{ 
// get user kernel number of arguments. 
constexpr int I = sizeof...(Types); 
printf("size of Args = %d\n",I); 

printf("callback() address = %x\n",callback); 
printf("func_d() address = %x\n",func_d); 

dim3 nblocks = 100; 
int nthread = 100; 
kernel<Types...><<<nblocks,nthread>>>(callback); 
} 


__host__ 
int main(int argc, char** argv) 
{ 
host_func(func_d); 
} 

Answer 1

我的問題是：爲什麼用裝置的回調函數指針內核不工作？

可能有幾個問題需要解決。但最簡單的答案是因爲在主機代碼中取得設備實體的地址是非法的。對於設備變量以及設備功能來說，這是正確的。現在，你可以取這些實體的地址。但地址是垃圾。它不能在主機或設備上使用。如果你嘗試使用它們，你會在設備上產生未定義的行爲，這通常會讓你的內核停頓。

在主機代碼中可能會出現主機地址。設備地址可以在設備代碼中觀察到。任何其他行爲都需要API干預。

1. 你似乎是使用nvidia-smi利用查詢作爲東西是否被正確運行的措施。我建議您改爲使用proper cuda error checking，並且您也可以使用cuda-memcheck運行您的代碼。

2. 「爲什麼func_d的地址與callback的地址匹配？「因爲你正在服用都地址在主機代碼，和地址都是垃圾說服自己這一點，在你的內核的最後添加一行是這樣的：

   if ((!threadIdx.x)&&(!blockIdx.x)) printf("in-kernel func_d() address = %x\n",func_d); 
   

   ，你會看到它打印出與主機上打印的東西不同的內容

3. 「設備使用情況如何？」只要設備遇到錯誤，內核就會終止並且利用率爲零。爲你解釋這句話：「另一個奇怪的是，如果在kernel（）中取消註釋callback（）的調用，那麼GPU的使用率會回到0％，即使使用fu nc_d（）調用還在裏面......「

4. ‘我怎樣才能解決這個問題？’我不知道一個偉大的方式來解決這個問題。如果您在編譯時已知有限數量的CUDA函數，並希望用戶能夠從中進行選擇，那麼恰當的事情可能只是創建適當的索引，然後使用它來選擇函數。如果你真的想要的話，你可以運行一個初步的/設置的內核，它將獲取你關心的函數的地址，然後你可以將這些地址傳遞迴主機代碼，並在隨後的內核調用中用它們作爲參數，應該允許你的機制工作。但我不明白它是如何防止需要通過編譯時已知的一組預定義函數進行索引的。如果你正在前進的方向是，你希望用戶能夠在運行時提供用戶自定義函數我想你會覺得這很困難的時刻與CUDA運行時API做（我懷疑這很可能在未來改變），我提供了一個相當扭曲的機制，試圖做到這一點here（讀取整個問答; talonmies回答有內容，也）。另一方面，如果您願意使用CUDA驅動程序API，那麼它應該是可能的，儘管有些參與，因爲這正是PyCUDA中非常優雅的方式。

5. 在未來，請縮進代碼。

下面是一個完整的工作示例，演示了上述幾個想法。特別是，我顯示在一個相當粗糙的方式，即func_d地址可在設備代碼中存在時，則傳回主機，然後用作未來內核參數來成功地選擇/調用該設備的功能。

$ cat t595.cu 
// compiled with: 
// nvcc -g -G -O0 -std=c++11 -arch=sm_20 -x cu sample.cpp 
// 
// Nvidia Quadro 6000 (compute capability 2.0) 
// CUDA 6.5 (V6.5.12), 
// Arch Linux, Nvidia driver 343.22-4, gcc 4.9.1 
// Nov, 2014 


#include <stdio.h> 

__device__ 
void func_d(double* vol) 
{ 
    if ((!threadIdx.x) && (!blockIdx.x)) printf("value = %f\n", *vol); 
    *vol += 5.4321f; 
} 

template <typename... Types> 
__global__ void setup_kernel(void (**my_callback)(Types*...)){ 
    *my_callback = func_d; 
} 

// CUDA kernel function 
template <typename... Types> 
__global__ void kernel(void (*callback)(Types*...)) 
{ 
    double val0 = 1.2345f; 

// // does not use gpu (0% gpu utilization) 
// for (int i = 0; i < 1000000; i++) { 
    callback(&val0); 
// } 

    val0 = 0.0f; 
// uses gpu (99% gpu utilization) 
// for (int i = 0; i < 10000000; i++) { 
    func_d(&val0); 
// } 
    if ((!threadIdx.x)&&(!blockIdx.x)) printf("in-kernel func_d() address = %x\n",func_d); 
} 


// host function 
template <typename... Types> 
void host_func(void (*callback)(Types*...)) 
{ 
// get user kernel number of arguments. 
    constexpr int I = sizeof...(Types); 
    printf("size of Args = %d\n",I); 

    printf("callback() address = %x\n",callback); 
    printf("func_d() address = %x\n",func_d); 

    dim3 nblocks = 100; 
    int nthread = 100; 
    unsigned long long *d_callback, h_callback; 
    cudaMalloc(&d_callback, sizeof(unsigned long long)); 
    setup_kernel<<<1,1>>>((void (**)(Types*...))d_callback); 
    cudaMemcpy(&h_callback, d_callback, sizeof(unsigned long long), cudaMemcpyDeviceToHost); 
    kernel<Types...><<<nblocks,nthread>>>((void (*)(Types*...))h_callback); 
    cudaDeviceSynchronize(); 
} 


__host__ 
int main(int argc, char** argv) 
{ 
    host_func(func_d); 
} 
$ nvcc -std=c++11 -arch=sm_20 -o t595 t595.cu 
$ cuda-memcheck ./t595 
========= CUDA-MEMCHECK 
size of Args = 1 
callback() address = 4025dd 
func_d() address = 4025dd 
value = 1.234500 
value = 0.000000 
in-kernel func_d() address = 4 
========= ERROR SUMMARY: 0 errors 
$