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

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

diff --git a/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp b/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp
index 3e5b56fdf309e8cd8f8c46f190b42486d3bca814..2f7ae81b4c3b1a9821f96fb2505bd73f8cefc9ab 100644 (file)
--- a/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp
+++ b/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp
@@ -69,19 +69,19 @@ auto settleKernel(cl::sycl::handler&                                           c
                   DeviceAccessor<Float3, mode::read_write>                     a_xp,
                   float                                                        invdt,
                   OptionalAccessor<Float3, mode::read_write, updateVelocities> a_v,
-                  OptionalAccessor<float, mode_atomic, computeVirial>          a_virialScaled,
+                  OptionalAccessor<float, mode::read_write, computeVirial>     a_virialScaled,
                   PbcAiuc                                                      pbcAiuc)
 {
-    cgh.require(a_settles);
-    cgh.require(a_x);
-    cgh.require(a_xp);
+    a_settles.bind(cgh);
+    a_x.bind(cgh);
+    a_xp.bind(cgh);
     if constexpr (updateVelocities)
     {
-        cgh.require(a_v);
+        a_v.bind(cgh);
     }
     if constexpr (computeVirial)
     {
-        cgh.require(a_virialScaled);
+        a_virialScaled.bind(cgh);
     }
 
     // shmem buffer for i x+q pre-loading
@@ -340,7 +340,7 @@ auto settleKernel(cl::sycl::handler&                                           c
             // First 6 threads in the block add the 6 components of virial to the global memory address
             if (threadIdx < 6)
             {
-                atomicFetchAdd(a_virialScaled, threadIdx, sm_threadVirial[threadIdx * blockSize]);
+                atomicFetchAdd(a_virialScaled[threadIdx], sm_threadVirial[threadIdx * blockSize]);
             }
         }
     };