SYCL: Fully switch to atomic_ref

[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..6a32e856c0dd3e60a311e15567151567dcd159a5 100644 (file)
--- a/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp
+++ b/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp
@@ -69,7 +69,7 @@ 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);
@@ -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]);
             }
         }
     };