内核启动后用cudaDeviceSynchronize做CUDA错误检查是否存在遗漏？

为什么cudaGetLastError()能捕获cudaDeviceSynchronize()检测不到的错误？ #

CUDA内核启动的错误分为两类，对应完全不同的错误报告逻辑：

• 内核启动前置错误：比如网格/块维度超限、参数类型不匹配、设备未初始化、共享内存分配超限等，这类错误是在主机端发起内核调用的瞬间就触发的，不会进入CUDA全局错误队列。
• 内核执行过程错误：比如内存越界、非法指令等，这类错误会在内核运行时产生，进入全局错误队列等待同步调用检测。

cudaDeviceSynchronize()仅会检测进入全局错误队列的执行阶段错误，而cudaGetLastError()会直接返回当前线程上下文里的最后一个错误（包括启动前置错误），并且会清除这个错误状态——这就是它能抓到同步调用检测不到的错误的核心原因。而cudaPeekAtLastError()只是查看不清除错误，后续API调用的错误可能会覆盖它，所以更推荐使用cudaGetLastError()。

内核启动后调用cudaDeviceSynchronize()的错误检查方式会在哪些场景失效？ #

以下场景中，你的错误检查方法会遗漏关键错误：

• 内核启动参数非法：比如传递的网格维度超过设备支持的最大值、核函数参数类型与定义不符（如把主机指针传给设备指针参数）、共享内存申请大小超出设备限制，这类错误在启动瞬间就产生，不会等到内核执行，同步调用无法检测。
• 设备未就绪或资源耗尽：未初始化CUDA设备就启动内核、当前设备流的硬件资源（如线程块、寄存器）被耗尽无法启动新内核，这类启动失败的错误同样不会被同步调用捕获。
• 错误被后续API覆盖：如果内核启动后、调用cudaDeviceSynchronize()之前执行了其他CUDA API调用，且这些调用产生了新错误，同步调用只会返回最后一个错误，之前的内核启动错误会被覆盖丢失。

内容的提问来源于stack exchange，提问作者zaphodxvii