測量CUDA程序和CUDA內核耗用時間的問題

Question

我目前有三種測量已用時間的方法，兩種使用CUDA事件，另一種使用開始和結束UNIX。使用CUDA事件的人測量兩件事，一個測量整個外部循環時間，另一個測量所有內核執行時間。

下面的代碼：

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是我在這裏找到在計算器上：

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時間？

Answer 1

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版本1.5-2.2 CUDA跟蹤活動或Visual探查運行4.0+所有的時間都可用。 GPU時間將比使用cudaEvents API收集的時間更準確。使用nvtxRangePush來測量CPU時間範圍是可選的。這也可以通過從示例中的第一個CUDA API測量到cudaDeviceSynchronize結束來完成。