Add copy_contender benchmark, Makefile, and copy_testbench header

Benchmark supports:
=> Generating a chain of copies in a single context
    - No timing pause between copies
=> Copies to, from, or between device(s)
=> Variable size and number of copies
=> Logging how many milliseconds each copy takes

*** Warning ***
* Relies on implicit synchronization of unidirectional copies across
* streams. Verify that this is still the case with CUDA before
* trusting results. One can check by verifying that the runtime for
* one iteration roughly matches the time for each iteration in a
* multi-iteration run.
***

3 files changed, 248 insertions, 0 deletions

diff --git a/copy_experiments/Makefile b/copy_experiments/Makefile
new file mode 100644
index 0000000..2358c60
--- /dev/null
+++ b/copy_experiments/Makefile
@@ -0,0 +1,3 @@
+
+copy_contender: copy_contender.cu copy_testbench.h
+        /usr/local/cuda/bin/nvcc copy_contender.cu -o $@ -lcuda -g
diff --git a/copy_experiments/copy_contender.cu b/copy_experiments/copy_contender.cu
new file mode 100644
index 0000000..5514552
--- /dev/null
+++ b/copy_experiments/copy_contender.cu
@@ -0,0 +1,148 @@
+/* Copyright 2023 Joshua Bakita
+ * Serially dispatch copies in one direction and record the transfer times
+ * In order to both record individual copy times, and maintain constant
+ * pending copies, this program relies on the implict synchronization which
+ * occurs in CUDA when unidirectional copies are dispatched into seprat
+ * streams.
+ *
+ * Parameters include the size of the copy to perform and the number of times
+ * to perform said copy.
+ */
+#include <stdio.h>
+#include <cuda.h>
+
+#include "copy_testbench.h"
+
+typedef enum {
+        MODE_TO_GPU,
+        MODE_FROM_GPU,
+        MODE_PEER,
+} copy_mode_t;
+
+void usage(char** argv) {
+        fprintf(stderr, "Usage: %s <# of 4KiB pages to copy> <# of iterations>"
+               "<direction [to/from/peer] GPU>\n", argv[0]);
+}
+
+int main(int argc, char** argv) {
+        cudaError_t err;
+        int dev1, dev2;
+        size_t i;
+        char *pinned_hostmem, *devmem, *devmem_d2;
+        cudaStream_t even_stream, odd_stream;
+        struct timespec even_finish, odd_finish;
+        copy_mode_t mode;
+
+        if (argc != 4) {
+                usage(argv);
+                return 1;
+        }
+
+        // Size of each copy to the GPU, in bytes
+        const size_t COPY_SIZE = strtoul(argv[1], NULL, 10) * PG_SZ;
+        // Total copies = NUM_ITERS * 2 + 1
+        const size_t NUM_ITERS = strtoul(argv[2], NULL, 10);
+        if (strcmp(argv[3], "to") == 0)
+                mode = MODE_TO_GPU;
+        else if (strcmp(argv[3], "from") == 0)
+                mode = MODE_FROM_GPU;
+        else if (strcmp(argv[3], "peer") == 0)
+                mode = MODE_PEER;
+        else {
+                usage(argv);
+                return 1;
+        }
+        fprintf(stderr, "Copying %s %lu times %s.\n", human_readable_bytes(COPY_SIZE),
+                NUM_ITERS * 2 + 1, argv[3]);
+
+        // When doing GPU-to-GPU copy, what is the source device?
+        if (mode == MODE_PEER) {
+                SAFE(cudaGetDevice(&dev1));
+                if (dev1 == 0)
+                        dev2 = 1;
+                else
+                        dev2 = 0;
+        }
+
+        SAFE(cudaMallocHost(&pinned_hostmem, COPY_SIZE));
+        // Populate pinned_hostmem with random data
+        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));
+        if (mode == MODE_PEER) {
+                err = cudaSetDevice(dev2);
+                if (err)
+                        fprintf(stderr, "Warning: No other device available. Treating 'peer' as an internal device transfer.\n");
+                SAFE(cudaMalloc(&devmem_d2, COPY_SIZE));
+                SAFE(cudaMemset(devmem_d2, 0, COPY_SIZE));
+                SAFE(cudaSetDevice(dev1));
+        }
+
+        // If copying from device (to CPU or peer), first copy over initialization data
+        if (mode != MODE_TO_GPU) {
+                // Use a synchronizing memory copy
+                SAFE(cudaMemcpy(devmem, pinned_hostmem, COPY_SIZE, cudaMemcpyHostToDevice));
+        }
+
+        // Initialize two streams for alternating dispatch and synchronization
+        SAFE(cudaStreamCreate(&even_stream));
+        SAFE(cudaStreamCreate(&odd_stream));
+
+        // Reconfigure depending on the copy type
+        void *target, *source;
+        cudaMemcpyKind direction;
+        switch (mode) {
+                case MODE_TO_GPU:
+                        target = devmem;
+                        source = pinned_hostmem;
+                        direction = cudaMemcpyHostToDevice;
+                        break;
+                case MODE_FROM_GPU:
+                        target = pinned_hostmem;
+                        source = devmem;
+                        direction = cudaMemcpyDeviceToHost;
+                        break;
+                case MODE_PEER:
+                        target = devmem;
+                        source = devmem_d2;
+                        direction = cudaMemcpyDeviceToDevice;
+                        break;
+                default:
+                        fprintf(stderr, "FATAL ERROR: 'mode' is not a valid enum value.\n");
+                        return 1;
+        }
+
+        // Kick off the first copy that doesn't include a log
+        clock_gettime(CLOCK_MONOTONIC_RAW, &even_finish);
+        SAFE(cudaMemcpyAsync(target, source, COPY_SIZE, direction, odd_stream));
+        // Primary timing and dispatch loop. NUM_ITERS * 2 + 1 copies will be done
+        for (i = 1; NUM_ITERS == (size_t) -1 || i < NUM_ITERS + 2;) {
+                // Odd copy is underway. Queue up even copy so that there's no break.
+                SAFE(cudaMemcpyAsync(target, source, COPY_SIZE, direction, even_stream));
+                SAFE(cudaStreamSynchronize(odd_stream));
+                // Odd copy has finished. Record and print time.
+                clock_gettime(CLOCK_MONOTONIC_RAW, &odd_finish);
+                // Time for the 'odd' copy: odd_finish (end) - even_finish (start)
+                printf("%s/%lu: %ldms\n", argv[3], i++, ns2ms(timediff(even_finish, odd_finish)));
+                // Even copy is underway. Queue up odd copy so that there's no break.
+                SAFE(cudaMemcpyAsync(target, source, COPY_SIZE, direction, odd_stream));
+                SAFE(cudaStreamSynchronize(even_stream));
+                // Even copy has finished. Record and print time.
+                clock_gettime(CLOCK_MONOTONIC_RAW, &even_finish);
+                // Time for the 'even' copy: even_finish (end) - odd_finish (start)
+                printf("%s/%lu: %ldms\n", argv[3], i++, ns2ms(timediff(odd_finish, even_finish)));
+        }
+        // Record the time for the last copy to end
+        SAFE(cudaStreamSynchronize(odd_stream));
+        clock_gettime(CLOCK_MONOTONIC_RAW, &odd_finish);
+        printf("%s/%lu: %ldms\n", argv[3], i, ns2ms(timediff(even_finish, odd_finish)));
+        fprintf(stderr, "Done. Shutting down...\n");
+
+        cudaFree(pinned_hostmem);
+        cudaFree(devmem_d2);
+        cudaFree(devmem);
+        return 0;
+}
+
diff --git a/copy_experiments/copy_testbench.h b/copy_experiments/copy_testbench.h
new file mode 100644
index 0000000..e1c4aa5
--- /dev/null
+++ b/copy_experiments/copy_testbench.h
@@ -0,0 +1,97 @@
+/* Copyright 2023 Joshua Bakita
+ * Header-only library that provides an ultra-low-overhead GPU timing function,
+ * a copy monitoring kernel, and other useful sundries.
+ */
+#include <stdio.h>
+#include <stdint.h>
+
+// 4k pages (4096 bytes)
+#define PG_SZ 0x1000
+
+#define SAFE(x) \
+        if ((err = (cudaError_t)(x)) != 0) { \
+                printf("CUDA error %d! %s\n", err, cudaGetErrorString(err)); \
+                printf("Suspect line: %s\n", #x); \
+                exit(1); \
+        }
+
+#define s2ns(s) ((s)*1000l*1000l*1000l)
+#define ns2ms(s) ((s)/(1000l*1000l))
+#define timediff(start, end) ((s2ns((end).tv_sec) + (end).tv_nsec) - (s2ns((start).tv_sec) + (start).tv_nsec))
+
+// Compiles to 5 inline SASS instructions on sm_52
+static __device__ inline uint64_t gclock64_fast(void) {
+        uint32_t lo_bits;
+        uint32_t hi_bits;
+        uint32_t hi_bits_2;
+        uint64_t ret;
+        // Upper bits may rollover between our first and 2nd read
+        asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(hi_bits));
+        asm volatile("mov.u32 %0, %%globaltimer_lo;" : "=r"(lo_bits));
+        asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(hi_bits_2));
+        // If upper bits rolled over, lo_bits = 0
+        lo_bits = (hi_bits != hi_bits_2) ? 0 : lo_bits;
+        // As sm_52 SASS is naively 32-bit, the following ops get optimized out
+        ret = hi_bits_2;
+        ret <<= 32;
+        ret |= lo_bits;
+        return ret;
+}
+
+/* Old copy function. Copyright Nathan Otterness
+// Returns the value of CUDA's global nanosecond timer.
+// Compiles to 9 inline SASS instructions on sm_52
+static __device__ inline uint64_t gclock64(void) {
+  // Due to a bug in CUDA's 64-bit globaltimer, the lower 32 bits can wrap