Fix timestamping bugs and clean up CPU side of mon_cross_ctx_copies

- Add a more reliable heuristic for how long to run the CPU-side
  copy monitor.
- Check that copies sizes are appropriately divisible.
- Fix parsing issue for copies larger than 8192 GiB.
- Clean up, and document bugs in, the time synchronization code.
- Wait for user before logging copies.

1 files changed, 57 insertions, 21 deletions

diff --git a/copy_experiments/mon_cross_ctx_copies.cu b/copy_experiments/mon_cross_ctx_copies.cu
index e836179..1163b95 100644
--- a/copy_experiments/mon_cross_ctx_copies.cu
+++ b/copy_experiments/mon_cross_ctx_copies.cu
@@ -7,9 +7,20 @@
  * - GPU clock ticks at a constant rate
  * - GPU clock ticks even while GPU is idle
  * - Different contexts view the same underlying GPU clock
-
-WARNING: By default, assumes that GPU and CPU clocks tick at the same rate.
+ * - GPU and CPU clocks tick at the same rate
+ *
+ * BUGS:
+ * - GPU and CPU clocks *do not* tick at the same rate
+ * - CPU clocks may be at different offsets and rates on different cores
+ * TODO:
+ * - Disable migrations while consistent timestamps are needed
+ * - When using GPU monitoring, synchronize clocks in each thread, and
+ *   do this serially
+ * - Support configuring copy direction
  *
+ * Note that only CPU-side monitoring is used in the RTAS'24 paper, so this
+ * tool is still correct and known-bug-free for the purposes of artifact
+ * evaluation.
  */
 #include "copy_testbench.h"
 #include "../task_host_utilities.cu"
@@ -110,7 +121,6 @@ void* copy_thread(void* args_raw) {
 
         uint64_t dev_ns, dev_ns2;
         double host_s, host_s2;
-        //GetHostDeviceTimeOffset(dev, &host_s, &dev_ns);
         if (GPU_COMPUTE && SKEW_CHECK)
                 InternalReadGPUNanoseconds(dev, &host_s, &dev_ns);
 
@@ -147,11 +157,18 @@ void* copy_thread(void* args_raw) {
                         for (int i = 0; i < GPU_MON_THREADS; i++)
                                 ready &= barrier[i];
                 }
+        } else {
+                if ((COPY_SIZE / PG_SZ) % CPU_MON_DIVISOR != 0) {
+                        fprintf(stderr, "copy_size must be divisible by %d when using CPU-monitored copy threads.\n", CPU_MON_DIVISOR);
+                        exit(1);
+                }
         }
 
         // Tell our parent we're ready
         args->is_ready = 1;
-        // Wait for our parent to tell us to go
+        // Wait for our parent to tell us to go (spinning here should also cause
+        // Linux's load-balancing logic to implictly move each monitoring thread to
+        // a separate core)
         while (!READY)
                 continue;
 
@@ -176,8 +193,10 @@ void* copy_thread(void* args_raw) {
                 // Wait for copy monitor (and hence copy) to complete
                 SAFE(cudaStreamSynchronize(stream2));
         } else {
-                // Guess number of needed CPU-only monitoring cycles
-                cpu_copy_mon(COPY_SIZE / PG_SZ, pinned_hostmem);
+                // Guess number of needed CPU-only monitoring cycles (this heuristic
+                // ensures that monitoring runs for roughly the same amount of time,
+                // no matter the timestamping granularity).
+                cpu_copy_mon(CPU_MON_DIVISOR * COPY_SIZE / PG_SZ, pinned_hostmem);
                 // Make sure that the copy finished in case we guessed small
                 SAFE(cudaStreamSynchronize(stream1));
         }
