0

私には2つのタスクがあります。それぞれがデバイスへのコピー（D）、カーネル（R）の実行、およびホストへのコピー（H）の操作を実行します。task2（D2）のデバイスへのコピーをtask1（R1）の実行カーネルとオーバーラップしています。さらに、task2（R2）の実行カーネルとtask1（H1）のホストへのコピーをオーバーラップしています。

また、cudaEventRecordを使用して、各タスクのD、R、Hopsの開始時間と停止時間を記録します。

私はGeForceGT555M、CUDA 4.1、およびFedora16を持っています。

私には3つのシナリオがあります：

シナリオ1：タスクごとに1つのストリームを使用します。手術の直前/直後にスタート/ストップイベントを配置します。

シナリオ2：タスクごとに1つのストリームを使用します。重複する操作の2番目の開始イベントを最初の操作の開始前に配置します（つまり、開始R1を開始D2の前に配置し、開始H1を開始R2の前に配置します）。

シナリオ3：タスクごとに2つのストリームを使用します。cudaStreamWaitEventsを使用して、これら2つのストリームを同期します。1つのストリームはDおよびH（コピー）操作に使用され、もう1つのストリームはR操作に使用されます。手術の直前/直後にスタート/ストップイベントを配置します。

シナリオ1はopsのオーバーラップに失敗しますが（D2-R1もR2-H1もオーバーラップできません）、シナリオ2とシナリオ3は成功します。そして私の質問は、Scenerio1が失敗するのに、他のScenerio1は成功するのはなぜですか？

シナリオごとに、タスク1とタスク2を実行するための全体的な時間を測定します。R1とR2の両方を実行するには、それぞれ5ミリ秒かかります。シナリオ1は操作のオーバーラップに失敗するため、全体の時間はシナリオ2および3よりも10ミリ秒長くなります。

シナリオの擬似コードは次のとおりです。

シナリオ1（失敗）：タスク1にstream1を使用し、タスク2にstream2を使用します

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にstream1を使用し、タスク2にstream2を使用し、重複する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にstream1と3を使用し、タスク2にstream2と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= %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();
}

よろしくお願いします！

4

1 に答える 1

0

（注：私はTeslaシリーズのデバイスに精通しており、実際に実験するGT 555Mを持っていないため、結果は特にC2070を参照しています。555mのコピーエンジンの数はわかりませんが、以下に説明する問題が、表示されている動作の原因であると思います。）

問題は、cudaEventRecordsもCUDA操作であり、起動/実行する前にハードウェアキューの1つに配置する必要があるというあまり知られていない事実です。（複雑な要因は、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 *

また、ストリーム1にあるため、R1はER1が完了するまで実行されません。これは、D1とD2の両方がH2Dコピーキューでシリアル化されるため、完了するまで実行されません。

シナリオ2でcudaEventRecordER1を上に移動すると、R1より前のストリーム1のすべてのCUDA操作が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. パフォーマンス調査のためにカーネルのタイミングを調整する必要がある場合は、イベントを使用して実行できますが、同時実行性が失われます。これは開発には問題ありませんが、本番コードでは避ける必要があります。

ただし、今後のKepler GK110（Tesla K20）デバイスでは、32個のハードウェアキューを使用することで、誤ったシリアル化を大幅に削減できることに注意してください。詳細については、GK110ホワイトペーパー（17ページ）を参照してください。

お役に立てれば。

于 2012-08-28T18:09:38.763 に答える