summaryrefslogtreecommitdiffstats
path: root/copy_experiments
diff options
context:
space:
mode:
Diffstat (limited to 'copy_experiments')
-rw-r--r--copy_experiments/Makefile3
-rw-r--r--copy_experiments/copy_testbench.h15
2 files changed, 4 insertions, 14 deletions
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 @@
1 1copy_contender: copy_contender.cu copy_testbench.h ../testbench.h
2copy_contender: copy_contender.cu copy_testbench.h
3 /usr/local/cuda/bin/nvcc copy_contender.cu -o $@ -lcuda -g 2 /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 @@
1/* Copyright 2023 Joshua Bakita 1/* Copyright 2021-2023 Joshua Bakita
2 * Header-only library that provides an ultra-low-overhead GPU timing function, 2 * Header-only library that provides an ultra-low-overhead GPU timing function,
3 * a copy monitoring kernel, and other useful sundries. 3 * a copy monitoring kernel, and other useful sundries.
4 */ 4 */
5#include <stdio.h> 5#include <stdio.h>
6#include <stdint.h> 6#include <stdint.h>
7 7
8#include "../testbench.h"
9
8// 4k pages (4096 bytes) 10// 4k pages (4096 bytes)
9#define PG_SZ 0x1000 11#define PG_SZ 0x1000
10 12
11#define SAFE(x) \
12 if ((err = (cudaError_t)(x)) != 0) { \
13 printf("CUDA error %d! %s\n", err, cudaGetErrorString(err)); \
14 printf("Suspect line: %s\n", #x); \
15 exit(1); \
16 }
17
18#define s2ns(s) ((s)*1000l*1000l*1000l)
19#define ns2ms(s) ((s)/(1000l*1000l))
20#define timediff(start, end) ((s2ns((end).tv_sec) + (end).tv_nsec) - (s2ns((start).tv_sec) + (start).tv_nsec))
21
22// Compiles to 5 inline SASS instructions on sm_52 13// Compiles to 5 inline SASS instructions on sm_52
23static __device__ inline uint64_t gclock64_fast(void) { 14static __device__ inline uint64_t gclock64_fast(void) {
24 uint32_t lo_bits; 15 uint32_t lo_bits;