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. --- constant_cycles_kernel.cu | 56 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 56 insertions(+) create mode 100644 constant_cycles_kernel.cu (limited to 'constant_cycles_kernel.cu') 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; +} -- cgit v1.2.2