Accelerated Ray Tracing （六）

https://developer.nvidia.com/blog/how-overlap-data-transfers-cuda-cc/

重叠的数据传输

CUDA Streams

CUDA中的stream 是按主机代码发出的顺序在设备上执行的一系列操作。尽管保证流中的操作按规定的顺序执行，但是可以交错不同流中的操作，并且在可能的情况下甚至可以同时运行。

The default stream

CUDA中的所有设备操作（kernels 和数据传输）都在流中运行。如果未指定任何流，则使用默认流（也称为“null stream”）。默认流与其他流不同，因为它是关于设备上的操作的同步流：直到设备上其他流中的所有先前发布的操作完成为止，默认流中的任何操作才会开始，并且默认流中的操作必须先完成其他操作（设备上其他流中的其他操作）才能开始。

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

default stream, 从设备的角度看，所有这三个操作都被发布到相同的（默认）流中，并将按照它们发出的顺序执行。从主机的角度来看，隐式数据传输是阻塞传输或同步传输，而内核启动是异步的。由于第一行上的主机到设备数据传输是同步的，因此在主机到设备的传输完成之前，CPU线程不会到达第二行上的内核调用。发出内核后，CPU线程移至第三行，但是由于设备端的执行顺序，该行的传输无法开始。

从主机的角度来看，kernel 启动的asynchronous 行为使重叠的设备和主机计算变得非常简单。我们可以修改代码以添加一些独立的CPU计算，如下所示。

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
myCpuFunction(b)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

在上面的代码中，一旦在设备上启动了increment（）内核，CPU线程就会执行myCpuFunction（），将其在CPU上的执行与在GPU上的内核执行重叠。host function 或device kernel是先完成的，这不会影响随后的设备到主机的传输，只有在内核完成后才开始。从设备的角度来看，与前面的示例相比，没有任何变化。设备完全不知道myCpuFunction（）。

Non-default streams

CUDA C / C ++中的非默认流在主机代码中声明，创建和销毁，如下所示。

cudaStream_t stream1;
cudaError_t result;
result = cudaStreamCreate(&stream1)
result = cudaStreamDestroy(stream1)

要将数据传输到非默认流，我们使用cudaMemcpyAsync（）函数，该函数类似于上一篇文章中讨论的cudaMemcpy（）函数，但是将流标识符用作第五个参数。

result = cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1)

cudaMemcpyAsync（）在主机上是非阻塞的，因此控制在发出传输后立即返回到主机线程。此例程有cudaMemcpy2DAsync（）和cudaMemcpy3DAsync（）变体，可以在指定的流中异步传输2D和3D数组节。

要将kernel 发布给非默认流，我们将流标识符指定为第四个执行配置参数（第三个执行配置参数分配共享设备内存，我们将在后面讨论；现在使用0）。

increment<<<1,N,0,stream1>>>(d_a)

Synchronization with streams

由于非默认流中的所有操作都相对于host 代码是非阻塞的，因此您将遇到需要将host 代码与流中的操作同步的情况。有几种方法可以做到这一点。“heavy hammer”方法是使用cudaDeviceSynchronize（），它将阻塞host 代码，直到设备上所有先前发布的操作都完成为止。在大多数情况下，这太过分了，并且由于使整个设备和主机线程停止运行，实际上会损害性能。

CUDA stream  API具有多种较不严格的方法来将主机与流同步。函数cudaStreamSynchronize（stream）可用于阻止主机线程，直到指定流中所有以前发出的操作完成为止。函数cudaStreamQuery（stream）测试是否已完成向指定流发出的所有操作，而不会阻止主机执行。函数cudaEventSynchronize（event）和cudaEventQuery（event）的功能类似于其流对应项，不同之处在于它们的结果基于是否已记录指定的事件而不是指定的流是否空闲。您还可以使用cudaStreamWaitEvent（event）对特定事件在单个流中进行同步操作（即使该事件记录在其他流中或在其他设备上！）。

Overlapping Kernel Execution and Data Transfers

•  该设备必须能够“concurrent copy and execution”。可以从cudaDeviceProp结构的deviceOverlap字段中查询，也可以从CUDA SDK / Toolkit附带的deviceQuery示例的输出中查询。几乎所有具有1.1和更高计算能力的设备都具有此功能。 内核执行和要重叠的数据传输必须都发生在不同的非默认流中。
• 内核执行和要重叠的数据传输必须都发生在不同的非默认流中。
• 数据传输中涉及的主机存储器必须是固定存储器。

