使用循环缓冲区的 cuda 核外实现

Question

我正在尝试在 GPU 内存和 CPU 内存之间进行核外处理。例如，我有每个 1GB 的数据块，我需要按顺序处理 1000 个这样的数据块，每个数据块都是由内核启动完成的。假设必须一个一个处理，因为第n个内核启动需要使用第(n-1)个内核产生的结果，该结果存储在第(n-1)个块中，除了第一次内核启动。所以我正在考虑使用GPU上的循环缓冲区来存储最近的5个块，并使用事件在数据流和任务流之间进行同步。数据流准备数据，任务流启动内核。代码如下图所示。

const int N_CBUF = 5, N_TASK = 1000;

// Each pointer points to a data block of 1GB
float* d_cir_buf[N_CBUF];
float* h_data_blocks[N_TASK];

// The data stream for transfering data from host to device.
// The task stream for launching kernels to process the data.
cudaStream_t s_data, s_task;

// The data events for the completion of each data transfer.
// The task events for the completion of each kernel execution.
cudaEvent_t e_data[N_TASK], e_task[N_TASK];

// ... code for creating the streams and events.

for (int i = 0; i < N_TASK; i++) {
  // Data transfer should not overwritten the data needed by the kernels.
  if (i >= N_CBUF) {
    cudaStreamWaitEvent(s_data, e_task[i-N_CBUF+1]);
  }
  cudaMemcpyAsync(d_cir_buf[i % N_CBUF], h_data_blocks[i], ..., cudaMemcpyHostToDevice, s_data);
  cudaEventRecord(e_data[i], s_data);

  cudaStreamWaitEvent(s_task, e_data[i]);

  // Pass the current and the last data block to the kernel.
  my_kernel<<<..., s_task>>>(d_cir_buf[i % N_CBUF], 
    i == 0 ? 0 : d_cir_buf[(i+N_CBUF-1)%N_CBUF]);
  cudaEventRecord(e_task[i], s_task);
}

我想知道这是否是一个有效的想法，还是有什么完全错误的？另外，CUDA编程指南提到，如果从两个不同的主机内存地址到同一个设备地址有memcpy，那么就不会并发执行，这对我来说有关系吗？特别是d_cir_buf的内存如果被分配为一个大块，然后分成5块，会不会算“设备内存地址相同”，导致并发失败？此外，在我的情况下，第 (n+5) 次数据传输将转到与第 n 次数据传输相同的地址，但是，鉴于需要同步，不会同时执行两个这样的传输.那么这样可以吗？

Answer 1

我觉得你的问题最适合双缓冲：

• 两个流
• 在stream1中上传data1
• 在 stream1 中的 data1 上运行内核
• 在stream2中上传data2
• 在 stream2 中的 data2 上运行内核

...等等

stream2 中的内核可以与 strezm 1 中的数据传输重叠，反之亦然