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. --- copy_experiments/copy_testbench.h | 15 +++------------ 1 file changed, 3 insertions(+), 12 deletions(-) (limited to 'copy_experiments/copy_testbench.h') 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; -- cgit v1.2.2