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
/*
* This file is part of the GROMACS molecular simulation package.
*
/*
* 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.
* 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.
GMX_ASSERT(stat == cudaSuccess,
("cudaEventSynchronize failed. " + gmx::getDeviceErrorString(stat)).c_str());
}
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)
{
/*! \brief Enqueues a wait for the recorded event in stream \p stream */
inline void enqueueWaitEvent(const DeviceStream& deviceStream)
{
/*
* This file is part of the GROMACS molecular simulation package.
*
/*
* 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.
* 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.
+ /*! \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
/*! \brief Enqueues a wait for the recorded event in stream \p stream
*
* After enqueue, the associated event is released, so this method should
/*
* This file is part of the GROMACS molecular simulation package.
*
/*
* 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.
* 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.
event_->wait_and_throw();
event_.reset();
}
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.
*/
/*! \brief Enqueues a wait for the recorded event in stream \p deviceStream.
* As in the OpenCL implementation, the event is released.
*/