Add GMX_GPU_SYCL_NO_SYNCHRONIZE environment variable

diff --git a/docs/user-guide/environment-variables.rst b/docs/user-guide/environment-variables.rst
index 56ed051de9bebbf36001071635fe9be45b7750a8..ecd06e671e8c4296836a168b57254cd83795d6e5 100644 (file)
--- 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 95614359a5156034bc205db92f7986fa2878fdff..470c89aac7e915ae00bf0a5ed764deb918076d68 100644 (file)
--- 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<cl::sycl::event> 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