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

3 files changed, 68 insertions, 2 deletions

diff --git a/.gitignore b/.gitignore
index b0d597a..7cfda47 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,5 +1,6 @@
 **/.gdb_history
 preemption_logger
 constant_cycles_kernel
+measure_launch_oh
 copy_experiments/copy_contender
 copy_experiments/mon_cross_ctx_copies
diff --git a/Makefile b/Makefile
index 688d73d..1956d72 100644
--- a/Makefile
+++ b/Makefile
@@ -1,7 +1,7 @@
 .PHONY: all copy_experiments clean
 NVCC ?= nvcc
 
-all: constant_cycles_kernel preemption_logger copy_experiments
+all: constant_cycles_kernel preemption_logger copy_experiments measure_launch_oh
 
 constant_cycles_kernel: constant_cycles_kernel.cu testbench.h
         $(NVCC) constant_cycles_kernel.cu -o constant_cycles_kernel
@@ -9,9 +9,12 @@ constant_cycles_kernel: constant_cycles_kernel.cu testbench.h
 preemption_logger: task_host_utilities.cu preemption_logger.cu testbench.h
         $(NVCC) preemption_logger.cu -o preemption_logger -g
 
+measure_launch_oh: measure_launch_oh.cu testbench.h
+        $(NVCC) measure_launch_oh.cu -o measure_launch_oh
+
 copy_experiments:
         $(MAKE) -C $@
 
 clean:
-        rm -f constant_cycles_kernel preemption_logger
+        rm -f constant_cycles_kernel preemption_logger measure_launch_oh
         $(MAKE) -C copy_experiments clean
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;
+}