cudaEventRecord 的位置和来自不同流的重叠操作

【问题标题】：location of cudaEventRecord and overlapping ops from different streamscudaEventRecord 的位置和来自不同流的重叠操作
【发布时间】：2012-08-16 23:15:03
【问题描述】：

我有两个任务。它们每个都执行复制到设备 (D)、运行内核 (R) 和复制到主机 (H) 操作。我将副本复制到 task2 (D2) 的设备与 task1 (R1) 的运行内核。此外，我将 task2 (R2) 的运行内核与复制到 task1 (H1) 的主机重叠。

我还使用 cudaEventRecord 记录每个任务的 D、R、H ops 的开始和停止时间。

我有 GeForce GT 555M、CUDA 4.1 和 Fedora 16。

我有三种情况：

场景 1：我为每个任务使用一个流。我在操作之前/之后放置开始/停止事件。

场景 2：我为每个任务使用一个流。我将第二个重叠操作的开始事件放在第一个开始之前（即，将 start R1 放在 start D2 之前，并将 start H1 em> 在开始 R2 之前）。

场景 3：我为每个任务使用两个流。我使用 cudaStreamWaitEvents 在这两个流之间进行同步。一个流用于 D 和 H（复制）操作，另一个用于 R op。我在操作之前/之后放置开始/停止事件。

Scenario1 无法重叠操作（D2-R1 和 R2-H1 都不能重叠），而 Scenario2 和 Scenario3 成功。 我的问题是：为什么 Scenerio1 失败了，而其他的成功了？

对于每个场景，我都会测量执行 Task1 和 Task2 的总时间。运行 R1 和 R2 分别需要 5 ms。由于 Scenario1 无法重叠 ops，因此总体时间比 Scenario 2 和 3 多 10ms。

以下是场景的伪代码：

场景 1（失败）：将 stream1 用于 task1，将 stream2 用于 task2

start overall 

start D1 on stream1 
D1 on stream1
stop D1 on stream1 

start D2 on stream2
D2 on stream2
stop D2 on stream2

start R1 on stream1
R1 on stream1
stop R1 on stream1

start R2 on stream2
R2 on stream2
stop R2 on stream2

start H1 on stream1
H1 on stream1
stop H1 on stream1

start H2 on stream2
H2 on stream2
stop H2 on stream2

stop overall

Scenario2 (SUCCEEDS)： task1 使用 stream1，task2 使用 stream2，上移第二个重叠操作的 start 事件。 

start overall

start D1 on stream1
D1 on stream1
stop D1 on stream1 

start R1 on stream1 //moved-up

start D2 on stream2
D2 on stream2
stop D2 on stream2

R1 on stream1
stop R1 on stream1

start H1 on stream1 //moved-up

start R2 on stream2
R2 on stream2
stop R2 on stream2

H1 on stream1
stop H1 on stream1

start H2 on stream2
H2 on stream2
stop H2 on stream2

stop overall

场景 3（成功）：将 stream1 和 3 用于 task1，将 stream2 和 4 用于 task2

start overall

start D1 on stream1
D1 on stream1
stop D1 on stream1 

start D2 on stream2
D2 on stream2
stop D2 on stream2

start R1 on stream3
R1 on stream3
stop R1 on stream3

start R2 on stream4
R2 on stream4
stop R2 on stream4

start H1 on stream1
H1 on stream1
stop H1 on stream1

start H2 on stream2
H2 on stream2
stop H2 on stream2

stop overall

以下是所有场景的总体时间信息： 场景 1 = 39.390240 方案 2 = 29.190241 场景3 = 29.298208

我还附上了下面的CUDA代码：

#include <stdio.h>
#include <cuda_runtime.h>
#include <sys/time.h>

__global__ void VecAdd(const float* A, const float* B, float* C, int N)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N)
        {
        C[i] = A[i] + B[N-i];
        C[i] = A[i] + B[i] * 2;
        C[i] = A[i] + B[i] * 3;
        C[i] = A[i] + B[i] * 4;
        C[i] = A[i] + B[i];
        }
}

