1 files changed, 12 insertions, 7 deletions

diff --git a/copy_experiments/mon_cross_ctx_copies.cu b/copy_experiments/mon_cross_ctx_copies.cu
index 65ccd4e..dfccdfc 100644
--- a/copy_experiments/mon_cross_ctx_copies.cu
+++ b/copy_experiments/mon_cross_ctx_copies.cu
@@ -88,6 +88,7 @@ void* copy_thread(void* args_raw) {
         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;
         uint64_t *devtimes, *times;
         cudaStream_t stream1, stream2;
@@ -114,9 +115,10 @@ void* copy_thread(void* args_raw) {
         if (GPU_COMPUTE) {
                 SAFE(cudaHostAlloc(&barrier, sizeof(int) * GPU_MON_THREADS, cudaHostAllocMapped));
                 SAFE(cudaMemset(barrier, 0, sizeof(int) * GPU_MON_THREADS));
+                SAFE(cudaHostGetDevicePointer(&barrier_dev_ptr, barrier, 0));
         }
 
-        SAFE(cudaMallocHost(&pinned_hostmem, COPY_SIZE));
+        SAFE(cudaHostAlloc(&pinned_hostmem, COPY_SIZE, cudaHostAllocDefault));
         for (i = 0; i < COPY_SIZE; i++)
                 // Don't allow 0 so that the copy detection logic works
                 pinned_hostmem[i] = max((rand() & 0xff), 1);
@@ -132,7 +134,7 @@ void* copy_thread(void* args_raw) {
 
         if (GPU_COMPUTE) {
                 // 1 thread can't go quick enough, even on the GV100
-                watch_devmem<<<1,GPU_MON_THREADS,0,stream2>>>(devmem, COPY_SIZE, devtimes, barrier);
+                watch_devmem<<<1,GPU_MON_THREADS,0,stream2>>>(devmem, COPY_SIZE, devtimes, barrier_dev_ptr);
                 // Wait for monitor(s) to initialize
                 bool ready = false;
                 while (!ready) {
@@ -249,11 +251,13 @@ int main(int argc, char**argv) {
         // 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, cudaDevAttrCanUseHostPointerForRegisteredMem, dev);
-        /*if (!tmp) {
-                fprintf(stderr, "Unsupported platform. Exiting...\n");
-                exit(1);
-        }*/
+        cudaDeviceGetAttribute(&tmp, cudaDevAttrCanMapHostMemory, dev);
+        if (!tmp) {
+                fprintf(stderr, "Unsupported platform. It must be possible to map host (CPU)"
+                                "DRAM into the GPU virtual address space for accurate clock"
+                                "synchronization. Exiting...\n");
+                return 1;
+        }
 
         struct argp argp = {opts, arg_parser, 0, desc};
         argp_parse(&argp, argc, argv, 0, 0, &g_args);
@@ -272,6 +276,7 @@ int main(int argc, char**argv) {
         InternalReadGPUNanoseconds(dev, &host_s, &dev_ns);
         GPU_TIME_OFFSET = dev_ns - s2ns(host_s);
 
+        // Copy buffers are filled with random numbers. Seed the RNG.
         srand(0);
         // Spawn and wait for children to intialize
         for (int tid = 0; tid < g_args.num_threads; tid++) {