10

一連のデバイス関数を呼び出す CUDA カーネルがあります。

各デバイス機能の実行時間を取得する最良の方法は何ですか?

デバイス関数の 1 つのコードのセクションの実行時間を取得する最良の方法は何ですか?

4

1 に答える 1

7

私自身のコードでは、このclock()関数を使用して正確なタイミングを取得しています。便宜上、私はマクロを持っています

enum {
    tid_this = 0,
    tid_that,
    tid_count
    };
__device__ float cuda_timers[ tid_count ];
#ifdef USETIMERS
 #define TIMER_TIC clock_t tic; if ( threadIdx.x == 0 ) tic = clock();
 #define TIMER_TOC(tid) clock_t toc = clock(); if ( threadIdx.x == 0 ) atomicAdd( &cuda_timers[tid] , ( toc > tic ) ? (toc - tic) : ( toc + (0xffffffff - tic) ) );
#else
 #define TIMER_TIC
 #define TIMER_TOC(tid)
#endif

次に、これらを使用して、次のようにデバイス コードをインストルメント化できます。

__global__ mykernel ( ... ) {

    /* Start the timer. */
    TIMER_TIC

    /* Do stuff. */
    ...

    /* Stop the timer and store the results to the "timer_this" counter. */
    TIMER_TOC( tid_this );

    }

cuda_timersその後、ホスト コードでを読み取ることができます。

いくつかのメモ:

  • タイマーはブロック単位で動作します。つまり、同じカーネルを実行する 100 個のブロックがある場合、すべての時間の合計が保存されます。
  • そうは言っても、タイマーはゼロ番目のスレッドがアクティブであると想定しているため、コードの分岐する可能性のある部分でこれらのマクロを呼び出さないようにしてください。
  • タイマーは、クロック ティックの数をカウントします。ミリ秒数を取得するには、これをデバイスの GHz 数で割り、1000 を掛けます。
  • タイマーはコードの速度を少し遅くする可能性がある#ifdef USETIMERSため、簡単にオフにできるように でラップしました。
  • clock()タイプ の整数値を返しますが、clock_t累積値を として保存しますfloat。そうしないと、数秒以上かかるカーネルの値がラップアラウンドします (すべてのブロックで累積されます)。
  • ( toc > tic ) ? (toc - tic) : ( toc + (0xffffffff - tic) ) )クロックカウンタがラップアラウンドする場合に備えて選択が必要です。

PS これは、この質問に対する私の回答のコピーです。必要なタイミングはカーネル全体であったため、多くのポイントを獲得できませんでした。

于 2012-06-26T14:12:10.643 に答える