2 files changed, 271 insertions, 0 deletions

diff --git a/copy_experiments/Makefile b/copy_experiments/Makefile
index b6b5a38..6ccec54 100644
--- a/copy_experiments/Makefile
+++ b/copy_experiments/Makefile
@@ -1,2 +1,10 @@
+.PHONY: all
+
+all: copy_contender mon_cross_ctx_copies
+
 copy_contender: copy_contender.cu copy_testbench.h ../testbench.h
         /usr/local/cuda/bin/nvcc copy_contender.cu -o $@ -lcuda -g
+
+mon_cross_ctx_copies: mon_cross_ctx_copies.cu copy_testbench.h ../testbench.h
+        /usr/local/cuda/bin/nvcc mon_cross_ctx_copies.cu -o $@ -lcuda -lpthread -g
+
diff --git a/copy_experiments/mon_cross_ctx_copies.cu b/copy_experiments/mon_cross_ctx_copies.cu
new file mode 100644
index 0000000..b38d037
--- /dev/null
+++ b/copy_experiments/mon_cross_ctx_copies.cu
@@ -0,0 +1,263 @@
+/* Copyright 2022 Joshua Bakita
+ * Starts two copies of configurable size in seperate contexts at the same time
+ * (by default) and logs the times that pages are copied to the GPU. Supports
+ * monitoring copy progress from either the GPU or the CPU. All output times are
+ * converted to ns.
+ * WARNING: By default, assumes that GPU and CPU clocks tick at the same rate.
+ */
+#include "copy_testbench.h"
+#include "../task_host_utilities.cu" // Just for GetHostDeviceTimeOffset()
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <unistd.h>
+#include <cuda.h>
+#include <pthread.h>
+
+// CPU-only monitoring can only go a fraction of the speed of GPU monitoring.
+// This sets how much slower it should go.
+// WARNING: Setting this too low will result in low clock resolution, but
+//          setting this too high will result in low unique sample count.
+static const int CPU_MON_DIVISOR = 50;
+// A single GPU thread cannot log page times as fast as they can be copied, so
+// spread the work across multiple threads.
+// WARNING: Setting this too high will result in extra memory contention, but
+//          setting this too low will result in inaccurate copy logs.
+// Known working settings:
+// | Chip:  | Num threads:
+// | GV100  | 8
+// | GV11b  | 32
+static const int GPU_MON_THREADS = 32;
+// If delay start is enabled on any thread, how great is the delay?
+static const struct timespec DELAY_TIME = {0, 50*1000l*1000l}; // 50ms (use 30ms for GV100)
+// Should each thread doing GPU monitoring check for GPU clock skew?
+static const bool SKEW_CHECK = false;
+
+// Size of each copy to the GPU, in bytes
+static unsigned int COPY_SIZE;
+// GPU time offset. CPU time = GPU time - offset
+static uint64_t GPU_TIME_OFFSET;
+// Copy start barrier
+static unsigned int READY = 0;
+
+typedef enum {
+        CPU_MONITORING,
+        GPU_MONITORING,
+        DELAY_START,  // For DELAY_TIME as set above
+} my_config_t;
+
+typedef struct {
+        int *cond_var;
+        my_config_t config;
+} copy_thread_args_t;
+
+// Repeatedly update a time counter in the last 8 bytes of every Nth page, where N is CPU_MON_DIVISOR
+void cpu_copy_mon(int loops, char* cpu_mem) {
+        struct timespec curr_time;
+        uint64_t curr_time_64;
+
+        for (int i = 0; i < loops; i++) {
+                clock_gettime(CLOCK_MONOTONIC_RAW, &curr_time);
+                curr_time_64 = s2ns(curr_time.tv_sec);
+                curr_time_64 += curr_time.tv_nsec;
+                // Updating pages is slower than the CE, so sparely timestamp
+                for (int offset = 0; offset < COPY_SIZE; offset += PG_SZ * CPU_MON_DIVISOR - 8)
+                        *(uint64_t*)(cpu_mem + offset) = curr_time_64;
+        }
+}
+
+// Run a single copy in it's own context (intended to be run as a thread)
+void* copy_thread(void* args_raw) {
+        CUcontext ctx;
+        int dev, i;
+        cudaError_t err;
+        int *barrier;
+        char *pinned_hostmem, *devmem;
+        uint64_t *devtimes, *times;
+        cudaStream_t stream1, stream2;
+        copy_thread_args_t *args = (copy_thread_args_t*)args_raw;
+
+        bool DELAY = args->config == DELAY_START;
+        bool GPU_COMPUTE = args->config == GPU_MONITORING;
+
+        // Explictly create a context (avoids creating a primary context implictly)
+        // This has been verified on CUDA 11.1 to give each thread a different context
+        // handle
+        cudaGetDevice(&dev);
+        SAFE(cuCtxCreate(&ctx, 0, dev));
+
+        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);
+
+        SAFE(cudaStreamCreate(&stream1));
+        SAFE(cudaStreamCreate(&stream2));
+
+        if (GPU_COMPUTE) {
+                SAFE(cudaHostAlloc(&barrier, sizeof(int) * GPU_MON_THREADS, cudaHostAllocMapped));
+                SAFE(cudaMemset(barrier, 0, sizeof(int) * GPU_MON_THREADS));
+        }
+
+        SAFE(cudaMallocHost(&pinned_hostmem, COPY_SIZE));
+        for (i = 0; i < COPY_SIZE; i++)
+                // Don't allow 0 so that the copy detection logic works
+                pinned_hostmem[i] = max((rand() & 0xff), 1);
+        SAFE(cudaMalloc(&devmem, COPY_SIZE));
+        SAFE(cudaMemset(devmem, 0, COPY_SIZE));
+
+        SAFE(cudaMalloc(&devtimes, 8 * COPY_SIZE / PG_SZ));
+        times = (uint64_t*)malloc(8 * COPY_SIZE / PG_SZ);
+        if (!times) {
+                fprintf(stderr, "Out of memory! Exiting...\n");
+                exit(1);
+        }
+
+        if (GPU_COMPUTE) {
+                // 1 thread can't go quick enough, even on the GV100
+                watch_devmem<<<1,GPU_MON_THREADS,0,stream2>>>(devmem, COPY_SIZE, devtimes, barrier);
+                // Wait for monitor(s) to initialize
+                bool ready = false;
+                while (!ready) {
+                        ready = true;
+                        for (int i = 0; i < GPU_MON_THREADS; i++)
+                                ready &= barrier[i];
+                }
+        }
+
+        // Tell our parent we're ready
+        *args->cond_var = 1;
+        // Wait for our parent to tell us to go
+        while (!READY)
+                continue;
+
+        if (DELAY)
+                nanosleep(&DELAY_TIME, NULL);
+
+        if (!GPU_COMPUTE) {
+                // If CPU-only monitoring, populate initial times
+                cpu_copy_mon(1, pinned_hostmem);
+        }
+
+        SAFE(cudaMemcpyAsync(devmem, pinned_hostmem, COPY_SIZE, cudaMemcpyHostToDevice, stream1));
+        // XXX start
+        //if (GPU_COMPUTE) {
+        //      const struct timespec DELAY_TIME = {0, 50*1000l*1000l}; // 50ms
+        //      nanosleep(&DELAY_TIME, NULL);
+        //      watch_devmem<<<1,8,0,stream2>>>(devmem, COPY_SIZE, devtimes, barrier);
+        //}
+        // XXX end
+
+        if (GPU_COMPUTE) {
+                // 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);
+                // Make sure that the copy finished in case we guessed small
+                SAFE(cudaStreamSynchronize(stream1));
+        }
+
+        // Pull back times from the GPU
+        if (GPU_COMPUTE) {
+                SAFE(cudaMemcpy(times, devtimes, 8 * COPY_SIZE / PG_SZ, cudaMemcpyDeviceToHost));
+                // Convert GPU times to CPU times
+                for (int i = 0; i < COPY_SIZE / PG_SZ; i++)
+                        times[i] -= GPU_TIME_OFFSET;
+        } else {
+                SAFE(cudaMemcpy(pinned_hostmem, devmem, COPY_SIZE, cudaMemcpyDeviceToHost));
+                int times_idx = 0;
+                // CPU-only monitoring has a fraction the # samples, so broadcast
+                for (int offset = PG_SZ * CPU_MON_DIVISOR - 8; offset < COPY_SIZE; offset += PG_SZ * CPU_MON_DIVISOR - 8) {
+                        uint64_t time = *(uint64_t*)(pinned_hostmem + offset);
+                        for (int i = 0; i < CPU_MON_DIVISOR; i++) {
+                                times[times_idx++] = time;
+                        }
+                }
+                // With CPU-only monitoring, we also record the very first page seperately
+                times[0] = *(uint64_t*)pinned_hostmem;
+        }
+
+        if (GPU_COMPUTE && SKEW_CHECK) {
+                InternalReadGPUNanoseconds(dev, &host_s2, &dev_ns2);
+                double host_diff = s2ns(host_s2 - host_s);
+                double dev_diff = dev_ns2 - dev_ns;
+                fprintf(stderr, "GPU clock ran at rate %f of CPU clock (%f ticks GPU vs %f ticks CPU)\n", dev_diff/host_diff, dev_diff, host_diff);
+        }