From 5ba80e14efb595688a3dc3d49c212bc393840513 Mon Sep 17 00:00:00 2001 From: Joshua Bakita Date: Fri, 9 Jun 2023 19:38:01 -0400 Subject: Add constant_cycles_kernel benchmark + header split and Makefile Benchmark supports: => Looping on the GPU for a specified number of iterations => Looping infinitely on the GPU => Logging the sum time (ns and ms) for all iterations to complete Also splits `testbench.h` out of `copy_testbench.h` for use in other benchmarks. --- Makefile | 7 ++++- constant_cycles_kernel.cu | 56 +++++++++++++++++++++++++++++++++++++++ copy_experiments/Makefile | 3 +-- copy_experiments/copy_testbench.h | 15 +++-------- testbench.h | 14 ++++++++++ 5 files changed, 80 insertions(+), 15 deletions(-) create mode 100644 constant_cycles_kernel.cu create mode 100644 testbench.h diff --git a/Makefile b/Makefile index e81353d..7eaf03d 100644 --- a/Makefile +++ b/Makefile @@ -1,4 +1,9 @@ -.PHONY: copy_experiments +.PHONY: all copy_experiments + +all: constant_cycles_kernel copy_experiments + +constant_cycles_kernel: constant_cycles_kernel.cu testbench.h + nvcc constant_cycles_kernel.cu -o constant_cycles_kernel copy_experiments: $(MAKE) -C $@ diff --git a/constant_cycles_kernel.cu b/constant_cycles_kernel.cu new file mode 100644 index 0000000..eda963d --- /dev/null +++ b/constant_cycles_kernel.cu @@ -0,0 +1,56 @@ +/* Copyright 2021-2023 Joshua Bakita + * Simple kernel that spins on the GPU for a specified number of iterations, + * while tracking and printing the necessary CPU time. + */ +#include +#include +#include + +#include "testbench.h" + +__global__ void loop_on_gpu(unsigned long iters, int *__unused) { + for (volatile int i = 0; i < iters; i++) (*__unused)++; +} + +int main(int argc, char **argv) { + cudaError_t err; + int res, *__unused; + struct timespec start, end; + + if (argc != 2) { + fprintf(stderr, "Usage: %s <# of millions of iterations, or -1 for infinite>\n", + argv[0]); + return 1; + } + + // Input is multiplied by one million, unless infinite + unsigned long num_iters = strtoul(argv[1], NULL, 10); + if (num_iters != (unsigned long)(-1)) + num_iters *= 1000 * 1000; + + // Initialize CUDA and a context (hack) + SAFE(cudaMalloc(&__unused, 8)); + + // Run iterations on a single thread + clock_gettime(CLOCK_MONOTONIC_RAW, &start); + loop_on_gpu<<<1,1>>>(num_iters, __unused); + SAFE(cudaGetLastError() /* Check successful launch */); + SAFE(cudaDeviceSynchronize()); + clock_gettime(CLOCK_MONOTONIC_RAW, &end); + + // Print detailed timing information + long elapsed = timediff(start, end); + fprintf(stderr, "Started at %ld ns, ended at %ld ns\n", + s2ns(start.tv_sec) + start.tv_nsec, s2ns(end.tv_sec) + end.tv_nsec); + fprintf(stderr, "%ld ns (%.2f ms) elapsed\n", elapsed, elapsed / (1000 * 1000.)); + + // Verify success (fool optimizer) + SAFE(cudaMemcpy(&res, __unused, 8, cudaMemcpyDeviceToHost)); + // Theoretically this can happen if `__unused` wraps around; maybe for a + // very small `long` type, or a very long run. More likely indicates an error. + if (!res) + fprintf(stderr, "CRITICAL: Zero iterations seem completed. Likely incorrect " + "arguments, internal error, or corruption of CUDA internal state.\n"); + + return 0; +} diff --git a/copy_experiments/Makefile b/copy_experiments/Makefile index 2358c60..b6b5a38 100644 --- a/copy_experiments/Makefile +++ b/copy_experiments/Makefile @@ -1,3 +1,2 @@ - -copy_contender: copy_contender.cu copy_testbench.h +copy_contender: copy_contender.cu copy_testbench.h ../testbench.h /usr/local/cuda/bin/nvcc copy_contender.cu -o $@ -lcuda -g diff --git a/copy_experiments/copy_testbench.h b/copy_experiments/copy_testbench.h index e1c4aa5..fed62ba 100644 --- a/copy_experiments/copy_testbench.h +++ b/copy_experiments/copy_testbench.h @@ -1,24 +1,15 @@ -/* Copyright 2023 Joshua Bakita +/* Copyright 2021-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 +#include "../testbench.h" + // 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; diff --git a/testbench.h b/testbench.h new file mode 100644 index 0000000..28d7f77 --- /dev/null +++ b/testbench.h @@ -0,0 +1,14 @@ +/* Copyright 2021-2023 Joshua Bakita + * Header for miscellaneous experimental helper functions. + */ + +#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)) -- cgit v1.2.2