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 the total time, which consequently is correct). Adding it before each Stats.sample fixes that, and and changes my measurement for integrate from 0.00* milliseconds to 3.* milliseconds, which makes a lot of sense.
Note though that cudaDeviceSynchronize() can slow things sometimes, but hasn't done so in my measurements.
I found two problems with the way kfusion measures how long the individual kernels take:
return double(std::clock())/CLOCKS_PER_SECThis is not an atomic clock, and worse it depends on the CPU activity. Changing it to
makes the timing more accurate, and also shows that kfusion is around 10% faster than measured with the old way.
cudaDeviceSynchronize(), but kfusion doesn't do that in many places (apart from thetotaltime, which consequently is correct). Adding it before eachStats.samplefixes that, and and changes my measurement forintegratefrom0.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.