Temporización de diferentes secciones en el kernel CUDA

Question

Tengo un kernel CUDA que llama a una serie de funciones del dispositivo.

¿Cuál es la mejor manera de obtener el tiempo de ejecución para cada una de las funciones del dispositivo?

¿Cuál es la mejor manera de obtener el tiempo de ejecución para una sección de código en una de las funciones del dispositivo?

Answer 1

En mi propio código, uso la función clock() para obtener sincronizaciones precisas. Para mayor comodidad, tengo las macros

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 

Estos pueden ser utilizados para instrumentar el código del dispositivo de la siguiente manera:

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

    } 

A continuación, puede leer el cuda_timers en el código de acogida.

Unas pocas notas:

• los temporizadores de trabajo en función de cada bloque, es decir, si tiene 100 bloques de ejecución del mismo núcleo, se almacena la suma de todos sus tiempos.
• Habiendo dicho eso, el temporizador supone que el subproceso zeroth está activo, así que asegúrese de no llamar a estas macros en una parte posiblemente divergente del código.
• Los temporizadores cuentan el número de marcas de reloj. Para obtener el número de milisegundos, divida esto por el número de GHz en su dispositivo y multiplique por 1000.
• Los temporizadores pueden ralentizar su código un poco, por eso los envolví en el #ifdef USETIMERS para que pueda apagarlos fácilmente.
• Aunque clock() devuelve valores enteros del tipo clock_t, almaceno los valores acumulados como float; de lo contrario, los valores se ajustarán para los kernels que tarden más que unos pocos segundos (acumulados en todos los bloques).
• La selección (toc > tic) ? (toc - tic) : (toc + (0xffffffff - tic))) es necesaria en caso de que el contador del reloj se ajuste.

P.S. Esta es una copia de mi respuesta al this question, que no obtuvo muchos puntos allí ya que el tiempo requerido era para todo el kernel.