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

[alexxy/gromacs.git] / src / gromacs / ewald / pme_solve_sycl.cpp

diff --git a/src/gromacs/ewald/pme_solve_sycl.cpp b/src/gromacs/ewald/pme_solve_sycl.cpp
index 633cf31e8e690b5cfacadcdb8726dcb5699a5156..5a0125c44d286ddcc71bc29f6dd5d80fd5024388 100644 (file)
--- a/src/gromacs/ewald/pme_solve_sycl.cpp
+++ b/src/gromacs/ewald/pme_solve_sycl.cpp
@@ -65,13 +65,16 @@ template<GridOrdering gridOrdering, bool computeEnergyAndVirial, int subGroupSiz
 auto makeSolveKernel(cl::sycl::handler&                            cgh,
                      DeviceAccessor<float, mode::read>             a_splineModuli,
                      DeviceAccessor<SolveKernelParams, mode::read> a_solveKernelParams,
-                     DeviceAccessor<float, mode::read_write>       a_virialAndEnergy,
-                     DeviceAccessor<float, mode::read_write>       a_fourierGrid)
+                     OptionalAccessor<float, mode::read_write, computeEnergyAndVirial> a_virialAndEnergy,
+                     DeviceAccessor<float, mode::read_write> a_fourierGrid)
 {
-    cgh.require(a_splineModuli);
-    cgh.require(a_solveKernelParams);
-    cgh.require(a_virialAndEnergy);
-    cgh.require(a_fourierGrid);
+    a_splineModuli.bind(cgh);
+    a_solveKernelParams.bind(cgh);
+    if constexpr (computeEnergyAndVirial)
+    {
+        a_virialAndEnergy.bind(cgh);
+    }
+    a_fourierGrid.bind(cgh);
 
     /* Reduce 7 outputs per warp in the shared memory */
     const int stride =
@@ -281,7 +284,7 @@ auto makeSolveKernel(cl::sycl::handler&                            cgh,
         }
 
         /* Optional energy/virial reduction */
-        if (computeEnergyAndVirial)
+        if constexpr (computeEnergyAndVirial)
         {
             /* A tricky shuffle reduction inspired by reduce_force_j_warp_shfl.
              * The idea is to reduce 7 energy/virial components into a single variable (aligned by