SYCL: Use acc.bind(cgh) instead of cgh.require(acc)

[alexxy/gromacs.git] / src / gromacs / mdlib / gpuforcereduction_impl_internal_sycl.cpp

diff --git a/src/gromacs/mdlib/gpuforcereduction_impl_internal_sycl.cpp b/src/gromacs/mdlib/gpuforcereduction_impl_internal_sycl.cpp
index e2ee7b46979d1bd12dccf8600faeed8f19553e51..b8a2bcbb0ca66f2f458080546ccd30ec08b3b91a 100644 (file)
--- a/src/gromacs/mdlib/gpuforcereduction_impl_internal_sycl.cpp
+++ b/src/gromacs/mdlib/gpuforcereduction_impl_internal_sycl.cpp
@@ -51,7 +51,7 @@
 #include "gromacs/gpu_utils/gmxsycl.h"
 #include "gromacs/gpu_utils/devicebuffer.h"
 #include "gromacs/gpu_utils/gpu_utils.h"
-#include "gromacs/gpu_utils/gpueventsynchronizer_sycl.h"
+#include "gromacs/gpu_utils/gpueventsynchronizer.h"
 #include "gromacs/utility/template_mp.h"
 
 //! \brief Class name for reduction kernel
@@ -63,6 +63,7 @@ namespace gmx
 
 using cl::sycl::access::mode;
 
+//! \brief Function returning the force reduction kernel lambda.
 template<bool addRvecForce, bool accumulateForce>
 static auto reduceKernel(cl::sycl::handler&                                 cgh,
                          DeviceAccessor<Float3, mode::read>                 a_nbnxmForce,
@@ -71,13 +72,13 @@ static auto reduceKernel(cl::sycl::handler&                                 cgh,
                          DeviceAccessor<int, cl::sycl::access::mode::read> a_cell,
                          const int                                         atomStart)
 {
-    cgh.require(a_nbnxmForce);
+    a_nbnxmForce.bind(cgh);
     if constexpr (addRvecForce)
     {
-        cgh.require(a_rvecForceToAdd);
+        a_rvecForceToAdd.bind(cgh);
     }
-    cgh.require(a_forceTotal);
-    cgh.require(a_cell);
+    a_forceTotal.bind(cgh);
+    a_cell.bind(cgh);
 
     return [=](cl::sycl::id<1> itemIdx) {
         // Set to nbnxnm force, then perhaps accumulate further to it
@@ -97,6 +98,7 @@ static auto reduceKernel(cl::sycl::handler&                                 cgh,
     };
 }
 
+//! \brief Force reduction SYCL kernel launch code.
 template<bool addRvecForce, bool accumulateForce>
 static void launchReductionKernel_(const int                   numAtoms,
                                    const int                   atomStart,
@@ -119,7 +121,7 @@ static void launchReductionKernel_(const int                   numAtoms,
     });
 }
 
-/*! \brief Select templated kernel and launch it. */
+/*! \brief Select templated Force reduction kernel and launch it. */
 void launchForceReductionKernel(int                  numAtoms,
                                 int                  atomStart,
                                 bool                 addRvecForce,