私には2つのタスクがあります。どちらもデバイスへのコピー（D）を実行し、カーネル（R）操作を実行します。タスクには異なるカーネルランタイムがあります。R1の完了にはR2の5倍の時間がかかります（R1 =〜17ミリ秒、R2 =〜3.5ミリ秒）。タスクのカーネルは待機操作を実行し、これらのカーネルを同時に実行できるようにします。各コピー操作には7ミリ秒かかります。
私はGeForceGT555M、CUDA 4.1、およびFedora16を持っています。
cudaEventRecordを使用して、各タスクのDおよびR操作の開始時刻と停止時刻を記録します。タスクごとに2つのストリームを使用します。cudaStreamWaitEventsを使用して、タスクのこれら2つのストリーム間で同期します。1つのストリームはタスクのDocに使用され、もう1つのストリームはタスクのRopに使用されます。私の目標は、D2をR1とオーバーラップさせることです。task1とtask2の全体的なタイミングを測定して、このオーバーラップが達成されているかどうかを判断します。
2つのシナリオがあります。Scenerio1では、「start R1」はカーネルの前に配置され、「startR2」はカーネルの間に配置されます。Scenerio2では、「startR1」と「startR2」の両方がカーネルの前に配置されます。
以下の特定の擬似コードでは、Scenario1とScenerio2は同じように動作しません。Scenerio2はD2とR1のオーバーラップに失敗しますが、Scenerio1はオーバーラップに成功します。だから私の質問は：R2がR1より短いのに、なぜD2をR1とオーバーラップさせるために、カーネル間に「開始R2」を（前ではなく）配置する必要があるのですか？（R1がR2より短いシナリオもテストしたことに注意してください。この場合、カーネルの前または間に「start R2」を配置しても違いはなく、どちらの場合も、D2をR1とオーバーラップさせることができます。 D2が完了したら、R1とR2を同時に実行することもできます。）
シナリオ1と2の擬似コードは次のとおりです（タスク1にはstream1とstream3を使用し、タスク2にはstream2とstream4を使用します）。
シナリオ1（成功）：
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 //longer
start R2 on stream4 // start R2 is in between kernels
R2 on stream4 //shorter
stop R2 on stream4
stop R1 on stream3
stop overall
シナリオ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 stream3
start R2 on stream4 // start R2 is before kernels
R1 on stream3 //longer
R2 on stream4 //shorter
stop R2 on stream4
stop R1 on stream3
stop overall
シナリオの全体的なタイミングは次のとおりです。
シナリオ1=24.109312
シナリオ2=31.194496
これらのシナリオで予想される全体的な実行時間は、D1 + R1 = 7 + 17 = 24です（D2をR1とオーバーラップさせ、同時にR1とR2を同時に実行できます）。シナリオ1はこのランタイムの達成に成功しましたが、Scenerio2はそれを達成できませんでした。これは、シナリオ2がD2とR1をオーバーラップできないためです。（D2には7ミリ秒かかるため、Scenario2の実行時間は24 + 7 = 31です）。
以下にCUDAコードも添付しました。
#include <stdio.h>
#include <cuda_runtime.h>
#include <sys/time.h>
__global__ void wait_k(long time_clocks)
{
long start_clock = clock();
long clock_offset = 0;
while( clock_offset < time_clocks) {
clock_offset = clock() - start_clock;
}
}
void shorterR2_D2_R1_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);
for (int i = 0; i<N; ++i)
{
h_A[i] = 1;
h_A2[i] = 5;
}
cudaStream_t csStream1, csStream2, csStream3, csStream4;
cudaStreamCreate(&csStream1);
cudaStreamCreate(&csStream2);
cudaStreamCreate(&csStream3);
cudaStreamCreate(&csStream4);
//allocate vars for dummy copy
float* h_pfDummy;
float* d_pfDummy;
size_t iMemSz = 10 * sizeof(float);
cudaMallocHost((void**) &h_pfDummy, iMemSz);
cudaMalloc((void**)&d_pfDummy, iMemSz);
cudaMemcpyAsync(d_pfDummy, h_pfDummy, iMemSz, cudaMemcpyHostToDevice, csStream1);
cudaMemcpyAsync(d_pfDummy, h_pfDummy, iMemSz, cudaMemcpyHostToDevice, csStream2);
//delete vars of dummy copy
cudaFree(d_pfDummy);
cudaFreeHost(h_pfDummy);
long time_clocks = 20000000;
long div = 5;
cudaEvent_t ceEvStart, ceEvStop;
cudaEventCreate( &ceEvStart );
cudaEventCreate( &ceEvStop );
//diff stream time events
cudaEvent_t ceEvStartCpyDev1, ceEvStopCpyDev1, ceEvStartKer1, ceEvStopKer1;
cudaEventCreate( &ceEvStartCpyDev1 );
cudaEventCreate( &ceEvStopCpyDev1 );
cudaEventCreate( &ceEvStartKer1 );
cudaEventCreate( &ceEvStopKer1 );
cudaEvent_t ceEvStartCpyDev2, ceEvStopCpyDev2, ceEvStartKer2, ceEvStopKer2;
cudaEventCreate( &ceEvStartCpyDev2 );
cudaEventCreate( &ceEvStopCpyDev2 );
cudaEventCreate( &ceEvStartKer2 );
cudaEventCreate( &ceEvStopKer2 );
//Scenario1: put start R1 before kernels and start R2 between kernels
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);
//insert runker1 start event before concurrent kernels
cudaStreamWaitEvent(csStream3, ceEvStopCpyDev1, 0);
cudaEventRecord(ceEvStartKer1, csStream3);
wait_k<<<1,1,0,csStream3>>>(time_clocks);
//insert runker2 start event between concurrent kernels
cudaStreamWaitEvent(csStream4, ceEvStopCpyDev2, 0);
cudaEventRecord(ceEvStartKer2, csStream4);
wait_k<<<1,1,0,csStream4>>>(time_clocks/div);
cudaEventRecord(ceEvStopKer2, csStream4);
cudaEventRecord(ceEvStopKer1, csStream3);
cudaEventRecord(ceEvStop, 0);
cudaDeviceSynchronize();
float fTim1;
cudaEventElapsedTime( &fTim1, ceEvStart, ceEvStop);
printf("Scenario1 overall runtime = %10f\n", fTim1);
//Scenario2: put start R1 before kernels and start R2 between kernels
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);
//insert runker1 start event before concurrent kernels
cudaStreamWaitEvent(csStream3, ceEvStopCpyDev1, 0);
cudaEventRecord(ceEvStartKer1, csStream3);
//insert runker2 start event before concurrent kernels
cudaStreamWaitEvent(csStream4, ceEvStopCpyDev2, 0);
cudaEventRecord(ceEvStartKer2, csStream4);
wait_k<<<1,1,0,csStream3>>>(time_clocks);
wait_k<<<1,1,0,csStream4>>>(time_clocks/div);
cudaEventRecord(ceEvStopKer2, csStream4);
cudaEventRecord(ceEvStopKer1, csStream3);
cudaEventRecord(ceEvStop, 0);
cudaDeviceSynchronize();
float fTim2;
cudaEventElapsedTime( &fTim2, ceEvStart, ceEvStop);
printf("Scenario2 overall runtime = %10f\n", fTim2);
}
int main()
{
shorterR2_D2_R1_Overlap();
}
よろしくお願いします！