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并明确地管理内核启动。