void overlap()
{

float* h_A;
float *d_A, *d_C;
float* h_A2;
float *d_A2, *d_C2;

int N = 10000000;
size_t size = N * sizeof(float); 

cudaMallocHost((void**) &h_A, size);
cudaMallocHost((void**) &h_A2, size);

// Allocate vector in device memory
cudaMalloc((void**)&d_A, size);
cudaMalloc((void**)&d_C, size);
cudaMalloc((void**)&d_A2, size);
cudaMalloc((void**)&d_C2, size);

float fTimCpyDev1, fTimKer1, fTimCpyHst1, fTimCpyDev2, fTimKer2, fTimCpyHst2;
float fTimOverall3, fTimOverall1, fTimOverall2;

for (int i = 0; i<N; ++i)
    {
    h_A[i] = 1;
    h_A2[i] = 5;
    }

int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

cudaStream_t csStream1, csStream2, csStream3, csStream4;
cudaStreamCreate(&csStream1);
cudaStreamCreate(&csStream2);
cudaStreamCreate(&csStream3);
cudaStreamCreate(&csStream4);

cudaEvent_t ceEvStart, ceEvStop; 
cudaEventCreate( &ceEvStart );
cudaEventCreate( &ceEvStop );

cudaEvent_t ceEvStartCpyDev1, ceEvStopCpyDev1, ceEvStartKer1, ceEvStopKer1, ceEvStartCpyHst1, ceEvStopCpyHst1;
cudaEventCreate( &ceEvStartCpyDev1 );
cudaEventCreate( &ceEvStopCpyDev1 );
cudaEventCreate( &ceEvStartKer1 );
cudaEventCreate( &ceEvStopKer1 );
cudaEventCreate( &ceEvStartCpyHst1 );
cudaEventCreate( &ceEvStopCpyHst1 );
cudaEvent_t ceEvStartCpyDev2, ceEvStopCpyDev2, ceEvStartKer2, ceEvStopKer2, ceEvStartCpyHst2, ceEvStopCpyHst2; 
cudaEventCreate( &ceEvStartCpyDev2 );
cudaEventCreate( &ceEvStopCpyDev2 );
cudaEventCreate( &ceEvStartKer2 );
cudaEventCreate( &ceEvStopKer2 );
cudaEventCreate( &ceEvStartCpyHst2 );
cudaEventCreate( &ceEvStopCpyHst2 );


//Scenario1

cudaDeviceSynchronize();

cudaEventRecord(ceEvStart, 0);

cudaEventRecord(ceEvStartCpyDev1, csStream1);
cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, csStream1);
cudaEventRecord(ceEvStopCpyDev1, csStream1);

cudaEventRecord(ceEvStartCpyDev2, csStream2);
cudaMemcpyAsync(d_A2, h_A2, size, cudaMemcpyHostToDevice, csStream2);
cudaEventRecord(ceEvStopCpyDev2, csStream2);

cudaEventRecord(ceEvStartKer1, csStream1); 
VecAdd<<<blocksPerGrid, threadsPerBlock, 0, csStream1>>>(d_A, d_A, d_C, N);
cudaEventRecord(ceEvStopKer1, csStream1); 

cudaEventRecord(ceEvStartKer2, csStream2); 
VecAdd<<<blocksPerGrid, threadsPerBlock, 0, csStream2>>>(d_A2, d_A2, d_C2, N);
cudaEventRecord(ceEvStopKer2, csStream2);

cudaEventRecord(ceEvStartCpyHst1, csStream1);
cudaMemcpyAsync(h_A, d_C, size, cudaMemcpyDeviceToHost, csStream1);
cudaEventRecord(ceEvStopCpyHst1, csStream1);

cudaEventRecord(ceEvStartCpyHst2, csStream2);
cudaMemcpyAsync(h_A2, d_C2, size, cudaMemcpyDeviceToHost, csStream2);
cudaEventRecord(ceEvStopCpyHst2, csStream2);

cudaEventRecord(ceEvStop, 0);
cudaDeviceSynchronize();

cudaEventElapsedTime( &fTimOverall1, ceEvStart, ceEvStop);
printf("Scenario1 overall time= %10f\n", fTimOverall1);


//Scenario2 

cudaDeviceSynchronize();

cudaEventRecord(ceEvStart, 0);

