From: Szilárd Páll <pall.szilard@gmail.com>
Date: Thu, 11 Feb 2021 10:57:05 +0000 (+0100)
Subject: Support checking event status in GpuEventSynchronizer
X-Git-Url: http://biod.pnpi.spb.ru/gitweb/?a=commitdiff_plain;h=2dd46f2842059bf9834547ecceef88872ff394d7;p=alexxy%2Fgromacs.git

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 fee73cdb04..b098f35498 100644
--- 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 128e0564dd..2b42308e84 100644
--- 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 1daa5b9f32..6740cc3a3d 100644
--- 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.
      */