1 files changed, 56 insertions, 0 deletions

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;
+}