From 19221c3d67046521adc8837ea594a4fa9697ae39 Mon Sep 17 00:00:00 2001 From: Andrey Alekseenko Date: Wed, 6 Oct 2021 16:06:15 +0200 Subject: [PATCH] Fix SYCL LINCS kernel --- src/gromacs/mdlib/lincs_gpu_internal_sycl.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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( - 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); -- 2.22.0