cudaEventRecord(ceEvStartCpyDev1, csStream1);
cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, csStream1);
cudaEventRecord(ceEvStopCpyDev1, csStream1);

cudaEventRecord(ceEvStartKer1, csStream1); //moved up 

cudaEventRecord(ceEvStartCpyDev2, csStream2);
cudaMemcpyAsync(d_A2, h_A2, size, cudaMemcpyHostToDevice, csStream2);
cudaEventRecord(ceEvStopCpyDev2, csStream2);

VecAdd<<<blocksPerGrid, threadsPerBlock, 0, csStream1>>>(d_A, d_A, d_C, N);
cudaEventRecord(ceEvStopKer1, csStream1); 

cudaEventRecord(ceEvStartCpyHst1, csStream1); //moved up

cudaEventRecord(ceEvStartKer2, csStream2); 
VecAdd<<<blocksPerGrid, threadsPerBlock, 0, csStream2>>>(d_A2, d_A2, d_C2, N);
cudaEventRecord(ceEvStopKer2, csStream2);

cudaMemcpyAsync(h_A, d_C, size, cudaMemcpyDeviceToHost, csStream1);
cudaEventRecord(ceEvStopCpyHst1, csStream1);

cudaEventRecord(ceEvStartCpyHst2, csStream2);
cudaMemcpyAsync(h_A2, d_C2, size, cudaMemcpyDeviceToHost, csStream2);
cudaEventRecord(ceEvStopCpyHst2, csStream2);

cudaEventRecord(ceEvStop, 0);
cudaDeviceSynchronize();


cudaEventElapsedTime( &fTimOverall2, ceEvStart, ceEvStop);
printf("Scenario2 overall time= %10f\n", fTimOverall2);

//Scenario3
cudaDeviceSynchronize();

cudaEventRecord(ceEvStart, 0);

cudaEventRecord(ceEvStartCpyDev1, csStream1);
cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, csStream1);
cudaEventRecord(ceEvStopCpyDev1, csStream1);

cudaEventRecord(ceEvStartCpyDev2, csStream2);
cudaMemcpyAsync(d_A2, h_A2, size, cudaMemcpyHostToDevice, csStream2);
cudaEventRecord(ceEvStopCpyDev2, csStream2);

cudaStreamWaitEvent(csStream3, ceEvStopCpyDev1, 0);
cudaEventRecord(ceEvStartKer1, csStream3); 
VecAdd<<<blocksPerGrid, threadsPerBlock, 0, csStream3>>>(d_A, d_A, d_C, N);
cudaEventRecord(ceEvStopKer1, csStream3);

cudaStreamWaitEvent(csStream4, ceEvStopCpyDev2, 0);
cudaEventRecord(ceEvStartKer2, csStream4); 
VecAdd<<<blocksPerGrid, threadsPerBlock, 0, csStream4>>>(d_A2, d_A2, d_C2, N);
cudaEventRecord(ceEvStopKer2, csStream4);

cudaStreamWaitEvent(csStream1, ceEvStopKer1, 0);
cudaEventRecord(ceEvStartCpyHst1, csStream1);
cudaMemcpyAsync(h_A, d_C, size, cudaMemcpyDeviceToHost, csStream1);
cudaEventRecord(ceEvStopCpyHst1, csStream1);

cudaStreamWaitEvent(csStream2, ceEvStopKer2, 0);
cudaEventRecord(ceEvStartCpyHst2, csStream2);
cudaMemcpyAsync(h_A2, d_C2, size, cudaMemcpyDeviceToHost, csStream2);
cudaEventRecord(ceEvStopCpyHst2, csStream2);

cudaEventRecord(ceEvStop, 0);
cudaDeviceSynchronize();

cudaEventElapsedTime( &fTimOverall3, ceEvStart, ceEvStop);
printf("Scenario3 overall time = %10f\n", fTimOverall3);

cudaStreamDestroy(csStream1);
cudaStreamDestroy(csStream2);
cudaStreamDestroy(csStream3);
cudaStreamDestroy(csStream4);

cudaFree(d_A);
cudaFree(d_C);
cudaFreeHost(h_A);
cudaFree(d_A2);
cudaFree(d_C2);
cudaFreeHost(h_A2);

}

int main()
{

  overlap();
}

