From e79b665d5c3fee78103813c7304996b29cc62800 Mon Sep 17 00:00:00 2001 From: Andrey Alekseenko Date: Wed, 28 Apr 2021 00:19:15 +0300 Subject: [PATCH] Add SYCL implementation of X/V scale kernel Refs #3929 --- .../update_constrain_gpu_internal_sycl.cpp | 38 ++++++++++++++++--- 1 file changed, 32 insertions(+), 6 deletions(-) diff --git a/src/gromacs/mdlib/update_constrain_gpu_internal_sycl.cpp b/src/gromacs/mdlib/update_constrain_gpu_internal_sycl.cpp index 72723bd99e..cf9439bf17 100644 --- a/src/gromacs/mdlib/update_constrain_gpu_internal_sycl.cpp +++ b/src/gromacs/mdlib/update_constrain_gpu_internal_sycl.cpp @@ -44,18 +44,44 @@ #include "update_constrain_gpu_internal.h" +#include "gromacs/gpu_utils/devicebuffer_sycl.h" +#include "gromacs/gpu_utils/gmxsycl.h" +#include "gromacs/gpu_utils/gputraits_sycl.h" #include "gromacs/utility/gmxassert.h" namespace gmx { -void launchScaleCoordinatesKernel(const int /* numAtoms */, - DeviceBuffer /* d_coordinates */, - const ScalingMatrix /* mu */, - const DeviceStream& /* deviceStream */) +static auto scaleKernel(cl::sycl::handler& cgh, + DeviceAccessor a_x, + const ScalingMatrix scalingMatrix) { - // SYCL_TODO - GMX_RELEASE_ASSERT(false, "Coordinates scaling is not yet implemented in SYCL."); + cgh.require(a_x); + + return [=](cl::sycl::id<1> itemIdx) { + Float3 x = a_x[itemIdx]; + x[0] = scalingMatrix.xx * x[0] + scalingMatrix.yx * x[1] + scalingMatrix.zx * x[2]; + x[1] = scalingMatrix.yy * x[1] + scalingMatrix.zy * x[2]; + x[2] = scalingMatrix.zz * x[2]; + a_x[itemIdx] = x; + }; +} + +void launchScaleCoordinatesKernel(const int numAtoms, + DeviceBuffer d_coordinates, + const ScalingMatrix mu, + const DeviceStream& deviceStream) +{ + const cl::sycl::range<1> rangeAllAtoms(numAtoms); + cl::sycl::queue queue = deviceStream.stream(); + + cl::sycl::event e = queue.submit([&](cl::sycl::handler& cgh) { + auto kernel = scaleKernel(cgh, d_coordinates, mu); + cgh.parallel_for(rangeAllAtoms, kernel); + }); + // TODO: Although this only happens on the pressure coupling steps, this synchronization + // can affect the performance if nstpcouple is small. See Issue #4018 + e.wait_and_throw(); } } // namespace gmx -- 2.22.0