5 files changed, 554 insertions, 2 deletions

diff --git a/.gitignore b/.gitignore
new file mode 100644
index 0000000..e589da8
--- /dev/null
+++ b/.gitignore
@@ -0,0 +1,4 @@
+preemption_logger
+constant_cycles_kernel
+copy_experiments/copy_contender
+copy_experiments/mon_cross_ctx_copies
diff --git a/Makefile b/Makefile
index 16e17dd..9f6da4d 100644
--- a/Makefile
+++ b/Makefile
@@ -1,13 +1,16 @@
 .PHONY: all copy_experiments clean
 
-all: constant_cycles_kernel copy_experiments
+all: constant_cycles_kernel preemption_logger copy_experiments
 
 constant_cycles_kernel: constant_cycles_kernel.cu testbench.h
         nvcc constant_cycles_kernel.cu -o constant_cycles_kernel
 
+preemption_logger: task_host_utilities.cu preemption_logger.cu testbench.h
+        nvcc preemption_logger.cu -o preemption_logger -g
+
 copy_experiments:
         $(MAKE) -C $@
 
 clean:
-        rm -f constant_cycles_kernel
+        rm -f constant_cycles_kernel preemption_logger
         $(MAKE) -C copy_experiments clean
diff --git a/preemption_logger.cu b/preemption_logger.cu
new file mode 100644
index 0000000..667c459
--- /dev/null
+++ b/preemption_logger.cu
@@ -0,0 +1,158 @@
+/* Copyright 2023 Joshua Bakita
+ * Tool that tracks preemptions as discontinuities in GPU time and logs them.
+ *
+ * The logged intervals are output to standard out as a CSV, with each row
+ * containing a single interval. The first column contains the interval start
+ * time in nanoseconds (ns), and the second column contains the interval end
+ * time in ns.
+ *
+ * All progress and error messages and printed to standard error.
+ *
+ * Assumptions:
+ * - The GPU clock continues to tick, regardless of if this is the current
+ *   application running.
+ */
+#include <stdio.h>
+#include <stdint.h>
+#include <sys/types.h>
+#include <time.h>
+#include <unistd.h>
+
+#include "task_host_utilities.cu"
+#include "testbench.h"
+
+struct interval {
+        uint64_t start;
+        uint64_t end;
+};
+
+// Minimum time discontinuity which indicates a gap between intervals
+// One clock tick is about 1ns
+#define MIN_PREEMPT_TICKS 2*1000 // ~2us
+
+// Watch for discontinuities in the GPU clock, indicating intervals during
+// which we we were preempted and another context ran.
+// @param ivls     Base address of array to log interval start and end times
+// @param num_ivls How many intervals should be logged before exiting?
+__global__ void loop_on_gpu(struct interval *ivls, int num_ivls) {
+        struct interval *curr_ivl = ivls;
+        uint64_t last_time = GlobalTimer64();
+        if (num_ivls <= 0)
+                return;
+        curr_ivl->start = last_time;
+        // While no preemptions are occuring, this repeatedly obtains the clock,
+        // storing the value in `last_time`. When we're switched away from and back
+        // to, there will appear to be a large gap between `now` and `last_time`.
+        // We record `last_time` as the end of the last interval, and `now` as the
+        // start of the next interval, then resume repeatedly obtaining the clock.
+
+        // Continue to log intervals until the pointer `curr_ivl` points to a
+        // location off the end of the `ivls` array.
+        while (curr_ivl < ivls + num_ivls) {
+                uint64_t now = GlobalTimer64();
+                if (now > last_time + MIN_PREEMPT_TICKS) {
+                        curr_ivl->end = last_time;
+                        curr_ivl++;
+                        curr_ivl->start = now;
+                }
+                last_time = now;
+        }
+}
+
+static const char* usage_msg = "\
+Usage: %s NUM_INTERVALS [-r]\n\
+Spin on the GPU, logging intervals to stdout during which we are scheduled.\n\
+  NUM_INTERVALS   Number of intervals of execution to log (logs one less than\n\
+                  this number of preemptions).\n\
+  -r, --raw       Print raw logged GPU times (skip conversion to CPU time).\n";
+
+int main(int argc, char **argv) {
+        cudaError_t err; // Needed for the SAFE() macro
+        struct interval *ivls_gpu, *ivls;
+        struct timespec start, end, end_ivls_only;
+        int num_ivls, skip_conversion;
+        uint64_t dev_ns = 0;
+        double d2h_scale = 1.0, host_s = 0;
+        pid_t pid = getpid();
+
+        // Argument parsing (could be improved)
+        if (argc != 2 && argc != 3) {
+                fprintf(stderr, usage_msg, argv[0]);
+                return 1;
+        }
+        num_ivls = atoi(argv[1]);
+        if (!num_ivls) {
+                fprintf(stderr, usage_msg, argv[0]);
+                return 1;
+        }
+        skip_conversion = argc == 3 &&
+            (!strcmp(argv[2], "-r") || !strcmp(argv[2], "--raw"));
+
+        // Synchronize the GPU and CPU clocks (if requested) using the utilities
+        // from task_host_utilities.cu.
+        if (!skip_conversion) {
+                // Note that this appears to almost always be 1.0 on recent GPUs
+                d2h_scale = GetGPUTimerScale(0);
+                // The following uses a subprocess for clock synchronization. Sketch?
+                GetHostDeviceTimeOffset(0, &host_s, &dev_ns);
+        }
+
+        // Initialize GPU and CPU memory to store the intervals of execution
+        // Note that we use GPU memory rather than pinned host memory in order to make
+        // the time-monitoring loop on-GPU as fast as possible. (Pinned host memory is
+        // slower to read/write from/to.)
+        SAFE(cudaMalloc(&ivls_gpu, num_ivls * sizeof(struct interval)));
+        ivls = (struct interval*)malloc(num_ivls * sizeof(struct interval));
+        if (!ivls) {
+                perror("While allocating memory");
+                return 1;
+        }
+
+        // Monitor execution intervals
+        fprintf(stderr, "(%d) Started monitoring execution for %d intervals...\n",
+            pid, num_ivls);
+        clock_gettime(CLOCK_MONOTONIC_RAW, &start);
+        loop_on_gpu<<<1,1>>>(ivls_gpu, num_ivls);
+        SAFE(cudaDeviceSynchronize());
+        clock_gettime(CLOCK_MONOTONIC_RAW, &end_ivls_only);
+
+        // Pull back results onto the CPU
+        SAFE(cudaMemcpy(ivls, ivls_gpu, num_ivls * sizeof(struct interval),
+            cudaMemcpyDeviceToHost));
+        clock_gettime(CLOCK_MONOTONIC_RAW, &end);
+
+        // Print out the intervals, the start and end times of our monitor, and some
+        // summary statistics
+        for (int i = 0; i < num_ivls; i++) {
+                if (skip_conversion) {
+                        printf("%lu, %lu\n", ivls[i].start, ivls[i].end);
+                } else {
+                        // Convert all the interval times into CPU time.
+                        // First compute how far this interval's starting point is, in CPU seconds,
+                        // from the point where the device and CPU clocks were synchronized
+                        uint64_t start_time_since_sync = (ivls[i].start - dev_ns) * d2h_scale;
+                        // Then convert into absolute CPU time via offset adjustment
+                        uint64_t start = s2ns(host_s) + start_time_since_sync;
+                        // Repeat the same process for interval end times
+                        uint64_t end_time_since_sync = (ivls[i].end - dev_ns) * d2h_scale;
+                        uint64_t end = s2ns(host_s) + end_time_since_sync;
+                        printf("%lu, %lu\n", start, end);
+                }
+        }
+        fprintf(stderr, "(%d) Total monitor runtime: %ld ms\n", pid,
+            ns2ms(timediff(start, end)));
+        fprintf(stderr, "(%d) CPU logging start:     %ld ns\n", pid, time2ns(start));
+        fprintf(stderr, "(%d) CPU logging end:       %ld ns\n", pid, time2ns(end));
+        fprintf(stderr, "(%d) CPU copy out end:      %ld ns\n", pid, time2ns(end));
+        if (!skip_conversion) {
+                // Difference between CPU time of the start of the first interval and launch time
+                int64_t launch_oh = s2ns(host_s) +
+                    (ivls[0].start - dev_ns) * d2h_scale - time2ns(start);
+                fprintf(stderr, "(%d) Aprox launch overhead: %ld ns\n", pid, launch_oh);
+                fprintf(stderr, "(%d) CPU clock - GPU clock: %ld tick gap\n", pid,
+                    (long)s2ns(host_s) - dev_ns);
+                fprintf(stderr, "(%d) 1 CPU tick/1 GPU tick: %.3f\n", pid, d2h_scale);
+        }
+
+        return 0;
+}
diff --git a/task_host_utilities.cu b/task_host_utilities.cu
new file mode 100644
index 0000000..e35292f
--- /dev/null
+++ b/task_host_utilities.cu
@@ -0,0 +1,384 @@
+// This file contains the implementation of the functions defined in
+// task_host_utilities.h--used by task_host_utilities.c to work with the GPU.
+#include <cuda_runtime.h>
+#include <errno.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+#include <sys/mman.h>
+#include <sys/types.h>
+#include <sys/wait.h>
+#include <time.h>
+#include <unistd.h>
+
+// The number of GPU nanoseconds to spin for GetGPUTimerScale. Increasing this
+// will both increase the accuracy and the time the function takes to return.
+#define TIMER_SPIN_DURATION (2ull * 1000 * 1000 * 1000)
+
+// This macro takes a cudaError_t value. It prints an error message and returns
+// 0 if the cudaError_t isn't cudaSuccess. Otherwise, it returns nonzero.
+#define CheckCUDAError(val) (InternalCUDAErrorCheck((val), #val, __FILE__, __LINE__))
+
+// Prints an error message and returns 0 if the given CUDA result is an error.