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/Makefile | 3 +-- copy_experiments/copy_testbench.h | 15 +++------------ 2 files changed, 4 insertions(+), 14 deletions(-) (limited to 'copy_experiments') 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; -- cgit v1.2.2