小贝子编程

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



我有两个任务。它们中的每一个都执行复制到设备 （D）、运行内核 （R） 和复制到主机 （H） 操作。我正在将复制到任务 2 （D2） 的设备与任务 1 （R1） 的运行内核重叠。此外，我将任务 2 （R2） 的运行内核与复制到任务 1 （H1） 的主机重叠。

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

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

我有三种情况：

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

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

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

方案

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

对于每个方案，我测量执行任务 1 和任务 2 的总时间。运行 R1 和 R2 各需要 5 毫秒。由于方案 1 无法重叠操作，因此总时间比方案 2 和方案 3 多 10 毫秒。

下面是方案的伪代码：

方案 1（失败）：将流 1 用于任务

1，将流 2 用于任务 2

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
场景 2（成功）：将流 1 用于任务 1，将流 2

用于任务 2，上移重叠操作的第二个的开始事件。 

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（

成功）：对任务 1 使用流 1 和 3，对任务 2 使用流 2 和 4

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= %10fn", 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= %10fn", 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 = %10fn", 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();
}

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

（注意，我对特斯拉系列设备更熟悉，实际上没有GT 555M可以试验，所以我的结果特别提到了C2070。 我不知道 555m 有多少个复制引擎，但我希望下面描述的问题会导致您看到的行为。

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

如果我可以扩展您的符号以使用"E"表示"事件记录"，并详细说明硬件队列的填充方式（类似于"CUDA C/C++ 流和并发"网络研讨会中所做的操作），那么在您的场景 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 *

您可以看到，R1 由于位于流 1 中，在 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

我的评论是

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

但是，您应该注意，即将推出的开普勒GK110（Tesla K20）设备通过使用32个硬件队列在减少错误序列化方面进行了重大改进。 有关详细信息，请参阅 GK110 白皮书（第 17 页）。

希望这有帮助。

相关内容

最新更新


热门标签：

javascript python java c# php android html jquery c++ css ios sql mysql arrays asp.net json python-3.x ruby-on-rails .net sql-server django objective-c excel regex ruby linux ajax iphone xml vba spring asp.net-mvc database wordpress string postgresql wpf windows xcode bash git oracle list vb.net multithreading eclipse algorithm macos powershell visual-studio image forms numpy scala function api selenium