Ao sincronizar os kernels CUDA, o seguinte não funciona porque o kernel não bloqueia a execução do programa da CPU enquanto executa:

start timer
kernel<<<g,b>>>();
end timer

Eu vi três maneiras básicas de (com sucesso) sincronizar kernels CUDA:

(1) Dois eventRecords CUDA.

float responseTime; //result will be in milliseconds
cudaEvent_t start; cudaEventCreate(&start); cudaEventRecord(start); cudaEventSynchronize(start);
cudaEvent_t stop;  cudaEventCreate(&stop);
kernel<<<g,b>>>();
cudaEventRecord(stop); cudaEventSynchronize(stop);
cudaEventElapsedTime(&responseTime, start, stop); //responseTime = elapsed time

(2) Um registro de evento CUDA.

float start = read_timer(); //helper function on CPU, in milliseconds
cudaEvent_t stop;  cudaEventCreate(&stop);
kernel<<<g,b>>>();
cudaEventRecord(stop); cudaEventSynchronize(stop);
float responseTime = read_timer() - start;

(3) deviceSynchronize em vez de eventRecord. (Provavelmente apenas útil ao usar a programação em um único fluxo.)

float start = read_timer(); //helper function on CPU, in milliseconds
kernel<<<g,b>>>();
cudaDeviceSynchronize();
float responseTime = read_timer() - start;

Eu verifiquei experimentalmente que essas três estratégias produzem o mesmo resultado de tempo.

Questões:

Você provavelmente pode usar sua imaginação para descobrir o queread_timer() faz. No entanto, não faz mal fornecer uma implementação de exemplo:

double read_timer(){
    struct timeval start;
    gettimeofday( &start, NULL ); //you need to include <sys/time.h>
    return (double)((start.tv_sec) + 1.0e-6 * (start.tv_usec))*1000; //milliseconds
}