5 files changed, 44 insertions, 15 deletions

diff --git a/constant_cycles_kernel.cu b/constant_cycles_kernel.cu
index eda963d..db92b08 100644
--- a/constant_cycles_kernel.cu
+++ b/constant_cycles_kernel.cu
@@ -13,7 +13,6 @@ __global__ void loop_on_gpu(unsigned long iters, int *__unused) {
 }
 
 int main(int argc, char **argv) {
-        cudaError_t err;
         int res, *__unused;
         struct timespec start, end;
 
diff --git a/copy_experiments/mon_cross_ctx_copies.cu b/copy_experiments/mon_cross_ctx_copies.cu
index d9749e0..93d0e4c 100644
--- a/copy_experiments/mon_cross_ctx_copies.cu
+++ b/copy_experiments/mon_cross_ctx_copies.cu
@@ -92,7 +92,6 @@ void cpu_copy_mon(int loops, char* cpu_mem) {
 void* copy_thread(void* args_raw) {
         CUcontext ctx;
         int dev, i;
-        cudaError_t err;
         int *barrier;
         int *barrier_dev_ptr; // On Pascal+ this will be the same as *barrier
         char *pinned_hostmem, *devmem;
@@ -106,8 +105,8 @@ void* copy_thread(void* args_raw) {
         // Explictly create a context (avoids creating a primary context implictly)
         // This has been verified on CUDA 11.1 to give each thread a different context
         // handle
-        cudaGetDevice(&dev);
-        SAFE(cuCtxCreate(&ctx, 0, dev));
+        SAFE(cudaGetDevice(&dev));
+        SAFE_D(cuCtxCreate(&ctx, 0, dev));
 
         uint64_t dev_ns, dev_ns2;
         double host_s, host_s2;
@@ -249,21 +248,27 @@ static error_t arg_parser(int key, char* arg, struct argp_state *state) {
 }
 
 int main(int argc, char**argv) {
-        int tmp, dev;
+        int tmp, dev = 0;
+        CUdevice dev_itrl;
         uint64_t *ctx_times[MAX_THREADS] = {0};
         pthread_t t[MAX_THREADS];
         global_args_t g_args = {0};
 
+        // Temporarially initialize CUDA to query device attributes
+        SAFE_D(cuInit(0));
+        SAFE_D(cuDeviceGet(&dev_itrl, dev));
         // Due to some laziness in how we handle barriers, this flag needs to be true
         /// XXX: Still seems to work fine if it isn't???
-        cudaGetDevice(&dev);
-        cudaDeviceGetAttribute(&tmp, cudaDevAttrCanMapHostMemory, dev);
+        SAFE_D(cuDeviceGetAttribute(&tmp, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, dev_itrl));
         if (!tmp) {
                 fprintf(stderr, "Unsupported platform. It must be possible to map host (CPU)"
-                                "DRAM into the GPU virtual address space for accurate clock"
+                                " DRAM into the GPU virtual address space for accurate clock "
                                 "synchronization. Exiting...\n");
                 return 1;
         }
+        // Terminate the context used for attrib check so it's not accidentially
+        // reused in subprocesses
+        SAFE_D(cuDevicePrimaryCtxRelease(dev_itrl));
 
         struct argp argp = {opts, arg_parser, 0, desc};
         argp_parse(&argp, argc, argv, 0, 0, &g_args);
diff --git a/preemption_logger.cu b/preemption_logger.cu
index 667c459..d93600e 100644
--- a/preemption_logger.cu
+++ b/preemption_logger.cu
@@ -27,7 +27,7 @@ struct interval {
 };
 
 // Minimum time discontinuity which indicates a gap between intervals
-// One clock tick is about 1ns
+// Clock resolution is about 1ns, but it only ticks every 1us pre-H100
 #define MIN_PREEMPT_TICKS 2*1000 // ~2us
 
 // Watch for discontinuities in the GPU clock, indicating intervals during
@@ -67,7 +67,6 @@ Spin on the GPU, logging intervals to stdout during which we are scheduled.\n\
   -r, --raw       Print raw logged GPU times (skip conversion to CPU time).\n";
 
 int main(int argc, char **argv) {
-        cudaError_t err; // Needed for the SAFE() macro
         struct interval *ivls_gpu, *ivls;
         struct timespec start, end, end_ivls_only;
         int num_ivls, skip_conversion;
diff --git a/task_host_utilities.cu b/task_host_utilities.cu
index 9bb95aa..1f76080 100644
--- a/task_host_utilities.cu
+++ b/task_host_utilities.cu
@@ -4,6 +4,7 @@
 // been copied and heavily modified for use in the gpu-microbench project.
 #include <cuda_runtime.h>
 #include <errno.h>
+#include <sched.h>
 #include <stdint.h>
 #include <stdio.h>
 #include <string.h>
@@ -191,8 +192,8 @@ static void InternalReadGPUNanoseconds(int cuda_device, double *cpu_time,
   if (!CheckCUDAError(cudaMemcpy(gpu_time, device_time, sizeof(device_time),
     cudaMemcpyDeviceToHost))) goto out;
   max_error = (cpu_end - cpu_start) / 2.0;
-  fprintf(stderr, "Time synchronized to a maximum error of +/- %f us.\n",
-    max_error * (1000.0 * 1000.0));
+  fprintf(stderr, "Time synchronized to a maximum error of +/- %f us on CPU%d.\n",
+    max_error * (1000.0 * 1000.0), sched_getcpu());
 out:
   cudaFree(device_time);
   cudaFree((void*)start_barrier);
@@ -200,6 +201,7 @@ out:
   cudaFree((void*)ready_barrier);
 }
 
+// Returns 0 on success, 1 on error
 int GetHostDeviceTimeOffset(int cuda_device, double *host_seconds,
   uint64_t *gpu_nanoseconds) {
   uint64_t *shared_gpu_time = NULL;
diff --git a/testbench.h b/testbench.h
index 5b47bd4..5e77410 100644
--- a/testbench.h
+++ b/testbench.h
@@ -2,10 +2,34 @@
  * Header for miscellaneous experimental helper functions.
  */
 
+// cudaError_t and CUResult can both safely be cast to an unsigned int
+static __thread unsigned int __SAFE_err;
+
+// The very strange cast in these macros is to satisfy two goals at tension:
+// 1. This file should be able to be included in non-CUDA-using files, and thus
+//    should use no CUDA types outside of this macro.
+// 2. We want to typecheck uses of these macros. The driver and runtime APIs
+//    do not have identical error numbers and/or meanings, so runtime library
+//    calls should use SAFE, and driver library calls should use SAFE_D.
+// Our design allows typechecking while keeping a non-CUDA per-thread error var.
+
+// For CUDA Runtime Library functions; typically those prefixed with `cuda`
 #define SAFE(x) \
-        if ((err = (cudaError_t)(x)) != 0) { \
-                printf("CUDA error %d! %s\n", err, cudaGetErrorString(err)); \
-                printf("Suspect line: %s\n", #x); \
+        if ((*(cudaError_t*)(&__SAFE_err) = (x)) != 0) { \
+                printf("(%s:%d) CUDA error %d: %s i.e. \"%s\" returned by %s. Aborting...\n", \
+                       __FILE__, __LINE__, __SAFE_err, cudaGetErrorName((cudaError_t)__SAFE_err), cudaGetErrorString((cudaError_t)__SAFE_err), #x); \
+                exit(1); \
+        }
+
+// For CUDA Driver Library functions; typically those prefixed with just `cu`
+#define SAFE_D(x) \
+        if ((*(CUresult*)&(__SAFE_err) = (x)) != 0) { \
+                const char* name; \
+                const char* desc; \
+                cuGetErrorName((CUresult)__SAFE_err, &name); \
+                cuGetErrorString((CUresult)__SAFE_err, &desc); \
+                printf("(%s:%d) CUDA error %d: %s i.e. \"%s\" returned by %s. Aborting...\n", \
+                       __FILE__, __LINE__, __SAFE_err, name, desc, #x); \
                 exit(1); \
         }