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.

2 files changed, 4 insertions, 14 deletions

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;