From e415e282be53a3d8264b67d22b4b15b518f71099 Mon Sep 17 00:00:00 2001 From: Joshua Bakita Date: Thu, 6 Jul 2023 11:45:59 -0400 Subject: Add preemption_logger benchmark, make target, and fix d82c60c Benchmark supports: => Spinning on-GPU until some number of preemptions have occured => Monitoring GPU scheduling by logging preempt and resume times => Logging the above as raw GPU, or converted CPU, times => Measuring launch overheads Also imports task_host_utilities.cu from cuda_scheduling_examiner as of commit 9fc0ca898ef64b82cbc3734c60405723650bc66f. This is needed for timing logic used in both preemption_logger and mon_cross_ctx_copies, thus the latter would not compile prior to this commit. Also adds a new testbench helper, and a .gitignore for binaries. --- preemption_logger.cu | 158 +++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 158 insertions(+) create mode 100644 preemption_logger.cu (limited to 'preemption_logger.cu') diff --git a/preemption_logger.cu b/preemption_logger.cu new file mode 100644 index 0000000..667c459 --- /dev/null +++ b/preemption_logger.cu @@ -0,0 +1,158 @@ +/* Copyright 2023 Joshua Bakita + * Tool that tracks preemptions as discontinuities in GPU time and logs them. + * + * The logged intervals are output to standard out as a CSV, with each row + * containing a single interval. The first column contains the interval start + * time in nanoseconds (ns), and the second column contains the interval end + * time in ns. + * + * All progress and error messages and printed to standard error. + * + * Assumptions: + * - The GPU clock continues to tick, regardless of if this is the current + * application running. + */ +#include +#include +#include +#include +#include + +#include "task_host_utilities.cu" +#include "testbench.h" + +struct interval { + uint64_t start; + uint64_t end; +}; + +// Minimum time discontinuity which indicates a gap between intervals +// One clock tick is about 1ns +#define MIN_PREEMPT_TICKS 2*1000 // ~2us + +// Watch for discontinuities in the GPU clock, indicating intervals during +// which we we were preempted and another context ran. +// @param ivls Base address of array to log interval start and end times +// @param num_ivls How many intervals should be logged before exiting? +__global__ void loop_on_gpu(struct interval *ivls, int num_ivls) { + struct interval *curr_ivl = ivls; + uint64_t last_time = GlobalTimer64(); + if (num_ivls <= 0) + return; + curr_ivl->start = last_time; + // While no preemptions are occuring, this repeatedly obtains the clock, + // storing the value in `last_time`. When we're switched away from and back + // to, there will appear to be a large gap between `now` and `last_time`. + // We record `last_time` as the end of the last interval, and `now` as the + // start of the next interval, then resume repeatedly obtaining the clock. + + // Continue to log intervals until the pointer `curr_ivl` points to a + // location off the end of the `ivls` array. + while (curr_ivl < ivls + num_ivls) { + uint64_t now = GlobalTimer64(); + if (now > last_time + MIN_PREEMPT_TICKS) { + curr_ivl->end = last_time; + curr_ivl++; + curr_ivl->start = now; + } + last_time = now; + } +} + +static const char* usage_msg = "\ +Usage: %s NUM_INTERVALS [-r]\n\ +Spin on the GPU, logging intervals to stdout during which we are scheduled.\n\ + NUM_INTERVALS Number of intervals of execution to log (logs one less than\n\ + this number of preemptions).\n\ + -r, --raw Print raw logged GPU times (skip conversion to CPU time).\n"; + +int main(int argc, char **argv) { + cudaError_t err; // Needed for the SAFE() macro + struct interval *ivls_gpu, *ivls; + struct timespec start, end, end_ivls_only; + int num_ivls, skip_conversion; + uint64_t dev_ns = 0; + double d2h_scale = 1.0, host_s = 0; + pid_t pid = getpid(); + + // Argument parsing (could be improved) + if (argc != 2 && argc != 3) { + fprintf(stderr, usage_msg, argv[0]); + return 1; + } + num_ivls = atoi(argv[1]); + if (!num_ivls) { + fprintf(stderr, usage_msg, argv[0]); + return 1; + } + skip_conversion = argc == 3 && + (!strcmp(argv[2], "-r") || !strcmp(argv[2], "--raw")); + + // Synchronize the GPU and CPU clocks (if requested) using the utilities + // from task_host_utilities.cu. + if (!skip_conversion) { + // Note that this appears to almost always be 1.0 on recent GPUs + d2h_scale = GetGPUTimerScale(0); + // The following uses a subprocess for clock synchronization. Sketch? + GetHostDeviceTimeOffset(0, &host_s, &dev_ns); + } + + // Initialize GPU and CPU memory to store the intervals of execution + // Note that we use GPU memory rather than pinned host memory in order to make + // the time-monitoring loop on-GPU as fast as possible. (Pinned host memory is + // slower to read/write from/to.) + SAFE(cudaMalloc(&ivls_gpu, num_ivls * sizeof(struct interval))); + ivls = (struct interval*)malloc(num_ivls * sizeof(struct interval)); + if (!ivls) { + perror("While allocating memory"); + return 1; + } + + // Monitor execution intervals + fprintf(stderr, "(%d) Started monitoring execution for %d intervals...\n", + pid, num_ivls); + clock_gettime(CLOCK_MONOTONIC_RAW, &start); + loop_on_gpu<<<1,1>>>(ivls_gpu, num_ivls); + SAFE(cudaDeviceSynchronize()); + clock_gettime(CLOCK_MONOTONIC_RAW, &end_ivls_only); + + // Pull back results onto the CPU + SAFE(cudaMemcpy(ivls, ivls_gpu, num_ivls * sizeof(struct interval), + cudaMemcpyDeviceToHost)); + clock_gettime(CLOCK_MONOTONIC_RAW, &end); + + // Print out the intervals, the start and end times of our monitor, and some + // summary statistics + for (int i = 0; i < num_ivls; i++) { + if (skip_conversion) { + printf("%lu, %lu\n", ivls[i].start, ivls[i].end); + } else { + // Convert all the interval times into CPU time. + // First compute how far this interval's starting point is, in CPU seconds, + // from the point where the device and CPU clocks were synchronized + uint64_t start_time_since_sync = (ivls[i].start - dev_ns) * d2h_scale; + // Then convert into absolute CPU time via offset adjustment + uint64_t start = s2ns(host_s) + start_time_since_sync; + // Repeat the same process for interval end times + uint64_t end_time_since_sync = (ivls[i].end - dev_ns) * d2h_scale; + uint64_t end = s2ns(host_s) + end_time_since_sync; + printf("%lu, %lu\n", start, end); + } + } + fprintf(stderr, "(%d) Total monitor runtime: %ld ms\n", pid, + ns2ms(timediff(start, end))); + fprintf(stderr, "(%d) CPU logging start: %ld ns\n", pid, time2ns(start)); + fprintf(stderr, "(%d) CPU logging end: %ld ns\n", pid, time2ns(end)); + fprintf(stderr, "(%d) CPU copy out end: %ld ns\n", pid, time2ns(end)); + if (!skip_conversion) { + // Difference between CPU time of the start of the first interval and launch time + int64_t launch_oh = s2ns(host_s) + + (ivls[0].start - dev_ns) * d2h_scale - time2ns(start); + fprintf(stderr, "(%d) Aprox launch overhead: %ld ns\n", pid, launch_oh); + fprintf(stderr, "(%d) CPU clock - GPU clock: %ld tick gap\n", pid, + (long)s2ns(host_s) - dev_ns); + fprintf(stderr, "(%d) 1 CPU tick/1 GPU tick: %.3f\n", pid, d2h_scale); + } + + return 0; +} -- cgit v1.2.2