From a195ef37e0545ae7fbc3a9af80672c17883e0918 Mon Sep 17 00:00:00 2001 From: Andrey Alekseenko Date: Wed, 14 Apr 2021 07:14:21 +0000 Subject: [PATCH] Add GMX_GPU_SYCL_NO_SYNCHRONIZE environment variable --- docs/user-guide/environment-variables.rst | 4 ++++ .../gpu_utils/gpueventsynchronizer_sycl.h | 19 ++++++++++++++++++- 2 files changed, 22 insertions(+), 1 deletion(-) diff --git a/docs/user-guide/environment-variables.rst b/docs/user-guide/environment-variables.rst index 56ed051de9..ecd06e671e 100644 --- a/docs/user-guide/environment-variables.rst +++ b/docs/user-guide/environment-variables.rst @@ -170,6 +170,10 @@ Performance and Run Control (for coordinate and force buffers) directly on GPU memory spaces, without the staging of data through CPU memory, where possible. +``GMX_GPU_SYCL_NO_SYNCHRONIZE`` + disable synchronizations between different GPU streams in SYCL build, instead relying on SYCL runtime to + do scheduling based on data dependencies. Experimental. + ``GMX_CYCLE_ALL`` times all code during runs. Incompatible with threads. diff --git a/src/gromacs/gpu_utils/gpueventsynchronizer_sycl.h b/src/gromacs/gpu_utils/gpueventsynchronizer_sycl.h index 95614359a5..470c89aac7 100644 --- a/src/gromacs/gpu_utils/gpueventsynchronizer_sycl.h +++ b/src/gromacs/gpu_utils/gpueventsynchronizer_sycl.h @@ -81,7 +81,10 @@ class GpuEventSynchronizer { public: //! A constructor. - GpuEventSynchronizer() = default; + GpuEventSynchronizer() + { + doNotSynchronizeBetweenStreams_ = (std::getenv("GMX_GPU_SYCL_NO_SYNCHRONIZE") != nullptr); + } //! A constructor from an existing event. GpuEventSynchronizer(const cl::sycl::event& event) : event_(event) {} //! A destructor. @@ -132,6 +135,11 @@ public: */ inline void enqueueWaitEvent(const DeviceStream& deviceStream) { + if (doNotSynchronizeBetweenStreams_) + { + event_.reset(); + return; + } # if GMX_SYCL_HIPSYCL deviceStream.stream().wait_and_throw(); // SYCL-TODO: Use CUDA/HIP-specific solutions # else @@ -148,6 +156,15 @@ public: private: std::optional event_ = std::nullopt; + /*! \brief Dev. setting to no-op enqueueWaitEvent + * + * In SYCL, dependencies between the GPU tasks are managed by the runtime, so manual + * synchronization between GPU streams should be redundant, but we keep it on by default. + * + * Setting this to \c true via \c GMX_GPU_SYCL_NO_SYNCHRONIZE environment variable will + * immediately return from \ref enqueueWaitEvent, without placing a barrier into the stream. + */ + bool doNotSynchronizeBetweenStreams_; }; #endif // !defined DOXYGEN -- 2.22.0