CUDA C++ 重叠 SERIAL 内核执行和数据传输

Posted 2023-03-23

【中文标题】CUDA C++ 重叠 SERIAL 内核执行和数据传输

【英文标题】：CUDA C++ overlapping SERIAL kernel execution and data transfer

【发布时间】：2020-12-04 20:59:22

【问题描述】：

因此，本指南here 展示了重叠内核执行和数据传输的一般方法。

cudaStream_t streams[nStreams];
for (int i = 0; i < nStreams; ++i) 
  cudaStreamCreate(&streams[i]);
  int offset = ...;
  cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i]);
  kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
  // edit: no deviceToHost copy

但是，内核是串行的。所以它必须处理 0->1000，然后 1000->2000，... 简而言之，在重叠数据传输时正确执行此内核的顺序是：

copy[a->b] 必须在 kernel[a->b] 之前发生 内核 [a->b] 必须发生在内核[b->c] 之前，其中 c > a, b

是否可以在不使用 cudaDeviceSynchronize() 的情况下做到这一点？如果没有，最快的方法是什么？

【问题讨论】：

为什么你认为你需要 cudaDeviceSynchronize() ？

cudaMemcpyAsync(d_a, h_a,..,cudaMemcpyHostToDevice,stream[i]); kernel1<<<...,stream[i]>>>(d_a, d_b,...); kernel2<<<...,stream[i]>>>(d_b, d_c, ...); cudaMemcpyAsync(h_c, d_c, ...,cudaMemcpyDeviceToHost,stream[i]);

在您的评论中，stream[i+1] 中的内核可能会在 stream[i] 中的内核完成之前被调用，这不符合我的要求

也许显示所有预期操作和所有依赖关系的序列图会有所帮助。例如，对于每个块，内核处理取决于该块的 H->D 副本，加上之前的块内核？每个块只有一个内核吗？这其中是否有 D->H 副本？总共有多少块？那么第一块呢？我猜那个内核不依赖于以前的内核？

第一个块必须始终首先运行。然后每个块都依赖于前一个块（当然还有那个块的内存传输）。所以确切的顺序是：进程数据 0->1000，进程数据 1000->2000，进程 2000->3000，...没有设备 -> 主机副本，我将对其进行编辑。

【参考方案1】：

所以每个内核都依赖于（不能开始直到）：

关联的 H->D 副本已完成 之前的内核执行完成

普通流语义无法处理这种情况（2 个独立的依赖项，来自 2 个独立的流），因此我们需要在其中放置一个额外的互锁。我们可以使用一组事件和cudaStreamWaitEvent()来处理它。

对于最一般的情况（不知道块的总数）我会推荐这样的东西：

$ cat t1783.cu
#include <iostream>
#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL

unsigned long long dtime_usec(unsigned long long start)

  timeval tv;
  gettimeofday(&tv, 0);
  return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;


template <typename T>
__global__ void process(const T * __restrict__ in, const T * __restrict__ prev, T * __restrict__ out, size_t ds)

  for (size_t i = threadIdx.x+blockDim.x*blockIdx.x; i < ds; i += gridDim.x*blockDim.x)
    out[i] = in[i] + prev[i];
    

const int nTPB = 256;
typedef int mt;
const int chunk_size = 1048576;
const int data_size = 10*1048576;
const int ns = 3;

int main()

  mt *din, *dout, *hin, *hout;
  cudaStream_t str[ns];
  cudaEvent_t  evt[ns];
  for (int i = 0; i < ns; i++) 
    cudaStreamCreate(str+i);
    cudaEventCreate( evt+i);
  cudaMalloc(&din, sizeof(mt)*data_size);
  cudaMalloc(&dout, sizeof(mt)*data_size);
  cudaHostAlloc(&hin,  sizeof(mt)*data_size, cudaHostAllocDefault);
  cudaHostAlloc(&hout, sizeof(mt)*data_size, cudaHostAllocDefault);
  cudaMemset(dout, 0, sizeof(mt)*chunk_size);  // for first loop iteration
  for (int i = 0; i < data_size; i++) hin[i] = 1;
  cudaEventRecord(evt[ns-1], str[ns-1]); // this event will immediately "complete"
  unsigned long long dt = dtime_usec(0);
  for (int i = 0; i < (data_size/chunk_size); i++)
    cudaStreamSynchronize(str[i%ns]); // so we can reuse event safely
    cudaMemcpyAsync(din+i*chunk_size, hin+i*chunk_size, sizeof(mt)*chunk_size, cudaMemcpyHostToDevice, str[i%ns]);
    cudaStreamWaitEvent(str[i%ns], evt[(i>0)?(i-1)%ns:ns-1], 0);
    process<<<(chunk_size+nTPB-1)/nTPB, nTPB, 0, str[i%ns]>>>(din+i*chunk_size, dout+((i>0)?(i-1)*chunk_size:0), dout+i*chunk_size, chunk_size);
    cudaEventRecord(evt[i%ns]);
    cudaMemcpyAsync(hout+i*chunk_size, dout+i*chunk_size, sizeof(mt)*chunk_size, cudaMemcpyDeviceToHost, str[i%ns]);
    
  cudaDeviceSynchronize();
  dt = dtime_usec(dt);
  for (int i = 0; i < data_size; i++) if (hout[i] != (i/chunk_size)+1) std::cout << "error at index: " << i << " was: " << hout[i] << " should be: " << (i/chunk_size)+1 << std::endl; return 0;
  std::cout << "elapsed time: " << dt << " microseconds" << std::endl;

$ nvcc -o t1783 t1783.cu
$ ./t1783
elapsed time: 4366 microseconds

这里的良好做法是使用分析器来验证预期的重叠场景。但是，我们可以根据经过的时间测量走捷径。

循环将总共 40MB 的数据传输到设备，并传回 40MB。经过的时间是4366us。这给出了 (40*1048576)/4366 或 9606 字节/us 的每个方向的平均吞吐量，即 9.6GB/s。这基本上使 Gen3 链路在两个方向上都饱和，因此我的块处理大致是背靠背的，并且我基本上完全重叠了 D->H 与 H->D 内存副本。这里的内核是微不足道的，所以它在配置文件中只显示为条子。

对于您的情况，您表示不需要 D->H 副本，但它不会增加额外的复杂性，因此我选择展示它。如果您将该行注释到循环之外，仍会发生所需的行为（尽管这会影响稍后的结果检查）。

对这种方法的一个可能的批评是，cudaStreamSynchronize() 调用是必要的，因此我们不会“超出”事件联锁，这意味着循环只会继续进行到 ns 的迭代次数超出当前正在设备上执行。因此，不可能异步启动更多的工作。如果您想立即启动所有工作并继续在 CPU 上执行其他操作，则此方法不会完全允许（当流处理从最后一个迭代达到 ns 迭代时，CPU 将继续通过循环)。

提供代码是为了从概念上说明一种方法。它不保证没有缺陷，我也不声称它适用于任何特定目的。

【讨论】：

谢谢！幸运的是，在所有内核完成之前，我不需要在 CPU 上做任何事情，所以这解决了我的问题。这应该比 cudaDeviceSynchronize() 快吧？

对于我在这里展示的代码，如果我只是在每个循环中使用cudaDeviceSynchronize() 调用，就在内核启动之前，它不会有太大的不同。原因是内核太短了。但是对于持续时间占数据传输成本很大一部分的内核，这种方法应该更好。