在内核运行时将设备全局内存复制到固定主机内存时会出现永久过时的值

Question

我有一个在后台运行半无限循环的内核（只要全局哨兵值为真，就会运行 while 循环）。这是一个接近最小示例的内容，为了简洁起见，省略了一些行和错误检查。假设设备支持任何当前启用 CUDA 的设备上存在的最大异步行为集（引擎计数等）：

int main() {
  [Create streams, declare variables, etc.]
  ...
  cudaMalloc(&running_d, sizeof(bool));      // Global device memory
  cudaMalloc(&output_d, sizeof(bool));       // Global device memory
  cudaMallocHost(&running_h, sizeof(bool));  // Pinned host memory
  cudaMallocHost(&output_h, sizeof(bool));   // Pinned host memory

  // Set sentinel true
  bool running_h = true;
  cudaMemcpy(running_d, running_h, sizeof(bool), cudaMemcpyHostToDevice);

  // Launch kernel
  kernel<<<1,1, 0, stream1>>>(running_d, output_d);

  // Copy memory concurrently with kernel execution until output is false
  *output_h = true;
  while (*output_h) {
    cudaMemcpyAsync(output_h, output_d, sizeof(bool), cudaMemcpyDeviceToHost, stream2);
    cudaStreamSynchronize(stream2);
  }

  return 0;
}

__global__
void kernel(bool* running_ptr, bool* output_ptr) {
  while (*running_ptr) {
    *output_ptr = liveFunction();  // Some function which eventually will always return false
  }
}

我的问题实际上相当于主机循环是否会退出。在一个天真的假设的系统模型中，来自

*output_ptr = liveFunction()

的写入最终将对主机可见，并且主机循环将退出。

但是，在我的测试（CUDA 12、RTX 4090）中，对

cudaMemcpyAsync(output_h, output_d, sizeof(bool), cudaMemcpyDeviceToHost, stream2)

的调用在内核运行时异步执行，但最终还是处理了

true

的过时值，即使在我与设备端确认之后也是如此

printf("%d\n", *output_ptr)

该值已设置为

false

。这种情况似乎永远持续（即至少几分钟）。如果我通过同时从主机将

running_d

设置为

false

来停止内核，内核将退出，然后主机循环复制

false

的更新值并退出。

我也尝试过使用

atomicCAS

设置

output_ptr

而不是赋值运算符，但我仍然得到相同的结果。我不确定

output_ptr

的

false

值存储在哪里，因为

atomicCAS

似乎意味着新值对其他设备线程可见，而对主机不可见。

就无限内核循环的目的而言，我很清楚这被认为是在设备代码中要避免的事情，内核充当无限生成器，我想从中提取中间结果并可能需要用户干预。我知道使用无限循环有两个优点，据我所知，这是任何其他方式都无法获得的：

• 局部内核变量可以保存在寄存器中。如果内核不断停止和重新启动，则每次内核停止和重新启动时都需要将本地数据复制到全局内存或从全局内存复制。
• 扭曲可以发散（彼此不同，同时保持其车道连贯），而不会产生任何后果。如果内核被停止并重新启动，每次停止都会创建一段时间，等待最延迟的扭曲完成，可能占用率较低。

我知道 2013 年的一个答案，声称当时 无法保证这种行为，而另一个答案 推荐这种哨兵机制 ，尽管只是为了欺骗编译器。