1

CUDA カーネルのタイミングを計るとき、カーネルは実行中に CPU プログラムの実行をブロックしないため、以下は機能しません。

start timer
kernel<<<g,b>>>();
end timer

CUDA カーネルのタイミングを (うまく) 計る 3 つの基本的な方法を見てきました。

(1) 2 つの CUDA eventRecords。

float responseTime; //result will be in milliseconds
cudaEvent_t start; cudaEventCreate(&start); cudaEventRecord(start); cudaEventSynchronize(start);
cudaEvent_t stop;  cudaEventCreate(&stop);
kernel<<<g,b>>>();
cudaEventRecord(stop); cudaEventSynchronize(stop);
cudaEventElapsedTime(&responseTime, start, stop); //responseTime = elapsed time

(2) 1 つの CUDA eventRecord。

float start = read_timer(); //helper function on CPU, in milliseconds
cudaEvent_t stop;  cudaEventCreate(&stop);
kernel<<<g,b>>>();
cudaEventRecord(stop); cudaEventSynchronize(stop);
float responseTime = read_timer() - start;

(3) eventRecord の代わりに deviceSynchronize。(おそらく、単一のストリームでプログラミングを使用する場合にのみ役立ちます。)

float start = read_timer(); //helper function on CPU, in milliseconds
kernel<<<g,b>>>();
cudaDeviceSynchronize();
float responseTime = read_timer() - start;

これら 3 つの戦略が同じタイミング結果を生成することを実験的に確認しました。

質問:

  • これらの戦略のトレードオフは何ですか? ここに隠された詳細はありますか？
  • 複数のストリームで多くのカーネルのタイミングを計ること以外に、2 つのイベント レコードとcudaEventElapsedTime()関数を使用する利点はありますか?

おそらく想像力を働かせて、何read_timer()が機能するかを理解できるでしょう。それにもかかわらず、実装例を提供することは害にはなりません:

double read_timer(){
    struct timeval start;
    gettimeofday( &start, NULL ); //you need to include <sys/time.h>
    return (double)((start.tv_sec) + 1.0e-6 * (start.tv_usec))*1000; //milliseconds
}
4

2 に答える 2

1

あなたが示した比較的単純なケースですべて同じ結果が得られると言って（おそらく正確ではないかもしれませんが、あなたの言いたいことは理解できます）、「タイミング（複雑なシーケンス）は別として. .." 最初のケースの方が明らかに優れています。

考えられる違いの 1 つは、Windows と Linux 間の移植性です。あなたの例の read_timer 関数は Linux 指向だと思います。おそらく「移植可能な」 read_timer 関数を作成できますが、cuda イベント システム (方法 1) はそのまま移植可能です。

于 2012-12-03T03:18:30.240 に答える
0

オプション（1）は、cudaEventRecordを使用してCPUの時間を計測します。これは非常に非効率的であり、この目的でcudaEventRecordを使用することはお勧めしません。cudaEventRecordを使用して、GPUプッシュバッファー時間をカーネルを実行する時間を次のように計測できます。

float responseTime; //result will be in milliseconds
cudaEvent_t start;
cudaEvent_t stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start);
kernel<<<g,b>>>();
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&responseTime, start, stop); //responseTime = elapsed time

複数の作業項目を複数のストリームに送信する場合は、コードを少し変更する必要があります。NVVPとカウンターによって報告された時間の違いに対する答えを読むことをお勧めします

オプション（2）と（3）は、与えられた例と同様です。オプション（2）はより柔軟にすることができます。

于 2012-12-04T04:45:11.947 に答える