Rewrite clock rate synchronization code to increase accuracy

Mirrors architecture used for instantaneous clock synchronization.

The PPM (parts-per-million) error after a two-second
synchronization period (using results from new implementation after
a 100-second synchronization period as ground truth, as after 100
seconds, the new and old implementation disagree by at most 2 PPM).

GPU (machine name)              Before  After
------------------------------  ------  ------
GTX 1060 3 GiB (jbakita-old):   12      1
GTX 970 (jbakita-old):          9       2
GTX 1080 Ti (yamaha):           5       1
RTX 6000 Ada (yamaha):          4       2
(An error of 1 PPM is 1 microsecond per second.)

Also modify `preemption_logger` to print the rate of skew as PPM,
rather than just as a multiplier.

1 files changed, 75 insertions, 22 deletions

diff --git a/task_host_utilities.cu b/task_host_utilities.cu
index 1f76080..1afdaf0 100644
--- a/task_host_utilities.cu
+++ b/task_host_utilities.cu
@@ -40,6 +40,17 @@ static double CurrentSeconds(void) {
   return ((double) ts.tv_sec) + (((double) ts.tv_nsec) / 1e9);
 }
 
+// Note that CLOCK_MONOTONIC_RAW is slow before Linux 5.3 because it's not
+// supported by vDSO.
+static uint64_t CurrentNanoseconds(void) {
+  struct timespec ts;
+  if (clock_gettime(CLOCK_MONOTONIC_RAW, &ts) != 0) {
+    printf("Error getting time.\n");
+    exit(1);
+  }
+  return ((uint64_t) ts.tv_sec) * 1000*1000*1000 + ((uint64_t) ts.tv_nsec);
+}
+
 // GlobalTimer64: Get 64-bit counter of current time on GPU
 // ***This is duplicated in benchmark_gpu_utilities.h***
 #if __CUDA_ARCH__ >= 300 // Kepler+
@@ -299,7 +310,7 @@ int GetMaxResidentThreads(int cuda_device) {
 }
 
 #if __CUDA_ARCH__ >= 300 // Kepler+
-static __global__ void TimerSpin(uint64_t ns_to_spin) {
+static __device__ void TimerSpin(uint64_t ns_to_spin) {
   uint64_t start_time = GlobalTimer64();
   while ((GlobalTimer64() - start_time) < ns_to_spin) {
     continue;
@@ -314,7 +325,7 @@ static __device__ inline uint32_t Clock32(void) {
 }
 
 // 'clock' can easily roll over, so handle that for ancient architectures
-static __global__ void TimerSpin(uint64_t ns_to_spin) {
+static __device__ void TimerSpin(uint64_t ns_to_spin) {
   uint64_t total_time = 0;
   uint32_t last_time = Clock32();
   while (total_time < ns_to_spin) {
@@ -335,28 +346,70 @@ static __global__ void TimerSpin(uint64_t ns_to_spin) {
 #error Fermi-based GPUs (sm_2x) are unsupported!
 #endif
 
-// This function is intended to be run in a child process. Returns -1 on error.
+// Waits on and sets barriers for the CPU before and after spinning for a
+// specified number of clock ticks.
+static __global__ void BarrierTimerSpin(uint64_t ns_to_spin, volatile uint32_t *ready_barrier,
+    volatile uint32_t *start_barrier, volatile uint32_t *end_barrier) {
+  *ready_barrier = 1;
+  while (!*start_barrier)
+    continue;
+  TimerSpin(ns_to_spin);
+  *end_barrier = 1;
+}
+
+// This function returns the number of CPU ticks that pass per GPU tick. This is
+// intended to be run in a child process. Returns -1 on error.
 static double InternalGetGPUTimerScale(int cuda_device) {
-  struct timespec start, end;
-  uint64_t nanoseconds_elapsed;
+  volatile uint32_t *gpu_start_barrier, *start_barrier = NULL;
+  volatile uint32_t *gpu_end_barrier, *end_barrier = NULL;
+  volatile uint32_t *gpu_ready_barrier, *ready_barrier = NULL;
+  volatile uint64_t cpu_start, cpu_end;
+  double rate = -1;
   if (!CheckCUDAError(cudaSetDevice(cuda_device))) return -1;
-  // Run the kernel once to warm up the GPU.
-  TimerSpin<<<1, 1>>>(1000);
-  if (!CheckCUDAError(cudaDeviceSynchronize())) return -1;
-  // After warming up, do the actual timing.
-  if (clock_gettime(CLOCK_MONOTONIC_RAW, &start) != 0) {
-    printf("Failed getting start time.\n");
-    return -1;
-  }
-  TimerSpin<<<1, 1>>>(TIMER_SPIN_DURATION);
-  if (!CheckCUDAError(cudaDeviceSynchronize())) return -1;
-  if (clock_gettime(CLOCK_MONOTONIC_RAW, &end) != 0) {
-    printf("Failed getting end time.\n");
-    return -1;
-  }
-  nanoseconds_elapsed = end.tv_sec * 1e9 + end.tv_nsec;
-  nanoseconds_elapsed -= start.tv_sec * 1e9 + start.tv_nsec;
-  return ((double) nanoseconds_elapsed) / ((double) TIMER_SPIN_DURATION);
+  if (!CheckCUDAError(cudaHostAlloc(&start_barrier, sizeof(*start_barrier),
+    cudaHostAllocMapped))) goto out;
+  if (!CheckCUDAError(cudaHostAlloc(&end_barrier, sizeof(*end_barrier),
+    cudaHostAllocMapped))) goto out;
+  if (!CheckCUDAError(cudaHostAlloc(&ready_barrier, sizeof(*ready_barrier),
+    cudaHostAllocMapped))) goto out;
+  // Setup device pointers for all the barriers
+  if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_start_barrier,
+    (uint32_t*)start_barrier, 0))) goto out;
+  if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_end_barrier,
+    (uint32_t*)end_barrier, 0))) goto out;
+  if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_ready_barrier,
+    (uint32_t*)ready_barrier, 0))) goto out;
+  // Run the kernel a first time to warm up the GPU.
+  BarrierTimerSpin<<<1, 1>>>(100, gpu_ready_barrier, gpu_start_barrier,
+    gpu_end_barrier);
+  // Barrier flows works very similarly here as in InternalReadGPUNanoseconds(),
+  // except we spin for the specified number of ticks between the start and end
+  // barriers.
+  *start_barrier = 1;
+  if (!CheckCUDAError(cudaDeviceSynchronize())) goto out;
+  // Now run the actual time-checking kernel.
+  *start_barrier = 0;
+  *end_barrier = 0;
+  *ready_barrier = 0;
+  BarrierTimerSpin<<<1, 1>>>(TIMER_SPIN_DURATION, gpu_ready_barrier, gpu_start_barrier,
+    gpu_end_barrier);
+  // Wait for kernel to initialize
+  while (!*ready_barrier)
+    continue;
+  // Immediately record CPU time and tell GPU kernel to start spinning
+  cpu_start = CurrentNanoseconds();
+  *start_barrier = 1;
+  // Wait for kernel to finish spinning, and immediately record CPU time again
+  while (!*end_barrier)
+    continue;
+  cpu_end = CurrentNanoseconds();
+  // The rate is number of CPU ticks per GPU tick
+  rate = (cpu_end - cpu_start) / ((double) TIMER_SPIN_DURATION);
+out:
+  cudaFree((void*)start_barrier);
+  cudaFree((void*)end_barrier);
+  cudaFree((void*)ready_barrier);
+  return rate;
 }
 
 double GetGPUTimerScale(int cuda_device) {