Rozmieszczenie różnych sekcji w jądrze CUDA

Question

Mam jądro CUDA, które wywołuje szereg funkcji urządzenia.

Jaki jest najlepszy sposób uzyskania czasu wykonania dla każdej funkcji urządzenia?

Jaki jest najlepszy sposób uzyskania czasu wykonania sekcji kodu w jednej z funkcji urządzenia?

Answer 1

W moim własnym kodzie używam funkcji clock(), aby uzyskać dokładny pomiar czasu. Dla wygody, mam makra

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 

te mogą być następnie wykorzystane do instrumentu kod urządzenia w następujący sposób:

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

    } 

Następnie można odczytać cuda_timers w kodzie gospodarza.

Kilka uwag:

• Prace czasomierze na zasadzie per-bloku, to znaczy, jeśli masz 100 bloków wykonania tego samego jądra, suma wszystkich swoich czasów będą przechowywane.
• Po tym, zegar zakłada, że ​​zerowy wątek jest aktywny, więc upewnij się, że nie wywołasz tych makr w prawdopodobnie rozbieżnej części kodu.
• Liczniki liczą liczbę tyknięć zegara. Aby uzyskać liczbę milisekund, podziel to przez liczbę GHz na swoim urządzeniu i pomnóż przez 1000.
• Timery mogą nieco spowolnić twój kod, dlatego zapakowałem je w #ifdef USETIMERS, abyś mógł je wyłączyć z łatwością.
• Chociaż zwraca wartości całkowite typu clock_t, przechowuję skumulowane wartości jako float, w przeciwnym razie wartości zostaną zawinięte dla jądra, które trwają dłużej niż kilka sekund (nagromadzone na wszystkich blokach).
• Wybór (toc > tic) ? (toc - tic) : (toc + (0xffffffff - tic))) jest konieczny w przypadku, gdy licznik zegara zawija się.

P.S. To jest kopia mojej odpowiedzi na this question, która nie dostała wielu punktów, ponieważ czas wymagany był dla całego jądra.