From a4450b292e0e69c0079d494710cf782dff89277c Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Tue, 11 Jul 2023 17:21:47 -0700 Subject: [PATCH] Remove explicit device synchronization --- src/chai/ArrayManager.hpp | 17 ----------------- src/chai/managed_ptr.hpp | 24 +----------------------- 2 files changed, 1 insertion(+), 40 deletions(-) diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 8963cfd3..eca8848e 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -396,23 +396,6 @@ class ArrayManager */ void disableCallbacks() { m_callbacks_active = false; } - /*! - * \brief Turn on device synchronization after every kernel. - */ - [[deprecated("Set the CUDA_LAUNCH_BLOCKING or HIP_LAUNCH_BLOCKING environment variable instead")]] - void enableDeviceSynchronize() { m_device_synchronize = true; } - - /*! - * \brief Turn off device synchronization after every kernel. - */ - [[deprecated("Set the CUDA_LAUNCH_BLOCKING or HIP_LAUNCH_BLOCKING environment variable instead")]] - void disableDeviceSynchronize() { m_device_synchronize = false; } - - /*! - * \brief Turn on device synchronization after every kernel. - */ - bool deviceSynchronize() { return m_device_synchronize; } - /*! * \brief synchronize the device if there hasn't been a synchronize since the last kernel */ diff --git a/src/chai/managed_ptr.hpp b/src/chai/managed_ptr.hpp index 67fbeedd..a527437c 100644 --- a/src/chai/managed_ptr.hpp +++ b/src/chai/managed_ptr.hpp @@ -101,7 +101,7 @@ namespace chai { /// Be aware that CHAI checks every CUDA API call for GPU errors by default. To /// turn off GPU error checking, pass -DCHAI_ENABLE_GPU_ERROR_CHECKING=OFF as /// an argument to cmake when building CHAI. To turn on synchronization after - /// every kernel, call ArrayManager::getInstance()->enableDeviceSynchronize(). + /// every kernel, set the appropriate environment variable (e.g. CUDA_LAUNCH_BLOCKING or HIP_LAUNCH_BLOCKING). /// Alternatively, call cudaDeviceSynchronize() after any call to make_managed, /// make_managed_from_factory, or managed_ptr::free, and check the return code /// for errors. If your code crashes in the constructor/destructor of T, then @@ -502,11 +502,6 @@ namespace chai { if (pointer) { destroy_on_device(temp); m_gpu_pointer = nullptr; -#ifndef CHAI_DISABLE_RM - if (ArrayManager::getInstance()->deviceSynchronize()) { - synchronize(); - } -#endif } break; @@ -535,11 +530,6 @@ namespace chai { if (pointer) { destroy_on_device(pointer); m_gpu_pointer = nullptr; -#ifndef CHAI_DISABLE_RM - if (ArrayManager::getInstance()->deviceSynchronize()) { - synchronize(); - } -#endif } break; @@ -964,12 +954,6 @@ namespace chai { hipLaunchKernelGGL(detail::make_on_device, 1, 1, 0, 0, gpuBuffer, args...); #endif -#ifndef CHAI_DISABLE_RM - if (ArrayManager::getInstance()->deviceSynchronize()) { - synchronize(); - } -#endif - // Allocate space on the CPU for the pointer and copy the pointer to the CPU T** cpuBuffer = (T**) malloc(sizeof(T*)); gpuMemcpy(cpuBuffer, gpuBuffer, sizeof(T*), gpuMemcpyDeviceToHost); @@ -1025,12 +1009,6 @@ namespace chai { hipLaunchKernelGGL(detail::make_on_device_from_factory, 1, 1, 0, 0, gpuBuffer, f, args...); #endif -#ifndef CHAI_DISABLE_RM - if (ArrayManager::getInstance()->deviceSynchronize()) { - synchronize(); - } -#endif - // Allocate space on the CPU for the pointer and copy the pointer to the CPU T** cpuBuffer = (T**) malloc(sizeof(T*)); gpuMemcpy(cpuBuffer, gpuBuffer, sizeof(T*), gpuMemcpyDeviceToHost);