From f0ed971d47b3a48fc796adb9f36f0bfb5225026a Mon Sep 17 00:00:00 2001 From: Tom Lin Date: Tue, 13 Aug 2024 11:55:40 +0100 Subject: [PATCH] Disable possible UB use of async_work_group_copy in SYCL --- src/sycl/fasten.hpp | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/src/sycl/fasten.hpp b/src/sycl/fasten.hpp index 74850f0..b99fff4 100644 --- a/src/sycl/fasten.hpp +++ b/src/sycl/fasten.hpp @@ -55,10 +55,13 @@ template class IMPL_CLS final : public Bude { // XXX async_work_group_copy takes only gentypes, so no FFParams, // casting *_ptr parameter requires first converting to void and then to gentype // although probably free, there must be a better way of doing this - sycl::device_event event = item.async_work_group_copy( - sycl::local_ptr(sycl::local_ptr(local_forcefield.get_pointer())), - sycl::global_ptr(sycl::global_ptr(forcefields.get_pointer())), - ntypes * sizeof(FFParams) / sizeof(float)); + // sycl::device_event event = item.async_work_group_copy( + // sycl::local_ptr(sycl::local_ptr(local_forcefield.get_pointer())), + // sycl::global_ptr(sycl::global_ptr(forcefields.get_pointer())), + // ntypes * sizeof(FFParams) / sizeof(float)); + + for (int i = lid; i < ntypes; i += lrange) + local_forcefield[i] = forcefields[i]; // Compute transformation matrix to private memory const size_t lsz = lrange; @@ -88,7 +91,8 @@ template class IMPL_CLS final : public Bude { etot[i] = ZERO; } - item.wait_for(event); + // item.wait_for(event); + item.barrier(sycl::access::fence_space::local_space); // Loop over ligand atoms for (size_t il = 0; il < ligands.get_count(); il++) {