Временные ядра CUDA, которые должны выполняться более одного раза

Question

Я хочу посчитать время ядра, которое должно быть запущено более 1 раза, обрабатываемые данные различны для каждого выполняемого ядра. Мой код ниже, на время cudaMemcpy не следует учитывать.

1 cudaEvent_t start; 
2 error = cudaEventCreate(&start); 
3 cudaEvent_t stop; 
4 error = cudaEventCreate(&stop); 
6 float msecTotal = 0.0f; 
7 int nIter = 300; 
8 for (int j = 0; j < nIter; j++) 
9 {    
10  cudaMemcpy(...); 
     // Record the start event 
11  error = cudaEventRecord(start, NULL); 
12  matrixMulCUDA1<<< grid, threads >>>(...); 
     // Record the stop event 
13  error = cudaEventRecord(stop, NULL); 
14  error = cudaEventSynchronize(stop); 
15  float msec = 0.0f; 
16  error = cudaEventElapsedTime(&msec, start, stop); 
17  msecTotal+=msec; 
18 } 
19 cout<<"Total time = "<<msecTotal<<endl; 

Чтобы быть справедливым, алгоритм контраст должен быть ниже:

1 cudaEvent_t start; 
2 error = cudaEventCreate(&start); 
3 cudaEvent_t stop; 
4 error = cudaEventCreate(&stop); 
6 float msecTotal = 0.0f; 
7 int nIter = 300; 
8 for (int j = 0; j < nIter; j++) 
9 { 
     // Record the start event  
11  error = cudaEventRecord(start, NULL); 
12  matrixMulCUDA2<<< grid, threads >>>(...); 
     // Record the stop event 
13  error = cudaEventRecord(stop, NULL); 
14  error = cudaEventSynchronize(stop); 
15  float msec = 0.0f; 
16  error = cudaEventElapsedTime(&msec, start, stop); 
17  msecTotal+=msec; 
18 } 
19 cout<<"Total time = "<<msecTotal<<endl; 

Мой вопрос заключается в том, что метод является правильным? потому что я не уверен. Очевидно, что время должно быть больше, чем обычно.

Answer 1

Вы должны получить похожие результаты в любом случае. Записывая события вокруг запуска ядра, вы определенно измеряете только время, проведенное в ядре, а не время, потраченное на memcpy.

Мое единственное, что при вызове cudaEventSynchronize() на каждой итерации цикла вы нарушаете параллелизм CPU/GPU, что на самом деле очень важно для получения хорошей производительности. Если вы должны время каждого вызова ядра отдельно (вместо того, чтобы помещать цикл цикла nIter вокруг вызова ядра в отличие от всей операции), вы можете выделить больше событий CUDA. Если вы идете по этому маршруту, вам не нужно 2 события на одну итерацию цикла - вам нужно скопировать операцию с двумя, и для записи требуется только одно событие CUDA для каждой итерации цикла. Тогда время для любого заданного вызова ядра можно вычислить, вызвав cudaEventElapsedTime() на смежные записанные события.

Для записи времени GPU между N событий:

cudaEvent_t events[N+2];

cudaEventRecord(events[0], NULL); // record first event 
for (j = 0; j < nIter; j++) { 
    // invoke kernel, or do something else you want to time 
    // cudaEventRecord(events[j+1], NULL); 
} 
cudaEventRecord(events[j], NULL); 
// to compute the time taken for operation i, call: 
float ms; 
cudaEventElapsedTime(&ms, events[i+1], events[i]);