From 717f5eaeb0e5ffbfe9372dbca041bd751c95410d Mon Sep 17 00:00:00 2001 From: Joshua Bakita Date: Wed, 21 Jun 2023 11:40:00 -0400 Subject: Fix backwards compatibility and checks in mon_cross_ctx_copies On-GPU copy monitoring may have been disfunctional on Maxwell-. --- copy_experiments/mon_cross_ctx_copies.cu | 19 ++++++++++++------- 1 file changed, 12 insertions(+), 7 deletions(-) (limited to 'copy_experiments') 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++) { -- cgit v1.2.2