Misurare il tempo di esecuzione di OpenCL kernel

Question

Ho il seguente ciclo che misura il tempo dei miei kernel:

double elapsed = 0; 
cl_ulong time_start, time_end; 
for (unsigned i = 0; i < NUMBER_OF_ITERATIONS; ++i) 
{ 
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, &event); checkErr(err, "Kernel run"); 
    err = clWaitForEvents(1, &event); checkErr(err, "Kernel run wait fro event"); 
    err = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL); checkErr(err, "Kernel run get time start"); 
    err = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL); checkErr(err, "Kernel run get time end"); 
    elapsed += (time_end - time_start); 
} 

Poi divido elapsed da NUMBER_OF_ITERATIONS per ottenere la stima finale. Tuttavia, temo che il tempo di esecuzione dei singoli kernel sia troppo piccolo e quindi possa introdurre incertezza nelle mie misurazioni. Come posso misurare il tempo impiegato da tutti i kernel NUMBER_OF_ITERATIONS combinati?

Puoi suggerire uno strumento di profilazione, che potrebbe essere d'aiuto, poiché non è necessario accedere a questi dati a livello di programmazione. Uso NVCLIA's OpenCL.

Answer 1

Il tempo che viene misurato viene restituito in nanosecondi, ma hai ragione: la risoluzione del timer è inferiore. Tuttavia, mi chiedo quale sia l'effettivo tempo di esecuzione del tuo kernel quando dici che il tempo è troppo breve per misurare con precisione (la mia sensazione istintiva è che la risoluzione dovrebbe essere nel range di microsecondi).

Il modo più appropriato per misurare il tempo totale di più iterazioni dipende da cosa "multiplo" significa qui. NUMBER_OF_ITERATIONS=5 o NUMBER_OF_ITERATIONS=500000? Se il numero di iterazioni è "grande", si può semplicemente usare l'orologio di sistema, possibilmente con funzioni specifiche del sistema come QueryPerformanceCounter su Windows (vedi anche, ad esempio, Is there a way to measure time up to micro seconds using C standard library?), ma, naturalmente, la precisione dell'orologio di sistema potrebbe essere inferiore a quella del dispositivo OpenCL, quindi se questo ha senso in realtà dipende dal numero di iterazioni.

E 'un peccato che NVIDIA rimosso il supporto OpenCL dal loro visiva Profiler, anche se ...

Answer 2

Su implementazione GPU OpenCL Intel ho avuto successo con il suo approccio (cronometraggio al kernel) e il dosaggio preferisco a un ruscello di NDRanges.

• Un approccio alternativo è quello di eseguire N volte con e misurare il tempo con eventi marcatori come l'approccio proposto nella this question (la non domanda la risposta).

• I tempi per i kernel brevi sono generalmente almeno nel campo dei microsecondi nella mia esperienza.

• È possibile controllare la risoluzione del timer utilizzando clGetDeviceInfo con CL_DEVICE_PROFILING_TIMER_RESOLUTION (ad esempio 80 ns sul mio setup).

Answer 3

La funzione di profilatura restituisce nano secondi, ed è molto preciso (~ 50ns), tuttavia, l'esecuzione ha diversi tempi di esecuzione, a seconda di altri problemi minori che non puoi controllare.

Questo riduce la problematica su ciò che si vuole misurare:

• Misurare il tempo di esecuzione del kernel: Il tuo approccio è corretto, la precisione del tempo medio di esecuzione misurata aumenta man mano che aumenta N. Questo spiega solo il tempo di esecuzione, nessun costo aggiuntivo preso in considerazione.

• Misurare il tempo di esecuzione del kernel + spese generali: Si dovrebbe usare gli eventi come pure, ma misurare dal CL_PROFILING_COMMAND_SUBMIT, per tenere conto di spese generali di esecuzione supplementare.

• Misurazione del tempo di esecuzione del lato host reale: È necessario utilizzare anche gli eventi ma misurare dal primo avvio dell'evento fino all'ultimo evento. L'utilizzo della misurazione della temporizzazione della CPU è un'altra possibilità. Se si desidera misurare questo, è necessario rimuovere il waitforevents dal ciclo, per consentire il massimo throughput al sistema OpenCL (e meno overhead possibile).

Rispondendo alla domanda Strumenti, consiglio di utilizzare nVIDIA visual profiler. BUt poiché non è più disponibile per OpenCL, è necessario utilizzare Visual Studio Add on o una versione precedente (CUDA 3.0) di nvprofiler.

Answer 4

è necessario seguire i seguenti passi per misurare il tempo di esecuzione di OpenCL tempo di esecuzione del kernel:

1.Create una coda, profilatura necessità stati abilitare quando viene creata la coda:

cl_command_queue command_queue; 
command_queue = clCreateCommandQueue(context, devices[deviceUsed], CL_QUEUE_PROFILING_ENABLE, &err); 

2.Link un evento quando lanciare un kernel

cl_event event; 
err=clEnqueueNDRangeKernel(queue, kernel, woridim, NULL, workgroupsize, NULL, 0, NULL, &event); 

3.Attendere per il kernel per finire

clWaitForEvents(1, &event); 

4.Attendere per tutte le attività accodate alla fine

clFinish(queue); 

5.Get Profiling dati e calcolare il tempo di esecuzione del kernel (restituito dal API OpenCL in nanosecondi)

cl_ulong time_start; 
cl_ulong time_end; 

clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL); 
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL); 

double nanoSeconds = time_end-time_start; 
printf("OpenCl Execution time is: %0.3f milliseconds \n",total_time/1000000.0);