Support checking event status in GpuEventSynchronizer

This commit adds an isReady() method that allows checking whether
the event has been triggered (or in SYCL associated barrier task has
completed).
The primary aim is porting the nbnxm module to the unified GPU layer.

Implements part of #2527

diff --git a/src/gromacs/gpu_utils/gpueventsynchronizer.cuh b/src/gromacs/gpu_utils/gpueventsynchronizer.cuh
index fee73cdb04b9c77dc503b97cac777afdeef4d822..b098f3549857c9394ba80c003cdb0ac2c103fb48 100644 (file)
--- a/src/gromacs/gpu_utils/gpueventsynchronizer.cuh
+++ b/src/gromacs/gpu_utils/gpueventsynchronizer.cuh
@@ -1,7 +1,7 @@
 /*
  * This file is part of the GROMACS molecular simulation package.
  *
- * Copyright (c) 2018,2019,2020, by the GROMACS development team, led by
+ * Copyright (c) 2018,2019,2020,2021, by the GROMACS development team, led by
  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
  * and including many others, as listed in the AUTHORS file in the
  * top-level source directory and at http://www.gromacs.org.
@@ -99,6 +99,14 @@ public:
         GMX_ASSERT(stat == cudaSuccess,
                    ("cudaEventSynchronize failed. " + gmx::getDeviceErrorString(stat)).c_str());
     }
+    /*! \brief Checks the completion of the underlying event and resets the object if it was. */
+    inline bool isReady()
+    {
+        cudaError_t gmx_used_in_debug stat = cudaEventQuery(event_);
+        GMX_ASSERT((stat == cudaSuccess) || (stat == cudaErrorNotReady),
+                   ("cudaEventQuery failed. " + gmx::getDeviceErrorString(stat)).c_str());
+        return (stat == cudaSuccess);
+    }
     /*! \brief Enqueues a wait for the recorded event in stream \p stream */
     inline void enqueueWaitEvent(const DeviceStream& deviceStream)
     {

diff --git a/src/gromacs/gpu_utils/gpueventsynchronizer_ocl.h b/src/gromacs/gpu_utils/gpueventsynchronizer_ocl.h
index 128e0564ddbfdaf6e4624f0e3f75cbdbb13b0ea7..2b42308e841006f07c5c8a96d26e329944d61018 100644 (file)
--- a/src/gromacs/gpu_utils/gpueventsynchronizer_ocl.h
+++ b/src/gromacs/gpu_utils/gpueventsynchronizer_ocl.h
@@ -1,7 +1,7 @@
 /*
  * This file is part of the GROMACS molecular simulation package.
  *
- * Copyright (c) 2018,2019,2020, by the GROMACS development team, led by
+ * Copyright (c) 2018,2019,2020,2021, by the GROMACS development team, led by
  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
  * and including many others, as listed in the AUTHORS file in the
  * top-level source directory and at http://www.gromacs.org.
@@ -108,6 +108,23 @@ public:
 
         releaseEvent();
     }
+    /*! \brief Checks the completion of the underlying event and resets the object if it was. */
+    inline bool isReady()
+    {
+        cl_int result;
+        cl_int clError = clGetEventInfo(
+                event_, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &result, nullptr);
+        if (CL_SUCCESS != clError)
+        {
+            GMX_THROW(gmx::InternalError("Failed to retrieve event info: " + ocl_get_error_string(clError)));
+        }
+        bool hasTriggered = (result == CL_COMPLETE);
+        if (hasTriggered)
+        {
+            releaseEvent();
+        }
+        return hasTriggered;
+    }
     /*! \brief Enqueues a wait for the recorded event in stream \p stream
      *
      *  After enqueue, the associated event is released, so this method should

diff --git a/src/gromacs/gpu_utils/gpueventsynchronizer_sycl.h b/src/gromacs/gpu_utils/gpueventsynchronizer_sycl.h
index 1daa5b9f32db44c215efa5df3b60c36b2fd4808b..6740cc3a3dad8d2fed5ecdd6012baee1e0a98c03 100644 (file)
--- a/src/gromacs/gpu_utils/gpueventsynchronizer_sycl.h
+++ b/src/gromacs/gpu_utils/gpueventsynchronizer_sycl.h
@@ -1,7 +1,7 @@
 /*
  * This file is part of the GROMACS molecular simulation package.
  *
- * Copyright (c) 2020, by the GROMACS development team, led by
+ * Copyright (c) 2020,2021, by the GROMACS development team, led by
  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
  * and including many others, as listed in the AUTHORS file in the
  * top-level source directory and at http://www.gromacs.org.
@@ -110,6 +110,17 @@ public:
         event_->wait_and_throw();
         event_.reset();
     }
+    /*! \brief Checks the completion of the underlying event and resets the object if it was. */
+    inline bool isReady()
+    {
+        auto info         = event_->get_info<cl::sycl::info::event::command_execution_status>();
+        bool hasTriggered = (info == cl::sycl::info::event_command_status::complete);
+        if (hasTriggered)
+        {
+            event_.reset();
+        }
+        return hasTriggered;
+    }
     /*! \brief Enqueues a wait for the recorded event in stream \p deviceStream.
      * As in the OpenCL implementation, the event is released.
      */