diff options
Diffstat (limited to 'copy_experiments')
| -rw-r--r-- | copy_experiments/Makefile | 3 | ||||
| -rw-r--r-- | copy_experiments/copy_testbench.h | 15 |
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 | 1 | copy_contender: copy_contender.cu copy_testbench.h ../testbench.h | |
| 2 | copy_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 |
| 23 | static __device__ inline uint64_t gclock64_fast(void) { | 14 | static __device__ inline uint64_t gclock64_fast(void) { |
| 24 | uint32_t lo_bits; | 15 | uint32_t lo_bits; |