非常感谢您提前抽出时间！

【问题讨论】：

    标签： cuda gpu nvidia


    【解决方案1】：

    （注意，我对Tesla系列设备比较熟悉，实际上并没有GT 555M可以试验，所以我的结果特指C2070。我不知道555m有多少复制引擎有，但我希望下面描述的问题是导致您看到的行为的原因。）

    问题是鲜为人知的事实，即 cudaEventRecords 也是 CUDA 操作，并且它们还必须在启动/执行之前放置在硬件队列之一中。 （一个复杂的因素是，由于 cudaEventRecord 既不是复制操作，也不是计算内核，它实际上可以进入任何硬件队列。我的理解是它们通常与同一流的前面 CUDA 操作进入相同的硬件队列，但由于文档中未指定，实际操作可能取决于设备/驱动程序。）

    如果我可以扩展您的符号以将“E”用于“事件记录”，并详细说明如何填充硬件队列（类似于“CUDA C/C++ Streams and Concurrency”网络研讨会中所做的），那么在您的场景 1 示例中，你有：

    Issue order for CUDA operations:
   ED1
   D1
   ED1
   ED2
   D2
   ED2
   ER1
   R1
   ER1
   ...

    这些填充队列如下：

    Hardware Queues:    copyH2D     Kernel
                    -------     ------
                    ED1       * R1
                    D1       /  ER1
                    ED1     /   ...
                    ED2    /
                    D2    /
                    ED2  /
                    ER1 *

    您可以看到，由于位于流 1 中，R1 在 ER1 完成之前不会执行，这在 D1 和 D2 都完成之前不会发生，因为它们都在 H2D 复制队列中序列化。

    通过在场景 2 中向上移动 cudaEventRecord ER1，您可以避免这种情况，因为流 1 中的所有 CUDA 操作在 R1 之前，在 D2 之前完成。这允许 R1 与 D2 同时启动。

    Hardware Queues:    copyH2D     Kernel
                    -------     ------
                    ED1      *  R1
                    D1      /   ER1
                    ED1    /    ...
                    ER1   *
                    ED2    
                    D2    
                    ED2

    在您的场景 3 中，ER1 被替换为 ER3。由于这是流 3 中的第一个操作，它可以去任何地方，并且（猜测）进入内核或复制 D2H 队列，它可以立即启动，（如果你没有 

    cudaStreamWaitEvent(csStream3, ceEvStopCpyDev1, 0);

    用于与流 1) 同步，因此不会导致与 D2 的错误序列化。

    Hardware Queues:    copyH2D     Kernel
                    -------     ------
                    ED1     *   ER3
                    D1     /    R3
                    ED1   *     ER3
                    ED2         ...
                    D2    
                    ED2

    我的 cmets 会是

    1. 在考虑并发性时，CUDA 操作的问题顺序非常重要
    2. cudaEventRecord 和类似操作会像其他所有操作一样被放置在硬件队列中，并可能导致错误的序列化。没有很好地描述它们如何被放置在硬件队列中，并且可能取决于设备/驱动程序。因此，为了获得最佳并发性，应将 cudaEventRecord 和类似操作的使用减少到必要的最低限度。
    3. 如果内核需要为性能研究计时，可以使用事件来完成，但会破坏并发性。这对于开发来说很好，但对于生产代码应该避免。

    但是您应该注意到，即将推出的 Kepler GK110 (Tesla K20) 设备通过使用 32 个硬件队列在减少错误序列化方面做出了重大改进。详情请参阅GK110 Whitepaper（第 17 页）。

    希望这会有所帮助。

    【讨论】：

      猜你喜欢
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 2021-11-06
      • 2017-11-23
      • 2015-02-25
      • 2018-08-25
      相关资源
      最近更新 更多
      热门标签
      Java Python linux javascript Mysql C# Docker 算法 前端 SpringBoot Redis Vue spring 设计模式 .net core .net kubernetes c++ 数据库 数据结构 大数据 js 机器学习 微服务 Android Go 程序员 面试 JVM ASP.net core 云原生 人工智能 后端 PHP git CSS golang k8s Nginx Django mybatis 深度学习 多线程 React 架构 devops 爬虫 云计算 Spring Boot LeetCode