这次回顾ECE408 Lecture 22，这次介绍了任务并发。

课程主页：

搬运视频：

序列化数据传输

至此，我们使用cudaMemcpy的方式序列化数据传输和GPU计算：

支持并发

  • 大多数CUDA设备支持设备overlap；
    • 在设备和主机内存之间执行复制的同时执行内核；
int dev_count;
cudaDeviceProp prop;

cudaGetDeviceCount(&dev_count);
for (int i = 0; i < dev_count; i++) {
    cudaGetDeviceProperties(&prop, i);

    if (prop.deviceOverlap) …

重叠（流水线）时序

  • 将大向量分成段；
  • 相邻段的传输和计算overlap；

使用CUDA流和异步Memcpy

  • CUDA支持使用流并行执行内核和cudaMemcpy；
  • 每个流都是一个操作队列（内核启动和 cudaMemcpy 的）；
  • 不同流中的操作（任务）
    • 可以并行执行；
    • 这是任务并行的一个版本；

流

  • 从主机代码发出的设备请求被放入队列；
    • 由驱动程序和设备异步处理的队列，称为“流”；
    • 驱动程序确保队列中的命令严格按顺序处理。内存复制在内核启动之前结束，等等。

  • 为了允许并发复制和内核执行，需要多个队列；

流的概念视图

Multi-Stream Host Code

代码：

cudaStream_t stream0, stream1;
cudaStreamCreate( &stream0);
cudaStreamCreate( &stream1);
float *d_A0, *d_B0, *d_C0; // device memory for stream 0
float *d_A1, *d_B1, *d_C1; // device memory for stream 1

// cudaMalloc for d_A0, d_B0, d_C0, d_A1, d_B1, d_C1 go here

for (int i=0; i<n; i+=SegSize*2) {
    cudaMemcpyAsync(d_A0, h_A+i, SegSize*sizeof(float),.., stream0);
    cudaMemcpyAsync(d_B0, h_B+i, SegSize*sizeof(float),.., stream0);
    vecAdd<<<SegSize/256, 256, 0, stream0>>>(d_A0, d_B0, …);
    cudaMemcpyAsync(h_C+i, d_C+0, SegSize*sizeof(float),.., stream0);
    
    cudaMemcpyAsync(d_A1, h_A+i+SegSize; SegSize*sizeof(float),.., stream1);
    cudaMemcpyAsync(d_B1, h_B+i+SegSize; SegSize*sizeof(float),.., stream1);
    vecAdd<<<SegSize/256, 256, 0, stream1>>>(d_A1, d_B1, …);
    cudaMemcpyAsync(h_C+i+SegSize, d_C1, SegSize*sizeof(float),.., stream1);
}

支持并发的简单队列

队列头部的任务在执行之前等待依赖项（arcs）；例如，内核1等待memcpy A.1和memcpy B.1完成。

不是我们想要的重叠

C.0在复制引擎队列中阻止A.1和B.1：

更好的Multi-Stream Host Code

for (int i=0; i<n; i+=SegSize*2) {
    cudaMemCpyAsync(d_A0, h_A+i; SegSize*sizeof(float),.., stream0);
    cudaMemCpyAsync(d_B0, h_B+i; SegSize*sizeof(float),.., stream0);
    cudaMemCpyAsync(d_A1, h_A+i+SegSize; SegSize*sizeof(float),.., stream1);
    cudaMemCpyAsync(d_B1, h_B+i+SegSize; SegSize*sizeof(float),.., stream1);
    
    vecAdd<<<SegSize/256, 256, 0, stream0>>>(d_A0, d_B0, …);
    vecAdd<<<SegSize/256, 256, 0, stream1>>>(d_A1, d_B1, …);
    cudaMemCpyAsync(d_C0, h_C+I; SegSize*sizeof(float),.., stream0);
    cudaMemCpyAsync(d_C1, h_C+i+SegSize; SegSize*sizeof(float),.., stream1);
}

现在是更好的调度

更好地与两个流重叠

  • C.0不再阻塞复制引擎队列中的A.1和B.1；
  • 然而，C.1仍然阻止A.2和B.2进行下一次迭代；
    • PCIe 仅用于一个方向；

连续流水线需要三个stream

  • 将大向量分成段；
  • 相邻段的重叠传输和计算；

Hyer Queue

  • 为每个引擎提供多个真实流队列；
  • 通过允许一些流为引擎取得进展而其他流被阻止来允许更多的并发性；

开普勒改进的并发性

Kepler允许32路并发；

  • 每个流一个工作队列；
  • 全流级别的并发；
  • 没有流间依赖；

流中的显式同步

  • cudaEventRecord(event, stream)；
    • 在流中开始事件记录；
  • cudaStreamWaitEvent(stream, event)；
    • 等待最近的cudaEventRecord调用完成；
  • cudaStreamSynchronize(stream)；
    • 等待流中所有未决操作完成；

较小的段减少边界效应

段应该有多小？

  • 如果我们overlap；
    • 传输段N的输入；
    • 段N – 1的计算；
    • 以及传输段N – 2的结果；
  • 我们在开始和结束时仍然有不重叠的工作；

那么段应该真的很小吗？

执行时间在理想情况下是线性的

为什么是上图这样？

  • 一些SM空闲；
  • warp太少，无法让SM忙碌；
  • 线程太少，无法填充warp；
  • 块越少负载平衡越差；
  • 内核启动需要时间；

使用适中的段大小和设备查询

数据传输：

  • 对于小尺寸具有相似的非线性；
  • 由于主机和DMA的启动成本；

那么段应该有多小？大小适中。最佳尺寸可能取决于GPU。