CUDA信号给主持人

Posted 2021-04-21

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

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

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

我看过的替代品。

1. ASM（ “陷阱;”）; ＆assert（）;这些将分别触发主机中的未知错误和cudaErrorAssert。不幸的是，它们是“粘性的”，因为使用CudaGetLastError无法重置错误。唯一的方法是使用cudaDeviceReset（）重置设备。
2. 使用CudaHostAllocMapped来避免CudaMemCpy这是没有用的，因为它没有提供任何时间优势超过标准的固定内存分配+ CudaMemCpy。 （Pg 460，MultiCore和GPU编程，综合方法，Morgran Kruffmann 2014）。

将会欣赏其他方法来克服这个问题。

答案

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

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

1. 使用映射/固定内存。您在问题的第2项中的声明不受支持，IMO，没有比我们许多人可能无法查看的书籍的页面引用更多的上下文。
2. 使用动态并行。将内核启动过程移动到发出子内核的CUDA父内核。子内核设置的布尔值将在父内核中立即被发现，而不需要cudaMemcpy操作或映射/固定内存。
3. 使用流水线算法，并为每个管道阶段将推测内核启动与布尔的device->主机副本重叠。

我认为上面的前两个项目相当明显，所以我将为第3项提供一个有用的例子。基本的想法是我们将在两个流之间进行乒乓，将内核交替启动到一个流然后另一个。我们将有第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
", v_counter[0]);
}
$ nvcc -arch=sm_52 -o t267 t267.cu
$ ./t267
terminated at counter = 5
$

在上图中，我们看到5个内核启动很明显（实际上是6个），它们在两个流之间来回反弹。 （第6次内核启动，我们期望从代码组织和流水线操作，在上面的stream15结尾处是一个非常短的行。这个内核启动但立即见证stop是真的，所以它退出。）设备 - >主机副本位于第3个流中。如果我们密切关注从一次内核迭代到下一次迭代的切换：

我们看到即使是这些非常短的D-> H memcpy操作也基本上与下一次内核执行重叠。作为参考，上面的内核执行之间的差距大约是5us。

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