From 340457bd5dae3624d4c92a02b84fb8206c0348a9 Mon Sep 17 00:00:00 2001
From: Joshua Bakita <jbakita@cs.unc.edu>
Date: Mon, 26 Feb 2024 20:29:25 -0500
Subject: 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
---
 measure_launch_oh.cu | 62 ++++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 62 insertions(+)
 create mode 100644 measure_launch_oh.cu

(limited to 'measure_launch_oh.cu')

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;
+}
-- 
cgit v1.2.2