c++ - CUDA 信号到主机

Question

有没有办法在内核执行结束时向主机发出信号(成功/失败)？

我正在查看一个迭代过程，其中在设备中进行计算，并且在每次迭代之后，一个 bool 变量被传递到主机以告知该过程是否已经收敛。基于变量，主机决定停止迭代或进行另一轮迭代。

在每次迭代结束时复制单个 bool 变量会使通过并行化获得的时间增益无效。因此，我想找到一种方法让宿主知道收敛状态(成功/失败)，而不必每次都使用 CudaMemCpy。注意:使用固定内存传输数据后存在时间问题。

我看过的替代品。

1. asm("陷阱;"); &断言();这些将分别触发主机中的未知错误和 cudaErrorAssert。不幸的是，它们是“粘性的”，因为无法使用 CudaGetLastError 重置错误。唯一的方法是使用 cudaDeviceReset() 重置设备。

2. 使用 CudaHostAllocMapped 避免 CudaMemCpy 这是没有用的，因为它不提供任何基于时间的优势超过标准固定内存分配 + CudaMemCpy。 (第 460 页，多核和 GPU 编程，综合方法，Morgran Kruffmann 2014)。

将感谢其他解决此问题的方法。

Answer 1

我怀疑这里的真正问题是您的迭代内核运行时间非常短(大约 100us 或更少)，这意味着每次迭代的工作量非常小。最好的解决方案可能是尝试增加每次迭代的工作量(重构您的代码/算法，解决更大的问题等)

但是，这里有一些可能性:

1. 使用映射/固定内存。 IMO，你在问题第 2 项中的声明没有更多的背景信息，只是引用了一本书，而我们中的许多人可能都看不到这本书。

2. 使用动态并行。将您的内核启动过程移至发布子内核的 CUDA 父内核。子内核设置的任何 bool 值都将立即在父内核中被发现，无需任何 cudaMemcpy 操作或映射/固定内存。

3. 使用流水线算法，并在每个流水线阶段将推测性内核启动与 bool 值的设备-> 主机拷贝重叠。

我认为上面的前两项相当明显，所以我将为第 3 项提供一个可行的示例。基本思想是我们将在两个流之间来回切换，将内核交替启动到一个流然后另一个。我们将有第三个流，以便我们可以将设备-> 主机复制操作与下一次启动的执行重叠。由于 D->H 复制与内核执行的重叠，复制操作实际上没有“成本”，它被内核执行工作隐藏了。

这是一个完整的示例，加上一个 nvvp 时间轴:

$ cat t267.cu
#include <stdio.h>


const int stop_count = 5;
const long long tdelay = 1000000LL;

__global__ void test_kernel(int *icounter, bool *istop, int *ocounter, bool *ostop){

  if (*istop) return;
  long long start = clock64();
  while (clock64() < tdelay+start);
  int my_count = *icounter;
  my_count++;
  if (my_count >= stop_count) *ostop = true;
  *ocounter = my_count;
}

int main(){
  volatile bool *v_stop;
  volatile int *v_counter;
  bool *h_stop, *d_stop1, *d_stop2, *d_s1, *d_s2, *d_ss;
  int *h_counter, *d_counter1, *d_counter2, *d_c1, *d_c2, *d_cs;
  cudaStream_t s1, s2, s3, *sp1, *sp2, *sps;
  cudaEvent_t e1, e2, *ep1, *ep2, *eps;
  cudaStreamCreate(&s1);
  cudaStreamCreate(&s2);
  cudaStreamCreate(&s3);
  cudaEventCreate(&e1);
  cudaEventCreate(&e2);
  cudaMalloc(&d_counter1, sizeof(int));
  cudaMalloc(&d_stop1, sizeof(bool));
  cudaMalloc(&d_counter2, sizeof(int));
  cudaMalloc(&d_stop2, sizeof(bool));
  cudaHostAlloc(&h_stop, sizeof(bool), cudaHostAllocDefault);
  cudaHostAlloc(&h_counter, sizeof(int), cudaHostAllocDefault);
  v_stop = h_stop;
  v_counter = h_counter;
  int n_counter = 1;
  h_stop[0] = false;
  h_counter[0] = 0;
  cudaMemcpy(d_stop1, h_stop, sizeof(bool), cudaMemcpyHostToDevice);
  cudaMemcpy(d_stop2, h_stop, sizeof(bool), cudaMemcpyHostToDevice);
  cudaMemcpy(d_counter1, h_counter, sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_counter2, h_counter, sizeof(int), cudaMemcpyHostToDevice);
  sp1 = &s1;
  sp2 = &s2;
  ep1 = &e1;
  ep2 = &e2;
  d_c1 = d_counter1;
  d_c2 = d_counter2;
  d_s1 = d_stop1;
  d_s2 = d_stop2;
  test_kernel<<<1,1, 0, *sp1>>>(d_c1, d_s1, d_c2, d_s2);
  cudaEventRecord(*ep1, *sp1);
  cudaStreamWaitEvent(s3, *ep1, 0);
  cudaMemcpyAsync(h_stop, d_s2, sizeof(bool), cudaMemcpyDeviceToHost, s3);
  cudaMemcpyAsync(h_counter, d_c2, sizeof(int), cudaMemcpyDeviceToHost, s3);
  while (v_stop[0] == false){
    cudaStreamWaitEvent(*sp2, *ep1, 0);
    sps = sp1; // ping-pong
    sp1 = sp2;
    sp2 = sps;
    eps = ep1;
    ep1 = ep2;
    ep2 = eps;
    d_cs = d_c1;
    d_c1 = d_c2;
    d_c2 = d_cs;
    d_ss = d_s1;
    d_s1 = d_s2;
    d_s2 = d_ss;
    test_kernel<<<1,1, 0, *sp1>>>(d_c1, d_s1, d_c2, d_s2);
    cudaEventRecord(*ep1, *sp1);
    while (n_counter > v_counter[0]);
    n_counter++;
    if(v_stop[0]  == false){
      cudaStreamWaitEvent(s3, *ep1, 0);
      cudaMemcpyAsync(h_stop, d_s2, sizeof(bool), cudaMemcpyDeviceToHost, s3);
      cudaMemcpyAsync(h_counter, d_c2, sizeof(int), cudaMemcpyDeviceToHost, s3);
    }
  }
  cudaDeviceSynchronize();  // optional
  printf("terminated at counter = %d\n", v_counter[0]);
}
$ nvcc -arch=sm_52 -o t267 t267.cu
$ ./t267
terminated at counter = 5
$

在上图中，我们看到明显有 5 次内核启动(实际上是 6 次)并且它们在两个流之间来回跳动。 (我们从代码组织和流水线中期望的第 6 次内核启动是上面 stream15 末尾的非常短的一行。此内核启动但立即见证 stop 为真，因此它退出.) 设备 -> 主机拷贝位于第三个流中。如果我们仔细观察从一个内核迭代到下一个内核迭代的切换:

我们看到，即使是这些非常短的 D->H memcpy 操作也基本上与下一个内核执行重叠。作为引用，上述内核执行之间的间隔约为 5us。

请注意，这完全是在 Linux 上完成的。如果您在 Windows WDDM 上尝试此操作，由于 WDDM 命令批处理，可能很难实现类似的效果。然而，Windows TCC 应该大致复制 linux 行为。