From 5b7b92b20fe638a400f9918b8da6e98c942ff5d1 Mon Sep 17 00:00:00 2001 From: Joshua Bakita Date: Tue, 13 May 2025 00:55:24 -0400 Subject: Add a sched_yield() between each interation of measure_launch_oh This is to reduce the impact of scheduling noise when measuring a very large number of samples. This change is to support the ECRTS'25 partitioning paper. --- measure_launch_oh.cu | 17 ++++++++++++++--- 1 file 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 +#include #include #include #include @@ -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); -- cgit v1.2.2