当我尝试使用功能指针调用CUDA内核(__global__
函数(时,一切似乎都可以正常工作。但是,如果我在调用内核时忘记提供启动配置,则NVCC不会导致错误或警告,但是该程序会编译,然后如果我尝试运行它。
__global__ void bar(float x) { printf("foo: %fn", x); }
typedef void(*FuncPtr)(float);
void invoker(FuncPtr func)
{
func<<<1, 1>>>(1.0);
}
invoker(bar);
cudaDeviceSynchronize();
编译并运行上述。一切都会好起来的。然后,删除内核的启动配置(即,&lt;&lt;&lt; 1,1>>>(。该代码会很好地编译,但是当您尝试运行它时它会崩溃。
知道发生了什么事吗?这是一个错误,还是我不应该传递__global__
功能的指针?
CUDA版本:8.0
OS版本:Debian(测试回购(GPU:NVIDIA GEFORCE 750m
如果我们采用了您的repro的更复杂版本,并查看CUDA工具链前端发出的代码,则可以看到发生了什么:
#include <cstdio>
__global__ void bar_func(float x) { printf("foo: %fn", 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并明确管理内核启动。