到主机的 CUDA 信号

Question

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

我正在研究一个迭代过程，其中计算是在设备中进行的，并且在每次迭代之后，都会将一个布尔变量传递给主机，以告知该过程是否已经收敛。根据该变量，主机决定停止迭代或进行另一轮迭代。

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

我看过的替代方案。

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

2. 使用 CudaHostAllocMapped 来避免 CudaMemCpy 这没有用，因为与标准固定内存分配 + CudaMemCpy 相比，它没有任何基于时间的优势。 （第 460 页，多核和 GPU 编程，一种集成方法，Morgran Kruffmann 2014）。

希望有其他方法可以解决这个问题。

Answer 1

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

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

1. 使用映射/固定内存。国际海事组织，您在问题第 2 项中的主张不受支持，没有比我们许多人可能无法查看的书的页面引用更多的上下文。

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

3. 使用流水线算法，并将推测内核启动与布尔值的设备->主机副本重叠，用于每个流水线阶段。

我认为上面的前两项相当明显，因此我将为第 3 项提供一个工作示例。基本思想是我们将在两个流之间进行 ping-pong，将内核交替启动到一个流然后另一个流。我们将有第三个流，以便我们可以将设备->主机复制操作与下一次启动的执行重叠。由于 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 的行为。