Временно определяне на различни секции в ядрото на CUDA

В собствения си код използвам функцията clock(), за да получа точни времена. За удобство имам макросите

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

След това те могат да се използват за инструментиране на кода на устройството, както следва:

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

    }

След това можете да прочетете cuda_timers в кода на хоста.

Няколко бележки:

• Таймерите работят на базата на блок, т.е. ако имате 100 блока, изпълняващи едно и също ядро, сумата от всичките им времена ще бъде съхранена.
• Като каза това, таймерът приема, че нулевата нишка е активна, така че се уверете, че не извиквате тези макроси в евентуално различна част от кода.
• Таймерите отчитат броя на тиктаканията на часовника. За да получите броя милисекунди, разделете това на броя GHz на вашето устройство и умножете по 1000.
• Таймерите могат малко да забавят вашия код, поради което ги увих в #ifdef USETIMERS, за да можете лесно да ги изключите.
• Въпреки че clock() връща целочислени стойности от тип clock_t, аз съхранявам натрупаните стойности като float, в противен случай стойностите ще се обвиват за ядра, които отнемат повече от няколко секунди (натрупани за всички блокове).
• Изборът ( toc > tic ) ? (toc - tic) : ( toc + (0xffffffff - tic) ) ) е необходим в случай, че броячът на часовника се увива.

P.S. Това е копие от моя отговор на този въпрос, който не получи много точки там, тъй като необходимото време беше за цялото ядро.