From: Andrey Alekseenko <al42and@gmail.com>
Date: Wed, 6 Oct 2021 14:06:15 +0000 (+0200)
Subject: Fix SYCL LINCS kernel
X-Git-Url: http://biod.pnpi.spb.ru/gitweb/?a=commitdiff_plain;h=19221c3d67046521adc8837ea594a4fa9697ae39;p=alexxy%2Fgromacs.git

Fix SYCL LINCS kernel
---

diff --git a/src/gromacs/mdlib/lincs_gpu_internal_sycl.cpp b/src/gromacs/mdlib/lincs_gpu_internal_sycl.cpp
index 3c39f43f4b..9627ce8d2c 100644
--- a/src/gromacs/mdlib/lincs_gpu_internal_sycl.cpp
+++ b/src/gromacs/mdlib/lincs_gpu_internal_sycl.cpp
@@ -154,7 +154,7 @@ auto lincsKernel(cl::sycl::handler&                   cgh,
     // shmem buffer for right-hand-side values
     auto sm_rhs = [&]() {
         return cl::sycl::accessor<float, 1, mode::read_write, target::local>(
-                cl::sycl::range<1>(c_threadsPerBlock), cgh);
+                cl::sycl::range<1>(c_threadsPerBlock * 2), cgh);
     }();
 
     // shmem buffer for virial components
@@ -223,7 +223,7 @@ auto lincsKernel(cl::sycl::handler&                   cgh,
             rc         = rlen * dx;
         }
 
-        sm_r[threadIndex] = rc;
+        sm_r[threadInBlock] = rc;
         // Make sure that all r's are saved into shared memory
         // before they are accessed in the loop below
         itemIdx.barrier(fence_space::global_and_local);