From e63404cfda97050b9d516232ceb4f918f48edcfc Mon Sep 17 00:00:00 2001 From: Joshua Bakita Date: Fri, 9 Jun 2023 18:55:07 -0400 Subject: Add copy_contender benchmark, Makefile, and copy_testbench header Benchmark supports: => Generating a chain of copies in a single context - No timing pause between copies => Copies to, from, or between device(s) => Variable size and number of copies => Logging how many milliseconds each copy takes *** Warning *** * Relies on implicit synchronization of unidirectional copies across * streams. Verify that this is still the case with CUDA before * trusting results. One can check by verifying that the runtime for * one iteration roughly matches the time for each iteration in a * multi-iteration run. *** --- Makefile | 4 + copy_experiments/Makefile | 3 + copy_experiments/copy_contender.cu | 148 +++++++++++++++++++++++++++++++++++++ copy_experiments/copy_testbench.h | 97 ++++++++++++++++++++++++ 4 files changed, 252 insertions(+) create mode 100644 Makefile create mode 100644 copy_experiments/Makefile create mode 100644 copy_experiments/copy_contender.cu create mode 100644 copy_experiments/copy_testbench.h diff --git a/Makefile b/Makefile new file mode 100644 index 0000000..e81353d --- /dev/null +++ b/Makefile @@ -0,0 +1,4 @@ +.PHONY: copy_experiments + +copy_experiments: + $(MAKE) -C $@ diff --git a/copy_experiments/Makefile b/copy_experiments/Makefile new file mode 100644 index 0000000..2358c60 --- /dev/null +++ b/copy_experiments/Makefile @@ -0,0 +1,3 @@ + +copy_contender: copy_contender.cu copy_testbench.h + /usr/local/cuda/bin/nvcc copy_contender.cu -o $@ -lcuda -g diff --git a/copy_experiments/copy_contender.cu b/copy_experiments/copy_contender.cu new file mode 100644 index 0000000..5514552 --- /dev/null +++ b/copy_experiments/copy_contender.cu @@ -0,0 +1,148 @@ +/* Copyright 2023 Joshua Bakita + * Serially dispatch copies in one direction and record the transfer times + * In order to both record individual copy times, and maintain constant + * pending copies, this program relies on the implict synchronization which + * occurs in CUDA when unidirectional copies are dispatched into seprat + * streams. + * + * Parameters include the size of the copy to perform and the number of times + * to perform said copy. + */ +#include +#include + +#include "copy_testbench.h" + +typedef enum { + MODE_TO_GPU, + MODE_FROM_GPU, + MODE_PEER, +} copy_mode_t; + +void usage(char** argv) { + fprintf(stderr, "Usage: %s <# of 4KiB pages to copy> <# of iterations>" + "\n", argv[0]); +} + +int main(int argc, char** argv) { + cudaError_t err; + int dev1, dev2; + size_t i; + char *pinned_hostmem, *devmem, *devmem_d2; + cudaStream_t even_stream, odd_stream; + struct timespec even_finish, odd_finish; + copy_mode_t mode; + + if (argc != 4) { + usage(argv); + return 1; + } + + // Size of each copy to the GPU, in bytes + const size_t COPY_SIZE = strtoul(argv[1], NULL, 10) * PG_SZ; + // Total copies = NUM_ITERS * 2 + 1 + const size_t NUM_ITERS = strtoul(argv[2], NULL, 10); + if (strcmp(argv[3], "to") == 0) + mode = MODE_TO_GPU; + else if (strcmp(argv[3], "from") == 0) + mode = MODE_FROM_GPU; + else if (strcmp(argv[3], "peer") == 0) + mode = MODE_PEER; + else { + usage(argv); + return 1; + } + fprintf(stderr, "Copying %s %lu times %s.\n", human_readable_bytes(COPY_SIZE), + NUM_ITERS * 2 + 1, argv[3]); + + // When doing GPU-to-GPU copy, what is the source device? + if (mode == MODE_PEER) { + SAFE(cudaGetDevice(&dev1)); + if (dev1 == 0) + dev2 = 1; + else + dev2 = 0; + } + + SAFE(cudaMallocHost(&pinned_hostmem, COPY_SIZE)); + // Populate pinned_hostmem with random data + 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)); + if (mode == MODE_PEER) { + err = cudaSetDevice(dev2); + if (err) + fprintf(stderr, "Warning: No other device available. Treating 'peer' as an internal device transfer.\n"); + SAFE(cudaMalloc(&devmem_d2, COPY_SIZE)); + SAFE(cudaMemset(devmem_d2, 0, COPY_SIZE)); + SAFE(cudaSetDevice(dev1)); + } + + // If copying from device (to CPU or peer), first copy over initialization data + if (mode != MODE_TO_GPU) { + // Use a synchronizing memory copy + SAFE(cudaMemcpy(devmem, pinned_hostmem, COPY_SIZE, cudaMemcpyHostToDevice)); + } + + // Initialize two streams for alternating dispatch and synchronization + SAFE(cudaStreamCreate(&even_stream)); + SAFE(cudaStreamCreate(&odd_stream)); + + // Reconfigure depending on the copy type + void *target, *source; + cudaMemcpyKind direction; + switch (mode) { + case MODE_TO_GPU: + target = devmem; + source = pinned_hostmem; + direction = cudaMemcpyHostToDevice; + break; + case MODE_FROM_GPU: + target = pinned_hostmem; + source = devmem; + direction = cudaMemcpyDeviceToHost; + break; + case MODE_PEER: + target = devmem; + source = devmem_d2; + direction = cudaMemcpyDeviceToDevice; + break; + default: + fprintf(stderr, "FATAL ERROR: 'mode' is not a valid enum value.\n"); + return 1; + } + + // Kick off the first copy that doesn't include a log + clock_gettime(CLOCK_MONOTONIC_RAW, &even_finish); + SAFE(cudaMemcpyAsync(target, source, COPY_SIZE, direction, odd_stream)); + // Primary timing and dispatch loop. NUM_ITERS * 2 + 1 copies will be done + for (i = 1; NUM_ITERS == (size_t) -1 || i < NUM_ITERS + 2;) { + // Odd copy is underway. Queue up even copy so that there's no break. + SAFE(cudaMemcpyAsync(target, source, COPY_SIZE, direction, even_stream)); + SAFE(cudaStreamSynchronize(odd_stream)); + // Odd copy has finished. Record and print time. + clock_gettime(CLOCK_MONOTONIC_RAW, &odd_finish); + // Time for the 'odd' copy: odd_finish (end) - even_finish (start) + printf("%s/%lu: %ldms\n", argv[3], i++, ns2ms(timediff(even_finish, odd_finish))); + // Even copy is underway. Queue up odd copy so that there's no break. + SAFE(cudaMemcpyAsync(target, source, COPY_SIZE, direction, odd_stream)); + SAFE(cudaStreamSynchronize(even_stream)); + // Even copy has finished. Record and print time. + clock_gettime(CLOCK_MONOTONIC_RAW, &even_finish); + // Time for the 'even' copy: even_finish (end) - odd_finish (start) + printf("%s/%lu: %ldms\n", argv[3], i++, ns2ms(timediff(odd_finish, even_finish))); + } + // Record the time for the last copy to end + SAFE(cudaStreamSynchronize(odd_stream)); + clock_gettime(CLOCK_MONOTONIC_RAW, &odd_finish); + printf("%s/%lu: %ldms\n", argv[3], i, ns2ms(timediff(even_finish, odd_finish))); + fprintf(stderr, "Done. Shutting down...\n"); + + cudaFree(pinned_hostmem); + cudaFree(devmem_d2); + cudaFree(devmem); + return 0; +} + diff --git a/copy_experiments/copy_testbench.h b/copy_experiments/copy_testbench.h new file mode 100644 index 0000000..e1c4aa5 --- /dev/null +++ b/copy_experiments/copy_testbench.h @@ -0,0 +1,97 @@ +/* Copyright 2023 Joshua Bakita + * Header-only library that provides an ultra-low-overhead GPU timing function, + * a copy monitoring kernel, and other useful sundries. + */ +#include +#include + +// 4k pages (4096 bytes) +#define PG_SZ 0x1000 + +#define SAFE(x) \ + if ((err = (cudaError_t)(x)) != 0) { \ + printf("CUDA error %d! %s\n", err, cudaGetErrorString(err)); \ + printf("Suspect line: %s\n", #x); \ + exit(1); \ + } + +#define s2ns(s) ((s)*1000l*1000l*1000l) +#define ns2ms(s) ((s)/(1000l*1000l)) +#define timediff(start, end) ((s2ns((end).tv_sec) + (end).tv_nsec) - (s2ns((start).tv_sec) + (start).tv_nsec)) + +// Compiles to 5 inline SASS instructions on sm_52 +static __device__ inline uint64_t gclock64_fast(void) { + uint32_t lo_bits; + uint32_t hi_bits; + uint32_t hi_bits_2; + uint64_t ret; + // Upper bits may rollover between our first and 2nd read + asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(hi_bits)); + asm volatile("mov.u32 %0, %%globaltimer_lo;" : "=r"(lo_bits)); + asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(hi_bits_2)); + // If upper bits rolled over, lo_bits = 0 + lo_bits = (hi_bits != hi_bits_2) ? 0 : lo_bits; + // As sm_52 SASS is naively 32-bit, the following ops get optimized out + ret = hi_bits_2; + ret <<= 32; + ret |= lo_bits; + return ret; +} + +/* Old copy function. Copyright Nathan Otterness +// Returns the value of CUDA's global nanosecond timer. +// Compiles to 9 inline SASS instructions on sm_52 +static __device__ inline uint64_t gclock64(void) { + // Due to a bug in CUDA's 64-bit globaltimer, the lower 32 bits can wrap + // around after the upper bits have already been read. Work around this by + // reading the high bits a second time. Use the second value to detect a + // rollover, and set the lower bits of the 64-bit "timer reading" to 0, which + // would be valid, it's passed over during the duration of the reading. If no + // rollover occurred, just return the initial reading. + volatile uint64_t first_reading; + volatile uint32_t second_reading; + uint32_t high_bits_first; + asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(first_reading)); + high_bits_first = first_reading >> 32; + asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(second_reading)); + if (high_bits_first == second_reading) { + return first_reading; + } + // Return the value with the updated high bits, but the low bits set to 0. + return ((uint64_t) second_reading) << 32; +} +*/ + +// This works by monitoring and recording when the last byte of each +// 4k page of the `devmem_to_watch` becomes non-zero. +// @param devmem_len Length of devmem_to_watch in bytes +// @param times Arr to store fin time of each pg cpy. Len == devmem_len / PG_SZ +__global__ void watch_devmem(char* devmem_to_watch, int devmem_len, uint64_t* times, int* barrier) { + volatile char* cursor = devmem_to_watch; + // If more than one thread is launched, distribute work by offsetting threads + if (blockDim.x > 1) { + times += threadIdx.x; + cursor += threadIdx.x * PG_SZ; + } + barrier[threadIdx.x] = 1; + // Assume they show up in the right order + while (cursor < devmem_to_watch + devmem_len) { + // Wait until last byte of the page is written to + while (!*(cursor + PG_SZ - 1)) + continue; + *times = gclock64_fast(); // Log time + times += blockDim.x; + cursor += PG_SZ * blockDim.x; // Advance a page + } +} + +static inline char* human_readable_bytes(size_t bytes) { + static char out[100]; + if (bytes > 1024.*1024*1024) + snprintf(out, 100, "%.2fGiB", bytes/(1024.*1024*1024)); + else if (bytes > 1024.*1024) + snprintf(out, 100, "%.2fMiB", bytes/(1024.*1024)); + else if (bytes > 1024.) + snprintf(out, 100, "%.2fKiB", bytes/1024.); + return out; +} -- cgit v1.2.2