kfusion
kfusion copied to clipboard
Timing is inaccurate
I found two problems with the way kfusion measures how long the individual kernels take:
- Use of
return double(std::clock())/CLOCKS_PER_SEC
This is not an atomic clock, and worse it depends on the CPU activity. Changing it to
struct timespec clockData;
clock_gettime(CLOCK_MONOTONIC, &clockData);
return (double) clockData.tv_sec + clockData.tv_nsec / 1000000000.0;
makes the timing more accurate, and also shows that kfusion is around 10% faster than measured with the old way.
- CUDA kernels are asynchronous. Timing information can really only be measured after
cudaDeviceSynchronize()
, but kfusion doesn't do that in many places (apart from thetotal
time, which consequently is correct). Adding it before eachStats.sample
fixes that, and and changes my measurement forintegrate
from0.00*
milliseconds to3.*
milliseconds, which makes a lot of sense.
Note though that cudaDeviceSynchronize()
can slow things sometimes, but hasn't done so in my measurements.