#include <stdio.h>

// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
    if (result != cudaSuccess) {
        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
#endif
    return result;
}

__global__ void kernel(float* a, int offset)
{
    int i = offset + threadIdx.x + blockIdx.x * blockDim.x;
    float x = (float)i;
    float s = sinf(x);
    float c = cosf(x);
    a[i] = a[i] + sqrtf(s * s + c * c);
}

float maxError(float* a, int n)
{
    float maxE = 0;
    for (int i = 0; i < n; i++) {
        float error = fabs(a[i] - 1.0f);
        if (error > maxE) maxE = error;
    }
    return maxE;
}

int main(int argc, char** argv)
{
    const int blockSize = 256, nStreams = 4;//（非默认）流数
    const int n = 4 * 1024 * blockSize * nStreams;
    const int streamSize = n / nStreams;
    const int streamBytes = streamSize * sizeof(float);
    const int bytes = n * sizeof(float);

    int devId = 0;
    if (argc > 1) devId = atoi(argv[1]);

    cudaDeviceProp prop;
    checkCuda(cudaGetDeviceProperties(&prop, devId));
    printf("Device : %s\n", prop.name);
    checkCuda(cudaSetDevice(devId));

    // allocate pinned host memory and device memory
    float* a, * d_a;
    checkCuda(cudaMallocHost((void**)&a, bytes));      // host pinned
    checkCuda(cudaMalloc((void**)&d_a, bytes)); // device

    float ms; // elapsed time in milliseconds

    // create events and streams
    cudaEvent_t startEvent, stopEvent, dummyEvent;
    cudaStream_t stream[nStreams];
    checkCuda(cudaEventCreate(&startEvent));
    checkCuda(cudaEventCreate(&stopEvent));
    checkCuda(cudaEventCreate(&dummyEvent));
    for (int i = 0; i < nStreams; ++i)
        checkCuda(cudaStreamCreate(&stream[i]));

    // baseline case - sequential transfer and execute
    memset(a, 0, bytes);
    checkCuda(cudaEventRecord(startEvent, 0));
    checkCuda(cudaMemcpy(d_a, a, bytes, cudaMemcpyHostToDevice));
    kernel << <n / blockSize, blockSize >> > (d_a, 0);
    checkCuda(cudaMemcpy(a, d_a, bytes, cudaMemcpyDeviceToHost));
    checkCuda(cudaEventRecord(stopEvent, 0));
    checkCuda(cudaEventSynchronize(stopEvent));
    checkCuda(cudaEventElapsedTime(&ms, startEvent, stopEvent));
    printf("Time for sequential transfer and execute (ms): %f\n", ms);
    printf("  max error: %e\n", maxError(a, n));

    // asynchronous version 1: loop over {copy, kernel, copy}
    memset(a, 0, bytes);
    checkCuda(cudaEventRecord(startEvent, 0));
    for (int i = 0; i < nStreams; ++i) {//nStreams=N/streamSize 遍历数组每个块的所有操作。
        int offset = i * streamSize;
        checkCuda(cudaMemcpyAsync(&d_a[offset], &a[offset],
            streamBytes, cudaMemcpyHostToDevice,
            stream[i]));
        kernel << <streamSize / blockSize, blockSize, 0, stream[i] >> > (d_a, offset);//将大小为N的数组拆分为streamSize元素的块。由于内核在所有元素上独立运行，因此每个块都可以独立处理。
        checkCuda(cudaMemcpyAsync(&a[offset], &d_a[offset],
            streamBytes, cudaMemcpyDeviceToHost,
            stream[i]));
    }

/*
//另一种方法是将类似的操作分批处理，首先发出所有主机到设备的传输，然后发出所有内核启动，然后发出所有设备到主机的传输，如以下代码所示。
for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&d_a[offset], &a[offset], 
                  streamBytes, cudaMemcpyHostToDevice, cudaMemcpyHostToDevice, stream[i]);
}

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
}

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&a[offset], &d_a[offset], 
                  streamBytes, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToHost, stream[i]);
}
*/
//两种异步方法都能产生正确的结果，并且在两种情况下，依存操作都按照需要执行的顺序发给相同的流。
    checkCuda(cudaEventRecord(stopEvent, 0));
    checkCuda(cudaEventSynchronize(stopEvent));
    checkCuda(cudaEventElapsedTime(&ms, startEvent, stopEvent));
    printf("Time for asynchronous V1 transfer and execute (ms): %f\n", ms);
    printf("  max error: %e\n", maxError(a, n));

    // asynchronous version 2: 
    // loop over copy, loop over kernel, loop over copy
    memset(a, 0, bytes);
    checkCuda(cudaEventRecord(startEvent, 0));
    for (int i = 0; i < nStreams; ++i)
    {
        int offset = i * streamSize;
        checkCuda(cudaMemcpyAsync(&d_a[offset], &a[offset],
            streamBytes, cudaMemcpyHostToDevice,
            stream[i]));
    }
    for (int i = 0; i < nStreams; ++i)
    {
        int offset = i * streamSize;
        kernel << <streamSize / blockSize, blockSize, 0, stream[i] >> > (d_a, offset);
    }
    for (int i = 0; i < nStreams; ++i)
    {
        int offset = i * streamSize;
        checkCuda(cudaMemcpyAsync(&a[offset], &d_a[offset],
            streamBytes, cudaMemcpyDeviceToHost,
            stream[i]));
    }
    checkCuda(cudaEventRecord(stopEvent, 0));
    checkCuda(cudaEventSynchronize(stopEvent));
    checkCuda(cudaEventElapsedTime(&ms, startEvent, stopEvent));
    printf("Time for asynchronous V2 transfer and execute (ms): %f\n", ms);
    printf("  max error: %e\n", maxError(a, n));

    // cleanup
    checkCuda(cudaEventDestroy(startEvent));
    checkCuda(cudaEventDestroy(stopEvent));
    checkCuda(cudaEventDestroy(dummyEvent));
    for (int i = 0; i < nStreams; ++i)
        checkCuda(cudaStreamDestroy(stream[i]));
    cudaFree(d_a);
    cudaFreeHost(a);

    return 0;
}

 

