Add constant_cycles_kernel benchmark + header split and Makefile

Benchmark supports:
=> Looping on the GPU for a specified number of iterations
=> Looping infinitely on the GPU
=> Logging the sum time (ns and ms) for all iterations to complete

Also splits `testbench.h` out of `copy_testbench.h` for use in other
benchmarks.

5 files changed, 80 insertions, 15 deletions

diff --git a/Makefile b/Makefile
index e81353d..7eaf03d 100644
--- a/Makefile
+++ b/Makefile
@@ -1,4 +1,9 @@
-.PHONY: copy_experiments
+.PHONY: all copy_experiments
+
+all: constant_cycles_kernel copy_experiments
+
+constant_cycles_kernel: constant_cycles_kernel.cu testbench.h
+        nvcc constant_cycles_kernel.cu -o constant_cycles_kernel
 
 copy_experiments:
         $(MAKE) -C $@
diff --git a/constant_cycles_kernel.cu b/constant_cycles_kernel.cu
new file mode 100644
index 0000000..eda963d
--- /dev/null
+++ b/constant_cycles_kernel.cu
@@ -0,0 +1,56 @@
+/* Copyright 2021-2023 Joshua Bakita
+ * Simple kernel that spins on the GPU for a specified number of iterations,
+ * while tracking and printing the necessary CPU time.
+ */
+#include <time.h>
+#include <stdio.h>
+#include <cuda.h>
+
+#include "testbench.h"
+
+__global__ void loop_on_gpu(unsigned long iters, int *__unused) {
+        for (volatile int i = 0; i < iters; i++) (*__unused)++;
+}
+
+int main(int argc, char **argv) {
+        cudaError_t err;
+        int res, *__unused;
+        struct timespec start, end;
+
+        if (argc != 2) {
+                fprintf(stderr, "Usage: %s <# of millions of iterations, or -1 for infinite>\n",
+                        argv[0]);
+                return 1;
+        }
+
+        // Input is multiplied by one million, unless infinite
+        unsigned long num_iters = strtoul(argv[1], NULL, 10);
+        if (num_iters != (unsigned long)(-1))
+                num_iters *= 1000 * 1000;
+
+        // Initialize CUDA and a context (hack)
+        SAFE(cudaMalloc(&__unused, 8));
+
+        // Run iterations on a single thread
+        clock_gettime(CLOCK_MONOTONIC_RAW, &start);
+        loop_on_gpu<<<1,1>>>(num_iters, __unused);
+        SAFE(cudaGetLastError() /* Check successful launch */);
+        SAFE(cudaDeviceSynchronize());
+        clock_gettime(CLOCK_MONOTONIC_RAW, &end);
+
+        // Print detailed timing information
+        long elapsed = timediff(start, end);
+        fprintf(stderr, "Started at %ld ns, ended at %ld ns\n",
+                s2ns(start.tv_sec) + start.tv_nsec, s2ns(end.tv_sec) + end.tv_nsec);
+        fprintf(stderr, "%ld ns (%.2f ms) elapsed\n", elapsed, elapsed / (1000 * 1000.));
+
+        // Verify success (fool optimizer)
+        SAFE(cudaMemcpy(&res, __unused, 8, cudaMemcpyDeviceToHost));
+        // Theoretically this can happen if `__unused` wraps around; maybe for a
+        // very small `long` type, or a very long run. More likely indicates an error.
+        if (!res)
+                fprintf(stderr, "CRITICAL: Zero iterations seem completed. Likely incorrect "
+                        "arguments, internal error, or corruption of CUDA internal state.\n");
+
+        return 0;
+}
diff --git a/copy_experiments/Makefile b/copy_experiments/Makefile
index 2358c60..b6b5a38 100644
--- a/copy_experiments/Makefile
+++ b/copy_experiments/Makefile
@@ -1,3 +1,2 @@
-
-copy_contender: copy_contender.cu copy_testbench.h
+copy_contender: copy_contender.cu copy_testbench.h ../testbench.h
         /usr/local/cuda/bin/nvcc copy_contender.cu -o $@ -lcuda -g
diff --git a/copy_experiments/copy_testbench.h b/copy_experiments/copy_testbench.h
index e1c4aa5..fed62ba 100644
--- a/copy_experiments/copy_testbench.h
+++ b/copy_experiments/copy_testbench.h
@@ -1,24 +1,15 @@
-/* Copyright 2023 Joshua Bakita
+/* Copyright 2021-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>
 
+#include "../testbench.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;
diff --git a/testbench.h b/testbench.h
new file mode 100644
index 0000000..28d7f77
--- /dev/null
+++ b/testbench.h
@@ -0,0 +1,14 @@
+/* Copyright 2021-2023 Joshua Bakita
+ * Header for miscellaneous experimental helper functions.
+ */
+
+#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))