Timing di diverse sezioni nel kernel CUDA

Question

Ho un kernel CUDA che richiama una serie di funzioni del dispositivo.

Qual è il modo migliore per ottenere il tempo di esecuzione per ciascuna delle funzioni del dispositivo?

Qual è il modo migliore per ottenere il tempo di esecuzione per una sezione di codice in una delle funzioni del dispositivo?

Answer 1

Nel mio codice, utilizzo la funzione clock() per ottenere tempi precisi. Per comodità, ho le macro

enum { 
    tid_this = 0, 
    tid_that, 
    tid_count 
    }; 
__device__ float cuda_timers[ tid_count ]; 
#ifdef USETIMERS 
#define TIMER_TIC clock_t tic; if (threadIdx.x == 0) tic = clock(); 
#define TIMER_TOC(tid) clock_t toc = clock(); if (threadIdx.x == 0) atomicAdd(&cuda_timers[tid] , (toc > tic) ? (toc - tic) : (toc + (0xffffffff - tic))); 
#else 
#define TIMER_TIC 
#define TIMER_TOC(tid) 
#endif 

Questi possono poi essere utilizzati per strumento il codice dispositivo come segue:

__global__ mykernel (...) { 

    /* Start the timer. */ 
    TIMER_TIC 

    /* Do stuff. */ 
    ... 

    /* Stop the timer and store the results to the "timer_this" counter. */ 
    TIMER_TOC(tid_this); 

    } 

È quindi possibile leggere il cuda_timers nel codice host.

Alcune note:

• Il lavoro timer su una base per-block, vale a dire se si dispone di 100 blocchi che eseguono lo stesso kernel, la somma di tutti i tempi verrà memorizzato.
• Detto questo, il timer presuppone che il thread di zeroth sia attivo, quindi assicurati di non chiamare questi macro in una parte del codice che potrebbe divergere.
• I timer contano il numero di segni di clock. Per ottenere il numero di millisecondi, dividi questo per il numero di GHz sul tuo dispositivo e moltiplicalo per 1000.
• I timer possono rallentare il tuo codice un po ', motivo per cui li ho avvolti nello #ifdef USETIMERS in modo da poterli spegnere facilmente.
• Anche se clock() restituisce valori interi di tipo clock_t, memorizzo i valori accumulati come float, altrimenti i valori si avvolgeranno per i kernel che impiegano più di qualche secondo (accumulati su tutti i blocchi).
• La selezione (toc > tic) ? (toc - tic) : (toc + (0xffffffff - tic))) è necessaria nel caso in cui il contatore dell'orologio si avvolga.

P.S. Questa è una copia della mia risposta a this question, che non ha ottenuto molti punti lì poiché il tempo richiesto era per l'intero kernel.