From 5c65954998591bc61a4138024ba4895bed64a8a6 Mon Sep 17 00:00:00 2001 From: Joshua Bakita Date: Wed, 21 Feb 2024 14:05:34 -0500 Subject: Fix timestamping bugs and clean up CPU side of mon_cross_ctx_copies - Add a more reliable heuristic for how long to run the CPU-side copy monitor. - Check that copies sizes are appropriately divisible. - Fix parsing issue for copies larger than 8192 GiB. - Clean up, and document bugs in, the time synchronization code. - Wait for user before logging copies. --- copy_experiments/mon_cross_ctx_copies.cu | 78 +++++++++++++++++++++++--------- 1 file changed, 57 insertions(+), 21 deletions(-) diff --git a/copy_experiments/mon_cross_ctx_copies.cu b/copy_experiments/mon_cross_ctx_copies.cu index e836179..1163b95 100644 --- a/copy_experiments/mon_cross_ctx_copies.cu +++ b/copy_experiments/mon_cross_ctx_copies.cu @@ -7,9 +7,20 @@ * - GPU clock ticks at a constant rate * - GPU clock ticks even while GPU is idle * - Different contexts view the same underlying GPU clock - -WARNING: By default, assumes that GPU and CPU clocks tick at the same rate. + * - GPU and CPU clocks tick at the same rate + * + * BUGS: + * - GPU and CPU clocks *do not* tick at the same rate + * - CPU clocks may be at different offsets and rates on different cores + * TODO: + * - Disable migrations while consistent timestamps are needed + * - When using GPU monitoring, synchronize clocks in each thread, and + * do this serially + * - Support configuring copy direction * + * Note that only CPU-side monitoring is used in the RTAS'24 paper, so this + * tool is still correct and known-bug-free for the purposes of artifact + * evaluation. */ #include "copy_testbench.h" #include "../task_host_utilities.cu" @@ -110,7 +121,6 @@ void* copy_thread(void* args_raw) { uint64_t dev_ns, dev_ns2; double host_s, host_s2; - //GetHostDeviceTimeOffset(dev, &host_s, &dev_ns); if (GPU_COMPUTE && SKEW_CHECK) InternalReadGPUNanoseconds(dev, &host_s, &dev_ns); @@ -147,11 +157,18 @@ void* copy_thread(void* args_raw) { for (int i = 0; i < GPU_MON_THREADS; i++) ready &= barrier[i]; } + } else { + if ((COPY_SIZE / PG_SZ) % CPU_MON_DIVISOR != 0) { + fprintf(stderr, "copy_size must be divisible by %d when using CPU-monitored copy threads.\n", CPU_MON_DIVISOR); + exit(1); + } } // Tell our parent we're ready args->is_ready = 1; - // Wait for our parent to tell us to go + // Wait for our parent to tell us to go (spinning here should also cause + // Linux's load-balancing logic to implictly move each monitoring thread to + // a separate core) while (!READY) continue; @@ -176,8 +193,10 @@ void* copy_thread(void* args_raw) { // Wait for copy monitor (and hence copy) to complete SAFE(cudaStreamSynchronize(stream2)); } else { - // Guess number of needed CPU-only monitoring cycles - cpu_copy_mon(COPY_SIZE / PG_SZ, pinned_hostmem); + // Guess number of needed CPU-only monitoring cycles (this heuristic + // ensures that monitoring runs for roughly the same amount of time, + // no matter the timestamping granularity). + cpu_copy_mon(CPU_MON_DIVISOR * COPY_SIZE / PG_SZ, pinned_hostmem); // Make sure that the copy finished in case we guessed small SAFE(cudaStreamSynchronize(stream1)); } @@ -229,7 +248,7 @@ static error_t arg_parser(int key, char* arg, struct argp_state *state) { case 's': if (atol(arg) < CPU_MON_DIVISOR * 2) argp_error(state, "Please specify a larger copy size. It must be at least %d pages for accurate tracking.\n", CPU_MON_DIVISOR * 2); - COPY_SIZE = atol(arg) * PG_SZ; + COPY_SIZE = strtoull(arg, NULL, 0) * PG_SZ; break; case 'g': if (g_args->num_threads == MAX_THREADS) @@ -254,6 +273,16 @@ int main(int argc, char**argv) { pthread_t t[MAX_THREADS]; global_args_t g_args = {0}; + struct argp argp = {opts, arg_parser, 0, desc}; + argp_parse(&argp, argc, argv, 0, 0, &g_args); + + if (g_args.num_threads == 0) { + fprintf(stderr, "At least one copy thread must be specified with --gpu or --cpu arguments.\n"); + return 3; + } + + fprintf(stderr, "(%d) Synchronizing clocks and initializing copy threads...\n", getpid()); + // Temporarially initialize CUDA to query device attributes SAFE_D(cuInit(0)); SAFE_D(cuDeviceGet(&dev_itrl, dev)); @@ -261,7 +290,7 @@ int main(int argc, char**argv) { /// XXX: Still seems to work fine if it isn't??? SAFE_D(cuDeviceGetAttribute(&tmp, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, dev_itrl)); if (!tmp) { - fprintf(stderr, "Unsupported platform. It must be possible to map host (CPU)" + fprintf(stderr, "Unsupported GPU. It must be possible to map host (CPU)" " DRAM into the GPU virtual address space for accurate clock " "synchronization. Exiting...\n"); return 1; @@ -270,22 +299,25 @@ int main(int argc, char**argv) { // reused in subprocesses SAFE_D(cuDevicePrimaryCtxRelease(dev_itrl)); - struct argp argp = {opts, arg_parser, 0, desc}; - argp_parse(&argp, argc, argv, 0, 0, &g_args); - - if (g_args.num_threads == 0) { - fprintf(stderr, "At least one copy thread must be specified with --gpu or --cpu arguments.\n"); - return 3; - } - - // TODO: Rewrite this so that it's all (u)int64_t - // Skip, as this is 1-to-1 on GV100 and GV11B - //double d2h_scale = GetGPUTimerScale(dev); - // Get the offset of GPU time from CPU time + double d2h_scale, host_s; uint64_t dev_ns; - double host_s; + // Get the core-specific offset of GPU time from CPU time, and the + // core-specific difference in tick rates (typical variance of -13 to 60 + // microseconds per second). + // XXX: This is not sufficient for time synchronization, see "BUGS" at the + // top of this file. + // XXX: This creates an implict context, but should reuse the above. + d2h_scale = InternalGetGPUTimerScale(dev); InternalReadGPUNanoseconds(dev, &host_s, &dev_ns); + if (d2h_scale == -1 || (host_s == 0 && !dev_ns)) { + fprintf(stderr, "Unabled to synchronize time with the GPU. Aborting...\n"); + return 1; + } GPU_TIME_OFFSET = dev_ns - s2ns(host_s); + // Necessary to synchronize experiments running on different CPU cores + // (as CPU clocks are only semi-synchronized) + fprintf(stderr, "(%d) CPU clock - GPU clock: %ld tick gap\n", getpid(), (long)s2ns(host_s) - dev_ns); + fprintf(stderr, "(%d) 1 CPU tick/1 GPU tick: %.9f\n", getpid(), d2h_scale); // Copy buffers are filled with random numbers. Seed the RNG. srand(0); @@ -295,8 +327,12 @@ int main(int argc, char**argv) { while (!g_args.thread_args[tid].is_ready) continue; } + + fprintf(stderr, "(%d) Initialization completed. Press enter to start copies...", getpid()); + getc(stdin); // Wait for user // Tell children initialization is done and that they can go READY = 1; + // Wait for threads to finish, and determine the earliest recorded time uint64_t smallest = UINT64_MAX; for (int tid = 0; tid < g_args.num_threads; tid++) { -- cgit v1.2.2