CUDA：忘記內核啓動配置不會導致NVCC編譯器警告或錯誤

Question

當我嘗試使用函數指針調用CUDA內核（__global__函數）時，一切看起來都很好。但是，如果我在調用內核時忘記提供啓動配置，NVCC不會導致錯誤或警告，但是如果我嘗試運行該程序，程序將會編譯並崩潰。

__global__ void bar(float x) { printf("foo: %f\n", x); } 

typedef void(*FuncPtr)(float); 

void invoker(FuncPtr func) 
{ 
    func<<<1, 1>>>(1.0); 
} 

invoker(bar); 
cudaDeviceSynchronize(); 

編譯並運行以上。一切都會工作得很好。然後，刪除內核的啓動配置（即，< < < 1,1 >>>）。代碼會很好的編譯，但是當你嘗試運行它時會崩潰。

任何想法是怎麼回事？這是一個錯誤，或者我不應該傳遞__global__函數的指針嗎？

CUDA版本：8.0

OS的版本：Debian的（測試回購） GPU：NVIDIA的GeForce 750M

Answer 1

如果我們把你的清樣的一個稍微複雜的版本，並期待在所發射的代碼CUDA工具鏈的前端，可以看到發生了什麼：

#include <cstdio> 

__global__ void bar_func(float x) { printf("foo: %f\n", x); } 
typedef void(*FuncPtr)(float); 

void invoker(FuncPtr passed_func) 
{ 
#ifdef NVCC_FAILS_HERE 
    bar_func(1.0); 
#endif 
    bar_func<<<1,1>>>(1.0); 
    passed_func(1.0); 
    passed_func<<<1,1>>>(2.0); 
} 

讓我們編譯幾個方面：

$ nvcc -arch=sm_52 -c -DNVCC_FAILS_HERE invoker.cu 
invoker.cu(10): error: a __global__ function call must be configured 

即前端可以檢測到bar_func是一個全局函數並需要啓動參數。另一種嘗試：

$ nvcc -arch=sm_52 -c -keep invoker.cu 

如您所述，這不會產生編譯錯誤。讓我們來看看發生了什麼：

void bar_func(float x) ; 
# 5 "invoker.cu" 
typedef void (*FuncPtr)(float); 
# 7 "invoker.cu" 
void invoker(FuncPtr passed_func) 
# 8 "invoker.cu" 
{ 
# 12 "invoker.cu" 
(cudaConfigureCall(1, 1)) ? (void)0 : (bar_func)((1.0)); 
# 13 "invoker.cu" 
passed_func((2.0)); 
# 14 "invoker.cu" 
(cudaConfigureCall(1, 1)) ? (void)0 : passed_func((3.0)); 
# 15 "invoker.cu" 
} 

標準內核調用語法<<<>>>被擴大到內聯調用cudaConfigureCall，然後主機包裝函數被調用。主機包裝必須啓動內核所需的API內部：

void bar_func(float __cuda_0) 
# 3 "invoker.cu" 
{__device_stub__Z8bar_funcf(__cuda_0); } 

void __device_stub__Z8bar_funcf(float __par0) 
{ 
    if (cudaSetupArgument((void *)(char *)&__par0, sizeof(__par0), (size_t)0UL) != cudaSuccess) return; 
    { volatile static char *__f __attribute__((unused)); __f = ((char *)((void (*)(float))bar_func)); 
     (void)cudaLaunch(((char *)((void (*)(float))bar_func))); 
    }; 
} 

所以存根只處理參數，並通過cudaLaunch啓動內核。它不處理啓動配置

崩潰的根本原因（實際上未檢測到運行時API錯誤）是內核啓動發生時沒有事先配置。很明顯，這是因爲CUDA前端（和C++）在編譯時不能執行指針自檢，並檢測到函數指針是調用內核的存根函數。

我認爲描述這一點的唯一方法是運行時API和編譯器的「限制」。我不會說你在做什麼是錯誤的，但我可能會在這種情況下使用驅動程序API並明確地管理內核啓動。