CUDA 流并发和 D2H 数据传输重叠答案

【问题标题】：CUDA stream Concurrency and D2H data transfer overlapCUDA 流并发和 D2H 数据传输重叠
【发布时间】：2017-09-11 02:14:45
【问题描述】：

我曾尝试将内核执行与 memcpyasync D2H 重叠，但它不起作用。我有 N 组元素。每组有 64 个元素，可以并行处理。由于数据依赖性，组的处理必须是顺序的。即组i的元素必须在处理组i-1的元素之后处理。处理组中的每个元素会产生一个输出，该输出必须从 GPU 传输到 CPU。为了重叠这种 D2H 数据传输，我将一个组的元素划分为多个块，以便可以使用流重叠给定块上的内核执行和 D2H MemcpyAsync。我使用以下伪代码使用 K 流处理 N 组元素。

 groupId=0;
`while( groupId< N){`

    for(i=0;i<K;++i)

// all stream must wait to complete the kernel execution
 of last stream before starting of the processing of next group

if( groupId!=0)

cudaStreamWaitEvent(stream[K-1],syncEvent,0)
kernelA<<< >>>(----,----,----);
CUDAEventRecord(syncEvent,stream[K-1]);
cudaMemcpyAsync(,,,cudaMemcpyDeviceToHost,stream[i]);
}

groupId++

}

当我使用两个流时，会有一些重叠，而当我增加流的数量时，没有重叠，如下图所示。 Processing of 64 elements using two stream.

Processing of 64 elements using four stream

请解释为什么 D2H 数据传输没有完全重叠。此外，在四个流的情况下，每个流的内核使用 16 个线程块调用，每个线程块大小为 128 个线程。从概念上讲，两个流应该同时执行（每个在一个 SM 上），因为 GPU 上有足够的资源可用。但是，不同流的内核执行没有并发性（图 2）。这种情况下没有并发的原因是什么？

【问题讨论】：

标签： cuda

【解决方案1】：

您的 64 bytes 传输太短，无法与任何内容重叠 - 在 PCIe 2.0 全速（约 6GB/s）下，实际传输大约需要 10 nano秒。在您的屏幕截图比例下，这大约是分析器时间线上像素宽度的 1/1000。条形和间隙的有限宽度完全是由于每次传输（设置等）的开销。

您希望传输以兆字节为单位的传输，以便能够与计算重叠传输。

【讨论】：