c - 读取全局标志不适用于 CUDA 中的 CPU>GPU 数据交换

Question

我尝试使用 CUDA 在 CPU 和 K80 GPU 之间进行简单的单向通信。我想要一个位于全局内存中并由所有正在运行的 GPU/内核线程轮询的 bool 取消标志。该标志应默认为 false，并且可以在正在进行的计算期间由 CPU/主机线程设置为 true。然后 GPU/内核线程应该退出。

这是我尝试过的。我已经简化了代码。我删除了错误检查和应用程序逻辑(包括阻止并发访问 cancelRequested 的应用程序逻辑)。

在主机端，全局定义(.cpp):

// Host side thread safety of this pointer is covered by application logic
volatile bool* cancelRequested = nullptr; 

在计算线程 (.cpp) 的主机端:

initialize(&cancelRequested);
compute(cancelRequested);
finalize(&cancelRequested);

在主机端的主线程 (.cpp) 中:

cancel(cancelRequested); // Called after init is finished

主机例程(.cu 文件):

void initialize(volatile bool** pCancelRequested)
{
   cudaMalloc(const_cast<bool**>(pCancelRequested), sizeof(bool));
   const bool aFalse = false;
   cudaMemcpy(*const_cast<bool**>(pCancelRequested), &aFalse, sizeof(bool), cudaMemcpyHostToDevice);
}

void compute(volatile bool* pCancelRequested) 
{
   ....
   computeKernel<<<pBlocksPerGPU, aThreadsPerBlock>>>(pCancelRequested);
   cudaDeviceSynchronize(); // Non-busy wait
   ....
}

void finalize(volatile bool** pCancelRequested)
{
   cudaFree(*const_cast<bool**>(pCancelRequested));
   *pCancelRequested = nullptr;
}

void cancel(volatile bool* pCancelRequested)
{
   const bool aTrue = true;
   cudaMemcpy(const_cast<bool*>(pCancelRequested), &aTrue, sizeof(bool), cudaMemcpyHostToDevice);
}

设备例程(.cu 文件):

__global__ void computeKernel(volatile bool* pCancelRequested)
{
   while (someCondition)
   {
      // Computation step here

      if (*pCancelRequested)
      {
         printf("-> Cancel requested!\n");
         return;
      }
   }
}

代码运行良好。但它永远不会进入取消情况。我成功读回了 initialize() 和 cancel() 中的 false 和 true 值，并使用 gdb 检查了它们。 IE。写入全局标志效果很好，至少从主机端的角度来看是这样。但是，内核永远不会看到取消标志设置为 true 并从外部 while 循环正常退出。

知道为什么这不起作用吗？

Answer 1

我发现您的方法的根本问题是 cuda 流会阻止它工作。

CUDA 流有两个基本原则:

1. 发布到同一流中的项目不会重叠；他们会序列化。
2. 发布到单独创建的流中的项目有可能重叠； CUDA 提供的这些操作没有定义的顺序。

即使您没有明确使用流，您也会在“默认流”中进行操作，并且应用相同的流语义。

在这个简短的总结中，我并没有涵盖有关流的所有知识。您可以在this online training series的第7单元中了解有关CUDA流的更多信息。

由于 CUDA 流，这:

 computeKernel<<<pBlocksPerGPU, aThreadsPerBlock>>>(pCancelRequested);

还有这个:

 cudaMemcpy(const_cast<bool*>(pCancelRequested), &aTrue, sizeof(bool), cudaMemcpyHostToDevice);

不可能重叠(它们被发送到相同的“默认”CUDA 流中，因此上面的规则 1 表明它们不可能重叠)。但如果您想向正在运行的内核发出“信号”，重叠是必不可少的。我们必须允许cudaMemcpy操作在内核运行的同时发生。

我们可以通过直接应用 CUDA 流来解决这个问题(注意上面的规则 2)，将复制操作和计算(内核)操作放入单独创建的流中，以便允许它们重叠。当我们这样做时，事情就会按预期进行:

$ cat t2184.cu
#include <iostream>
#include <unistd.h>

__global__ void k(volatile int *flag){

  while (*flag != 0);
}

int main(){

  int *flag, *h_flag = new int;
  cudaStream_t s[2];
  cudaStreamCreate(s+0);
  cudaStreamCreate(s+1);
  cudaMalloc(&flag, sizeof(h_flag[0]));
  *h_flag = 1;
  cudaMemcpy(flag, h_flag, sizeof(h_flag[0]), cudaMemcpyHostToDevice);
  k<<<32, 256, 0, s[0]>>>(flag);
  sleep(5);
  *h_flag = 0;
  cudaMemcpyAsync(flag, h_flag, sizeof(h_flag[0]), cudaMemcpyHostToDevice, s[1]);
  cudaDeviceSynchronize();
}

$ nvcc -o t2184 t2184.cu
$ compute-sanitizer ./t2184
========= COMPUTE-SANITIZER
========= ERROR SUMMARY: 0 errors
$

注释:

• 虽然从静态文本打印输出中看不出来，但程序在退出之前大约花费了 5 秒。如果注释掉诸如*h_flag = 0;这样的行，那么程序将挂起，表明flag信号方法工作正常。
• 请注意 volatile 的使用。这是necessary为了指示编译器对该数据的任何访问都必须是实际访问，不允许编译器进行会阻止在预期位置发生内存读取或写入的修改。

这种主机->设备信号行为也可以在不显式使用流的情况下实现，但使用主机固定内存作为信令位置，因为它对主机和设备代码“同时”“可见”。这是一个例子:

#include <iostream>
#include <unistd.h>

__global__ void k(volatile int *flag){

  while (*flag != 0);
}

int main(){

  int *flag;
  cudaHostAlloc(&flag, sizeof(flag[0]), cudaHostAllocDefault);
  *flag = 1;
  k<<<32, 256>>>(flag);
  sleep(5);
  *flag = 0;
  cudaDeviceSynchronize();
}

对于其他信令示例，例如从设备到主机，其他读者可能会对 this 感兴趣。 .