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는 I에 유래에 여기 뭔가 : 여기

코드의 사람들은 실제 변수 이름도 올바른 커널 이름없는

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 

왜 그런 엄청난 차이가? 모든 커널 호출의 합계는 약 4ms이며, 다른 18ms는 어디입니까? CPU 시간?

Answer 1

cudaThreadSynchronize는 GPU의 모든 작업이 완료되기를 기다려야하므로 매우 높은 오버 헤드 작업입니다. Windows에서

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); 

시간을 측정 할 수있는 쉬운 방법이 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 비주얼 스튜디오 에디션 1.5-2.2 CUDA 추적 활동이나 비주얼 프로파일에서 실행하면

4.0 시대의 모든 사용할 수 있습니다. GPU 시간은 cudaEvents API를 사용하여 수집 할 수있는 것보다 더 정확합니다. nvtxRangePush를 사용하여 CPU 시간 범위를 측정하는 것은 선택 사항입니다.이 예제의 첫 번째 CUDA API에서 cudaDeviceSynchronize의 끝까지 측정하여 수행 할 수도 있습니다.