From d82c60cbfb35697b1dbd579eccc35e0d7902e453 Mon Sep 17 00:00:00 2001 From: Joshua Bakita Date: Tue, 13 Jun 2023 08:24:26 -0400 Subject: Add mon_cross_ctx_copies benchmark and make target Benchmark supports: => Generating parallel copies from CPU to GPU in seperate process => Granular timestamping of when each portion of the copy completes => Monitoring of copy process from either CPU or GPU => Tested on platforms including the GV100 and GV11B Formerly known as diff_ctx_diff_stream. *** Warning *** * The current implementation requires some hardcoded parameters * (such as the number of monitoring threads on the GPU) to be * updated and recompiled when switching test platform. *** --- copy_experiments/Makefile | 8 + copy_experiments/mon_cross_ctx_copies.cu | 263 +++++++++++++++++++++++++++++++ 2 files changed, 271 insertions(+) create mode 100644 copy_experiments/mon_cross_ctx_copies.cu diff --git a/copy_experiments/Makefile b/copy_experiments/Makefile index b6b5a38..6ccec54 100644 --- a/copy_experiments/Makefile +++ b/copy_experiments/Makefile @@ -1,2 +1,10 @@ +.PHONY: all + +all: copy_contender mon_cross_ctx_copies + copy_contender: copy_contender.cu copy_testbench.h ../testbench.h /usr/local/cuda/bin/nvcc copy_contender.cu -o $@ -lcuda -g + +mon_cross_ctx_copies: mon_cross_ctx_copies.cu copy_testbench.h ../testbench.h + /usr/local/cuda/bin/nvcc mon_cross_ctx_copies.cu -o $@ -lcuda -lpthread -g + diff --git a/copy_experiments/mon_cross_ctx_copies.cu b/copy_experiments/mon_cross_ctx_copies.cu new file mode 100644 index 0000000..b38d037 --- /dev/null +++ b/copy_experiments/mon_cross_ctx_copies.cu @@ -0,0 +1,263 @@ +/* Copyright 2022 Joshua Bakita + * Starts two copies of configurable size in seperate contexts at the same time + * (by default) and logs the times that pages are copied to the GPU. Supports + * monitoring copy progress from either the GPU or the CPU. All output times are + * converted to ns. + * WARNING: By default, assumes that GPU and CPU clocks tick at the same rate. + */ +#include "copy_testbench.h" +#include "../task_host_utilities.cu" // Just for GetHostDeviceTimeOffset() +#include +#include +#include +#include +#include +#include + +// CPU-only monitoring can only go a fraction of the speed of GPU monitoring. +// This sets how much slower it should go. +// WARNING: Setting this too low will result in low clock resolution, but +// setting this too high will result in low unique sample count. +static const int CPU_MON_DIVISOR = 50; +// A single GPU thread cannot log page times as fast as they can be copied, so +// spread the work across multiple threads. +// WARNING: Setting this too high will result in extra memory contention, but +// setting this too low will result in inaccurate copy logs. +// Known working settings: +// | Chip: | Num threads: +// | GV100 | 8 +// | GV11b | 32 +static const int GPU_MON_THREADS = 32; +// If delay start is enabled on any thread, how great is the delay? +static const struct timespec DELAY_TIME = {0, 50*1000l*1000l}; // 50ms (use 30ms for GV100) +// Should each thread doing GPU monitoring check for GPU clock skew? +static const bool SKEW_CHECK = false; + +// Size of each copy to the GPU, in bytes +static unsigned int COPY_SIZE; +// GPU time offset. CPU time = GPU time - offset +static uint64_t GPU_TIME_OFFSET; +// Copy start barrier +static unsigned int READY = 0; + +typedef enum { + CPU_MONITORING, + GPU_MONITORING, + DELAY_START, // For DELAY_TIME as set above +} my_config_t; + +typedef struct { + int *cond_var; + my_config_t config; +} copy_thread_args_t; + +// Repeatedly update a time counter in the last 8 bytes of every Nth page, where N is CPU_MON_DIVISOR +void cpu_copy_mon(int loops, char* cpu_mem) { + struct timespec curr_time; + uint64_t curr_time_64; + + for (int i = 0; i < loops; i++) { + clock_gettime(CLOCK_MONOTONIC_RAW, &curr_time); + curr_time_64 = s2ns(curr_time.tv_sec); + curr_time_64 += curr_time.tv_nsec; + // Updating pages is slower than the CE, so sparely timestamp + for (int offset = 0; offset < COPY_SIZE; offset += PG_SZ * CPU_MON_DIVISOR - 8) + *(uint64_t*)(cpu_mem + offset) = curr_time_64; + } +} + +// Run a single copy in it's own context (intended to be run as a thread) +void* copy_thread(void* args_raw) { + CUcontext ctx; + int dev, i; + cudaError_t err; + int *barrier; + char *pinned_hostmem, *devmem; + uint64_t *devtimes, *times; + cudaStream_t stream1, stream2; + copy_thread_args_t *args = (copy_thread_args_t*)args_raw; + + bool DELAY = args->config == DELAY_START; + bool GPU_COMPUTE = args->config == GPU_MONITORING; + + // Explictly create a context (avoids creating a primary context implictly) + // This has been verified on CUDA 11.1 to give each thread a different context + // handle + cudaGetDevice(&dev); + SAFE(cuCtxCreate(&ctx, 0, dev)); + + 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); + + SAFE(cudaStreamCreate(&stream1)); + SAFE(cudaStreamCreate(&stream2)); + + if (GPU_COMPUTE) { + SAFE(cudaHostAlloc(&barrier, sizeof(int) * GPU_MON_THREADS, cudaHostAllocMapped)); + SAFE(cudaMemset(barrier, 0, sizeof(int) * GPU_MON_THREADS)); + } + + SAFE(cudaMallocHost(&pinned_hostmem, COPY_SIZE)); + for (i = 0; i < COPY_SIZE; i++) + // Don't allow 0 so that the copy detection logic works + pinned_hostmem[i] = max((rand() & 0xff), 1); + SAFE(cudaMalloc(&devmem, COPY_SIZE)); + SAFE(cudaMemset(devmem, 0, COPY_SIZE)); + + SAFE(cudaMalloc(&devtimes, 8 * COPY_SIZE / PG_SZ)); + times = (uint64_t*)malloc(8 * COPY_SIZE / PG_SZ); + if (!times) { + fprintf(stderr, "Out of memory! Exiting...\n"); + exit(1); + } + + if (GPU_COMPUTE) { + // 1 thread can't go quick enough, even on the GV100 + watch_devmem<<<1,GPU_MON_THREADS,0,stream2>>>(devmem, COPY_SIZE, devtimes, barrier); + // Wait for monitor(s) to initialize + bool ready = false; + while (!ready) { + ready = true; + for (int i = 0; i < GPU_MON_THREADS; i++) + ready &= barrier[i]; + } + } + + // Tell our parent we're ready + *args->cond_var = 1; + // Wait for our parent to tell us to go + while (!READY) + continue; + + if (DELAY) + nanosleep(&DELAY_TIME, NULL); + + if (!GPU_COMPUTE) { + // If CPU-only monitoring, populate initial times + cpu_copy_mon(1, pinned_hostmem); + } + + SAFE(cudaMemcpyAsync(devmem, pinned_hostmem, COPY_SIZE, cudaMemcpyHostToDevice, stream1)); + // XXX start + //if (GPU_COMPUTE) { + // const struct timespec DELAY_TIME = {0, 50*1000l*1000l}; // 50ms + // nanosleep(&DELAY_TIME, NULL); + // watch_devmem<<<1,8,0,stream2>>>(devmem, COPY_SIZE, devtimes, barrier); + //} + // XXX end + + if (GPU_COMPUTE) { + // 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); + // Make sure that the copy finished in case we guessed small + SAFE(cudaStreamSynchronize(stream1)); + } + + // Pull back times from the GPU + if (GPU_COMPUTE) { + SAFE(cudaMemcpy(times, devtimes, 8 * COPY_SIZE / PG_SZ, cudaMemcpyDeviceToHost)); + // Convert GPU times to CPU times + for (int i = 0; i < COPY_SIZE / PG_SZ; i++) + times[i] -= GPU_TIME_OFFSET; + } else { + SAFE(cudaMemcpy(pinned_hostmem, devmem, COPY_SIZE, cudaMemcpyDeviceToHost)); + int times_idx = 0; + // CPU-only monitoring has a fraction the # samples, so broadcast + for (int offset = PG_SZ * CPU_MON_DIVISOR - 8; offset < COPY_SIZE; offset += PG_SZ * CPU_MON_DIVISOR - 8) { + uint64_t time = *(uint64_t*)(pinned_hostmem + offset); + for (int i = 0; i < CPU_MON_DIVISOR; i++) { + times[times_idx++] = time; + } + } + // With CPU-only monitoring, we also record the very first page seperately + times[0] = *(uint64_t*)pinned_hostmem; + } + + if (GPU_COMPUTE && SKEW_CHECK) { + InternalReadGPUNanoseconds(dev, &host_s2, &dev_ns2); + double host_diff = s2ns(host_s2 - host_s); + double dev_diff = dev_ns2 - dev_ns; + fprintf(stderr, "GPU clock ran at rate %f of CPU clock (%f ticks GPU vs %f ticks CPU)\n", dev_diff/host_diff, dev_diff, host_diff); + } + + cudaFreeHost(pinned_hostmem); + cudaFree(devtimes); + cudaFree(devmem); + + return times; +} + +int main(int argc, char**argv) { + int tmp, dev; + const int NUM_THREADS = 4; + pthread_t t[NUM_THREADS]; + int is_thread_ready[NUM_THREADS] = {0}; + uint64_t *ctx_times[NUM_THREADS] = {0}; + // Setup thread condition variables and configuration (my_config_t enum) + copy_thread_args_t thread_args[NUM_THREADS] = { + {&is_thread_ready[0], CPU_MONITORING}, + {&is_thread_ready[1], CPU_MONITORING}, + {&is_thread_ready[2], CPU_MONITORING}, + {&is_thread_ready[3], CPU_MONITORING}, + };//, {&is_thread_ready[1], GPU_MONITORING}};//, {&is_thread_ready[2], CPU_MONITORING}, {&is_thread_ready[3], CPU_MONITORING}, {&is_thread_ready[4], CPU_MONITORING}, {&is_thread_ready[5], CPU_MONITORING}, {&is_thread_ready[6], CPU_MONITORING}, {&is_thread_ready[7], CPU_MONITORING}, {&is_thread_ready[8], CPU_MONITORING}}; + + if (argc != 2) { + fprintf(stderr, "Usage: %s <# of 4KB pages to copy to GPU>\n", argv[0]); + return 1; + } + if (atol(argv[1]) < CPU_MON_DIVISOR * 2) { + fprintf(stderr, "Please specify a larger copy size. It must be at least %d pages for accurate tracking.\n", CPU_MON_DIVISOR * 2); + return 2; + } + COPY_SIZE = atol(argv[1]) * PG_SZ; + + // Due to some laziness in how we handle barriers, this flag needs to be true + /// XXX: Still seems to work fine if it isn't??? + cudaGetDevice(&dev); + cudaDeviceGetAttribute(&tmp, cudaDevAttrCanUseHostPointerForRegisteredMem, dev); + /*if (!tmp) { + fprintf(stderr, "Unsupported platform. Exiting...\n"); + exit(1); + }*/ + + // 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 + uint64_t dev_ns; + double host_s; + InternalReadGPUNanoseconds(dev, &host_s, &dev_ns); + GPU_TIME_OFFSET = dev_ns - s2ns(host_s); + + srand(0); + // Spawn and wait for children to intialize + for (int tid = 0; tid < NUM_THREADS; tid++) { + pthread_create(&t[tid], NULL, copy_thread, &thread_args[tid]); + while (!is_thread_ready[tid]) + continue; + } + // 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 < NUM_THREADS; tid++) { + pthread_join(t[tid], (void**)&ctx_times[tid]); + smallest = min(smallest, ctx_times[tid][0]); + } + // Rebase times to 0 and print + for (int tid = 0; tid < NUM_THREADS; tid++) { + // Summary to stderr + fprintf(stderr, "%lu, %lu\n", ctx_times[tid][0] - smallest, ctx_times[tid][COPY_SIZE / PG_SZ - 1] - smallest); + // Full dataset to stdout + for (tmp = 0; tmp < COPY_SIZE / PG_SZ - 1; tmp++) + printf("%lu, ", ctx_times[tid][tmp]); + printf("%lu\n", ctx_times[tid][COPY_SIZE / PG_SZ - 1]); + } + return 0; +} -- cgit v1.2.2