@@ -229,7 +248,7 @@ static error_t arg_parser(int key, char* arg, struct argp_state *state) {
                 case 's':
                         if (atol(arg) < CPU_MON_DIVISOR * 2)
                                 argp_error(state, "Please specify a larger copy size. It must be at least %d pages for accurate tracking.\n", CPU_MON_DIVISOR * 2);
-                        COPY_SIZE = atol(arg) * PG_SZ;
+                        COPY_SIZE = strtoull(arg, NULL, 0) * PG_SZ;
                         break;
                 case 'g':
                         if (g_args->num_threads == MAX_THREADS)
@@ -254,6 +273,16 @@ int main(int argc, char**argv) {
         pthread_t t[MAX_THREADS];
         global_args_t g_args = {0};
 
+        struct argp argp = {opts, arg_parser, 0, desc};
+        argp_parse(&argp, argc, argv, 0, 0, &g_args);
+
+        if (g_args.num_threads == 0) {
+                fprintf(stderr, "At least one copy thread must be specified with --gpu or --cpu arguments.\n");
+                return 3;
+        }
+
+        fprintf(stderr, "(%d) Synchronizing clocks and initializing copy threads...\n", getpid());
+
         // Temporarially initialize CUDA to query device attributes
         SAFE_D(cuInit(0));
         SAFE_D(cuDeviceGet(&dev_itrl, dev));
@@ -261,7 +290,7 @@ int main(int argc, char**argv) {
         /// XXX: Still seems to work fine if it isn't???
         SAFE_D(cuDeviceGetAttribute(&tmp, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, dev_itrl));
         if (!tmp) {
-                fprintf(stderr, "Unsupported platform. It must be possible to map host (CPU)"
+                fprintf(stderr, "Unsupported GPU. It must be possible to map host (CPU)"
                                 " DRAM into the GPU virtual address space for accurate clock "
                                 "synchronization. Exiting...\n");
                 return 1;
@@ -270,22 +299,25 @@ int main(int argc, char**argv) {
         // reused in subprocesses
         SAFE_D(cuDevicePrimaryCtxRelease(dev_itrl));
 
-        struct argp argp = {opts, arg_parser, 0, desc};
-        argp_parse(&argp, argc, argv, 0, 0, &g_args);
-
-        if (g_args.num_threads == 0) {
-                fprintf(stderr, "At least one copy thread must be specified with --gpu or --cpu arguments.\n");
-                return 3;
-        }
-
-        // TODO: Rewrite this so that it's all (u)int64_t
-        // Skip, as this is 1-to-1 on GV100 and GV11B
-        //double d2h_scale = GetGPUTimerScale(dev);
-        // Get the offset of GPU time from CPU time
+        double d2h_scale, host_s;
         uint64_t dev_ns;
-        double host_s;
+        // Get the core-specific offset of GPU time from CPU time, and the
+        // core-specific difference in tick rates (typical variance of -13 to 60
+        // microseconds per second).
+        // XXX: This is not sufficient for time synchronization, see "BUGS" at the
+        //      top of this file.
+        // XXX: This creates an implict context, but should reuse the above.
+        d2h_scale = InternalGetGPUTimerScale(dev);
         InternalReadGPUNanoseconds(dev, &host_s, &dev_ns);
+        if (d2h_scale == -1 || (host_s == 0 && !dev_ns)) {
+                fprintf(stderr, "Unabled to synchronize time with the GPU. Aborting...\n");
+                return 1;
+        }
         GPU_TIME_OFFSET = dev_ns - s2ns(host_s);
+        // Necessary to synchronize experiments running on different CPU cores
+        // (as CPU clocks are only semi-synchronized)
+        fprintf(stderr, "(%d) CPU clock - GPU clock: %ld tick gap\n", getpid(), (long)s2ns(host_s) - dev_ns);
+        fprintf(stderr, "(%d) 1 CPU tick/1 GPU tick: %.9f\n", getpid(), d2h_scale);
 
         // Copy buffers are filled with random numbers. Seed the RNG.
         srand(0);
@@ -295,8 +327,12 @@ int main(int argc, char**argv) {
                 while (!g_args.thread_args[tid].is_ready)
                         continue;
         }
+
+        fprintf(stderr, "(%d) Initialization completed. Press enter to start copies...", getpid());
+        getc(stdin); // Wait for user
         // Tell children initialization is done and that they can go
         READY = 1;
+
         // Wait for threads to finish, and determine the earliest recorded time
         uint64_t smallest = UINT64_MAX;
         for (int tid = 0; tid < g_args.num_threads; tid++) {