Device : Tesla C1060

Time for sequential transfer and execute (ms ): 12.92381
  max error : 2.3841858E -07
Time for asynchronous V1 transfer and execute (ms ): 13.63690 
  max error : 2.3841858E -07
Time for asynchronous V2 transfer and execute (ms ): 8.84588
  max error : 2.3841858E -07

第一个time的是顺序传输和使用blocking 传输的内核执行，我们将其用作异步加速比较的基准。

CUDA设备包含用于各种任务的引擎，这些引擎在发布操作时将操作排队。

 在原理图中，我们假设主机到设备的传输，内核执行和设备到主机的传输所需的时间大致相同（选择内核代码是为了实现这一点）。如顺序内核所预期的那样，任何操作都没有重叠。对于我们代码的第一个异步版本，复制引擎中的执行顺序为：H2D stream（1），D2H stream（1），H2D stream（2），D2H stream（2）等。这就是为什么在C1060上使用第一个异步版本时我们看不到任何加速的原因：任务以防止内核执行和数据传输重叠的顺序发布给复制引擎。但是，对于第二版，所有主机到设备的传输都在任何设备到主机的传输之前发出，因此重叠是可能的，这表明执行时间较短。从我们的示意图中，我们期望异步版本2的执行是顺序版本的8/12，即8.7毫秒，这在先前给出的时序结果中得到了证实。

在C2050上，两个特性相互作用导致了与C1060的行为差异。C2050有两个复制引擎，一个用于主机到设备的传输，另一个用于设备到主机的传输，还有一个单独的内核引擎。下图演示了我们的示例在C2050上的执行。

有两个复制引擎说明了为什么异步版本1在C2050上可以实现良好的加速：stream[i]中从设备到主机的数据传输不会像C1060上那样阻塞stream[i + 1]中的主机到设备数据传输，因为在每个复制方向上都有一个单独的引擎C2050。原理图预测执行时间相对于顺序版本将减少一半，这大致就是我们的时序结果所示。

但是，在C2050的异步版本2中观察到的性能下降如何呢？这与C2050可以同时运行多个内核的能力有关。当多个内核以不同的(非默认)流背对背发出时，调度器尝试启用这些内核的并发执行，并因此延迟通常在每个内核完成后发生的信号(负责启动设备到主机的传输)，直到所有内核完成。因此，尽管在第二版异步代码中主机到设备的传输和内核执行之间存在重叠，但内核执行和设备到主机的传输之间没有重叠。该示意图预测异步版本2的总时间为顺序版本的9/12的时间，即7.5毫秒，这由我们的计时结果证实。