Add measure_launch_oh benchmark and make target

Benchmark supports:
=> Warms up the GPU and kernel code before sampling
=> Uses a flag in zero-copy memory to track when a kernel begins
=> Outputs summary statistics to stderr, and sample data to stdout

1 files changed, 62 insertions, 0 deletions

diff --git a/measure_launch_oh.cu b/measure_launch_oh.cu
new file mode 100644
index 0000000..981448a
--- /dev/null
+++ b/measure_launch_oh.cu
@@ -0,0 +1,62 @@
+/* Copyright 2024 Joshua Bakita
+ * Simple kernel that clocks how long a kernel launch takes to the GPU.
+ * Prints samples to stdout and summary statistics to stderr.
+ */
+#include <time.h>
+#include <stdio.h>
+#include <stdint.h>
+#include <unistd.h>
+
+#include "testbench.h"
+
+__global__ void flag_on_gpu(volatile int *flag) {
+        *flag = 1;
+}
+
+// Get launch overhead
+long measure_launch_overhead() {
+        volatile int *barrier;
+        struct timespec start, end;
+        SAFE(cudaHostAlloc(&barrier, sizeof(*barrier), cudaHostAllocMapped));
+        *barrier = 0;
+        clock_gettime(CLOCK_MONOTONIC_RAW, &start);
+        flag_on_gpu<<<100,100>>>(barrier);
+        while (!*barrier) continue;
+        clock_gettime(CLOCK_MONOTONIC_RAW, &end);
+        SAFE(cudaDeviceSynchronize());
+        return timediff(start, end);
+}
+
+int main(int argc, char **argv) {
+        int *__unused, i;
+        unsigned long num_iters;
+        long time;
+        long double cumulative_time = 0;
+
+        if (argc != 2 || !strcmp(argv[1], "--help") || !strcmp(argv[1], "-h")) {
+                fprintf(stderr, "Usage: %s [# of samples]\n", argv[0]);
+                return 1;
+        }
+
+        num_iters = strtoul(argv[1], NULL, 10);
+
+        // Initialize CUDA and a context (hack)
+        SAFE(cudaMalloc(&__unused, 8));
+
+        // Run once to ensure the kernel is compiled
+        time = measure_launch_overhead();
+        fprintf(stderr, "(%d) %ld ns (%.2f ms) warmup launch overhead\n", getpid(),
+                time, ns2ms((double)time));
+
+        // Time several kernel launches
+        for (i = 0; i < num_iters; i++) {
+                time = measure_launch_overhead();
+                cumulative_time += time;
+                // Print one sample per line
+                fprintf(stdout, "%ld\n", time);
+        }
+
+        fprintf(stderr, "(%d) %.0Lf ns (%.2Lf ms) average launch overhead\n", getpid(),
+                cumulative_time / num_iters, ns2ms(cumulative_time / num_iters));
+        return 0;
+}