如何捕获cudaLaunchHostFunc回调函数中抛出的异常？

我需要在CUDA图内检查托管内存中由某条流上运行的内核写入的错误标志，并根据该标志抛出异常。原本可以直接同步该流后在主机端检查标志，但CUDA图场景下必须通过cudaLaunchHostFunc回调来实现主机端的错误检查。

官方文档未提及cudaLaunchHostFunc对异常的处理方式，是否有办法捕获该回调函数内抛出的异常？

最小可复现代码（MWE） #

#include <iostream>
#include <stdexcept>

__global__ void kern(){
  int id = blockIdx.x*blockDim.x + threadIdx.x;
  printf("Kernel\n");
  return;
}

void foo(void* data){
  std::cerr << "Callback" << std::endl;
  throw std::runtime_error("Error in callback");
}

void launch(){
  cudaStream_t st = 0;
  kern<<<1,1,0,st>>>();
  cudaHostFn_t fn = foo;
  cudaLaunchHostFunc(st, fn, nullptr);
  cudaDeviceSynchronize();
}

int main(){
  try{
    launch();
  }
  catch(...){
    std::cerr << "Catched exception" << std::endl;
  }
  return 0;
}

代码输出 #

Kernel
Callback
terminate called after throwing an instance of 'std::runtime_error'
  what():  Error in callback
Aborted (core dumped)

可以看到异常已抛出，但并未传播到launch函数中。我原本期望上述launch()函数的异常处理逻辑与以下代码等效：

对比代码 #

void launch(){
  cudaStream_t st = 0;
  kern<<<1,1,0,st>>>();
  cudaStreamSynchronize(st);
  foo(nullptr);
  // cudaHostFn_t fn = foo;
  // cudaLaunchHostFunc(st, fn, nullptr);
  cudaDeviceSynchronize();
}

对比代码输出 #

Kernel
Callback
Catched exception

此外，在第一种情况中，所有CUDA调用均返回cudaSuccess。

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