Синхронизация различных разделов в ядре 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 блоков, выполняющих одно и то же ядро, будет сохранена сумма всех их времен.
• Сказав это, таймер предполагает, что активен нулевой поток, поэтому убедитесь, что вы не вызываете эти макросы в возможно расходящейся части кода.
• Таймеры подсчитывают количество тактов часов. Чтобы получить количество миллисекунд, разделите его на количество ГГц на вашем устройстве и умножьте на 1000.
• Таймеры могут немного замедлить ваш код, поэтому я завернул их в #ifdef USETIMERS, чтобы вы могли легко их отключить.
• Хотя clock() возвращает целочисленные значения типа clock_t, я храню накопленные значения как float, иначе значения будут циклически повторяться для ядер, которые занимают больше нескольких секунд (накопляются по всем блокам).
• Выбор ( toc > tic ) ? (toc - tic) : ( toc + (0xffffffff - tic) ) ) необходим в случае, если счетчик часов зацикливается.

P.S. Это копия моего ответа на этот вопрос, который не получил много баллов там, поскольку требуемое время было для всего ядра.