0

現在、経過時間を測定する方法は 3 つあります。2 つは CUDA イベントを使用し、もう 1 つは UNIX の開始と終了を記録します。CUDA イベントを使用するものは 2 つのことを測定します。1 つは外側のループ時間全体を測定し、もう 1 つはすべてのカーネル実行時間を合計します。

コードは次のとおりです。

int64 x1, x2;

cudaEvent_t start;
cudaEvent_t end;
cudaEvent_t s1, s2;
float timeValue;


 #define timer_s cudaEventRecord(start, 0);
 #define timer_e cudaEventRecord(end, 0);   cudaEventSynchronize(end); cudaEventElapsedTime( &timeValue, start, end ); printf("time:  %f  ms \n", timeValue);


cudaEventCreate( &start );
cudaEventCreate( &end );
cudaEventCreate( &s1 );
cudaEventCreate( &s2 );

cudaEventRecord(s1, 0);   
x1 = GetTimeMs64();

for(int r = 0 ; r < 2 ; r++)
{
    timer_s
    kernel1<<<1, x>>>(gl_devdata_ptr);
    cudaThreadSynchronize();
    timer_e
    sum += timeValue;

    for(int j = 0 ; j < 5; j++)
    {
        timer_s
        kernel2<<<1,x>>>(gl_devdata_ptr);
        cudaThreadSynchronize();
        timer_e
        sum += timeValue;

        timer_s
        kernel3<<<1,x>>>(gl_devdata_ptr);
        cudaThreadSynchronize();
        timer_e
        sum += timeValue;
    }

    timer_s
    kernel4<<<y, x>>> (gl_devdata_ptr);
    cudaThreadSynchronize();
    timer_e
    sum += timeValue;
}

x2 = GetTimeMs64();

cudaEventRecord(s2, 0);   
cudaEventSynchronize(s2); 
cudaEventElapsedTime( &timeValue, s1, s2 ); 
printf("elapsed cuda :       %f  ms \n", timeValue);
printf("elapsed sum :       %f  ms \n", sum);
printf("elapsed win :       %d  ms \n", x2-x1);

GetTimeMs64 は、StackOverflow で見つけたものです。

int64 GetTimeMs64()
{
 /* Windows */
 FILETIME ft;
 LARGE_INTEGER li;
 uint64 ret;

 /* Get the amount of 100 nano seconds intervals elapsed since January 1, 1601 (UTC) and copy it
  * to a LARGE_INTEGER structure. */
 GetSystemTimeAsFileTime(&ft);
 li.LowPart = ft.dwLowDateTime;
 li.HighPart = ft.dwHighDateTime;

 ret = li.QuadPart;
 ret -= 116444736000000000LL; /* Convert from file time to UNIX epoch time. */
 ret /= 10000; /* From 100 nano seconds (10^-7) to 1 millisecond (10^-3) intervals */

 return ret;
}

これらは実際の変数名でも適切なカーネル名でもありません。コードを小さくするために一部を削除しただけです。

問題は、測定ごとに合計時間がまったく異なることです。

私が実行したいくつかの例：

elapsed cuda : 21.076832    
elapsed sum :  4.177984     
elapsed win :  27

では、なぜこれほど大きな違いがあるのでしょうか。すべてのカーネル呼び出しの合計は約 4 ミリ秒ですが、残りの 18 ミリ秒はどこにありますか? CPU時間？

4

1 に答える 1

0

cudaThreadSynchronize は、GPU 上のすべての作業が完了するまで待機する必要があるため、オーバーヘッドが非常に高い操作です。

次のようにコードを構成すると、正しい結果が得られるはずです。

int64 x1, x2;

cudaEvent_t start;
cudaEvent_t end;
const int k_maxEvents = 5 + (2 * 2) + (2 * 5 * 2);
cudaEvent_t events[k_maxEvents];
int eIdx = 0;
float timeValue;

for (int e = 0; e < 5; ++e)
{
    cudaEventCreate(&events[e]);
}

x1 = GetTimeMs64();
cudaEventRecord(events[eIdx++], 0);       
for(int r = 0 ; r < 2 ; r++)
{
    cudaEventRecord(events[eIdx++], 0);
    kernel1<<<1, x>>>(gl_devdata_ptr);

    for(int j = 0 ; j < 5; j++)
    {
        cudaEventRecord(events[eIdx++], 0);
        kernel2<<<1,x>>>(gl_devdata_ptr);

        cudaEventRecord(events[eIdx++], 0);
        kernel3<<<1,x>>>(gl_devdata_ptr);
    }

    cudaEventRecord(events[eIdx++], 0);
    kernel4<<<y, x>>> (gl_devdata_ptr);
}

cudaEventRecord(eIdx++, 0);   
cudaDeviceSynchronize(); 

x2 = GetTimeMs64();

cudaEventElapsedTime( &timeValue, events[0], events[k_maxEvents - 1] ); 
printf("elapsed cuda :       %f  ms \n", timeValue);
// TODO the time between each events is the time to execute each kernel.
// On WDDM a context switch may occur between any of the kernels leading
// to higher than expected results.
// printf("elapsed sum :       %f  ms \n", sum);
printf("elapsed win :       %d  ms \n", x2-x1);

Windows で時間を測定する簡単な方法は、QueryPerformanceCounter と QueryPerformanceFrequency を使用することです。

上記の例をイベントなしで次のように書くと

#include "NvToolsExt.h"
nvtxRangePushA("CPU Time");
for(int r = 0 ; r < 2 ; r++)
{
    kernel1<<<1, x>>>(gl_devdata_ptr);

    for(int j = 0 ; j < 5; j++)
    {
        kernel2<<<1,x>>>(gl_devdata_ptr); 
        kernel3<<<1,x>>>(gl_devdata_ptr);
    }
    kernel4<<<y, x>>> (gl_devdata_ptr);
}

cudaDeviceSynchronize(); 
nvtxRangePop();

Nsight Visual Studio Edition 1.5-2.2 CUDA Trace Activity または Visual Profiler 4.0+ で実行すると、いつでも利用できます。GPU 時間は、cudaEvents API を使用して収集できる時間よりも正確です。nvtxRangePush を使用して CPU 時間範囲を測定することはオプションです。これは、例の最初の CUDA API から cudaDeviceSynchronize の最後までを測定することによっても達成できます。

于 2012-08-21T04:58:20.947 に答える