Как посчитать чистое время выполнения на ядре?

Такой подход не работает...или работает, но невероятно как-то...

call system_clock( count=c1 )
пересылка на девайс
call system_clock( count=c2 )
call name_kernel<<>>( .....)
call system_clock( count=c3 )
получение с девайса
call system_clock( count=c4 )

В итоге с2 и с3 оказываются равны.

Кто знает, как замерить чистое время счета на девайсе?

Forums: 

call system_clock( count=c1

call system_clock( count=c1 )
пересылка на девайс
call cudaThreadSynchronize() <------
call system_clock( count=c2 )
call name_kernel<<>>( .....)
call cudaThreadSynchronize() <------
call system_clock( count=c3 )
получение с девайса
call cudaThreadSynchronize() <------
call system_clock( count=c4 )

У меня внутри девайсовой

У меня внутри девайсовой процедуры идет синхронизация в конце
call syncthreads()
Это не альтернатива?

Нет, syncthreads()

Нет, syncthreads() синхронизирует треды внутри блока.

вы имеете в виду внутри

вы имеете в виду внутри __global__функции вызывается __syncthreads()?
если да, то это совсем не то что вам надо.
вариант с event'ами, который вы описали выше, самый правильный

Тут есть две обычные

Тут есть две обычные грабли

1) Вы запускаете все асинхронно (это умолчание), поэтому у вас c2 и с3 практически равны
2) разрешение таймера мало.

Второе лечится таймером с высоким разрешением (QueryPerformanceCounter - если на виндах), первое - синхронным исполнением.

Второе врядли..там считать

Второе врядли..там считать много...
Сделал вот так, и сразу заработали все таймеры как положено.
type(cudaEvent) :: sstart, sstop
uu= cudaEventCreate(sstart)
uu= cudaEventCreate(sstop)
! Get the array sizes
call system_clock( count=c1 )
пересылка на девайс
uu= cudaEventRecord(sstart, 0)
call system_clock( count=c2 )
Счет на девайсе
uu= cudaEventRecord(sstop, 0)
uu= cudaEventSynchronize(sstop)
call system_clock( count=c3 )
Пересылка на хост
call system_clock( count=c4 )
uu= cudaEventElapsedTime(time, sstart, sstop)