Gromacs
2023-beta
|
#include "gromacs/gpu_utils/devicebuffer.h"
#include "gromacs/gpu_utils/gmxsycl.h"
#include "gromacs/gpu_utils/sycl_kernel_utils.h"
#include "gromacs/mdlib/lincs_gpu.h"
#include "gromacs/pbcutil/pbc_aiuc_sycl.h"
#include "gromacs/utility/gmxassert.h"
#include "gromacs/utility/template_mp.h"
#include "lincs_gpu_internal.h"
Implements LINCS kernels using SYCL.
This file contains SYCL kernels of LINCS constraints algorithm.
Functions | |
template<bool updateVelocities, bool computeVirial, bool haveCoupledConstraints> | |
auto | gmx::lincsKernel (sycl::handler &cgh, const int numConstraintsThreads, DeviceAccessor< AtomPair, mode::read > a_constraints, DeviceAccessor< float, mode::read > a_constraintsTargetLengths, OptionalAccessor< int, mode::read, haveCoupledConstraints > a_coupledConstraintsCounts, OptionalAccessor< int, mode::read, haveCoupledConstraints > a_coupledConstraintsIndices, OptionalAccessor< float, mode::read, haveCoupledConstraints > a_massFactors, OptionalAccessor< float, mode::read_write, haveCoupledConstraints > a_matrixA, DeviceAccessor< float, mode::read > a_inverseMasses, const int numIterations, const int expansionOrder, DeviceAccessor< Float3, mode::read > a_x, DeviceAccessor< Float3, mode::read_write > a_xp, const float invdt, OptionalAccessor< Float3, mode::read_write, updateVelocities > a_v, OptionalAccessor< float, mode::read_write, computeVirial > a_virialScaled, PbcAiuc pbcAiuc) |
Main kernel for LINCS constraints. More... | |
template<bool updateVelocities, bool computeVirial, bool haveCoupledConstraints, class... Args> | |
static void | gmx::launchLincsKernel (const DeviceStream &deviceStream, const int numConstraintsThreads, Args &&...args) |
template<class... Args> | |
static void | gmx::launchLincsKernel (bool updateVelocities, bool computeVirial, bool haveCoupledConstraints, Args &&...args) |
Select templated kernel and launch it. | |
void | gmx::launchLincsGpuKernel (LincsGpuKernelParameters *kernelParams, const DeviceBuffer< Float3 > &d_x, DeviceBuffer< Float3 > d_xp, bool updateVelocities, DeviceBuffer< Float3 > d_v, real invdt, bool computeVirial, const DeviceStream &deviceStream) |
Backend-specific function to launch LINCS kernel. More... | |