1 files changed, 14 insertions, 3 deletions

diff --git a/measure_launch_oh.cu b/measure_launch_oh.cu
index 981448a..0a253a2 100644
--- a/measure_launch_oh.cu
+++ b/measure_launch_oh.cu
@@ -1,8 +1,11 @@
 /* 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.
+ *
+ * For best results, pin to a CPU core and redirect interrupts.
  */
 #include <time.h>
+#include <sched.h>
 #include <stdio.h>
 #include <stdint.h>
 #include <unistd.h>
@@ -14,7 +17,7 @@ __global__ void flag_on_gpu(volatile int *flag) {
 }
 
 // Get launch overhead
-long measure_launch_overhead() {
+long measure_launch_overhead(bool first_run) {
         volatile int *barrier;
         struct timespec start, end;
         SAFE(cudaHostAlloc(&barrier, sizeof(*barrier), cudaHostAllocMapped));
@@ -24,6 +27,9 @@ long measure_launch_overhead() {
         while (!*barrier) continue;
         clock_gettime(CLOCK_MONOTONIC_RAW, &end);
         SAFE(cudaDeviceSynchronize());
+        if (first_run)
+                fprintf(stderr, "(%d) First launch completed at CLOCK_MONOTONIC_RAW = %ld ns\n",
+                                getpid(), time2ns(end));
         return timediff(start, end);
 }
 
@@ -44,13 +50,18 @@ int main(int argc, char **argv) {
         SAFE(cudaMalloc(&__unused, 8));
 
         // Run once to ensure the kernel is compiled
-        time = measure_launch_overhead();
+        time = measure_launch_overhead(true);
         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();
+                // We're measuring on the order of microseconds. Make sure the
+                // kernel scheduler does not interrupt us during a timing
+                // iteration by explicitly invoking it here.
+                sched_yield();
+                // Measure overhead
+                time = measure_launch_overhead(false);
                 cumulative_time += time;
                 // Print one sample per line
                 fprintf(stdout, "%ld\n", time);