settle_gpu_internal_sycl.cpp File Reference

#include "gmxpre.h"
#include "gromacs/gpu_utils/devicebuffer.h"
#include "gromacs/gpu_utils/sycl_kernel_utils.h"
#include "gromacs/pbcutil/pbc_aiuc_sycl.h"
#include "gromacs/utility/gmxassert.h"
#include "gromacs/utility/template_mp.h"
#include "settle_gpu_internal.h"

Include dependency graph for settle_gpu_internal_sycl.cpp:

Description

SYCL-specific routines for the GPU implementation of SETTLE constraints algorithm.

Author

Artem Zhmurov zhmur.nosp@m.ov@g.nosp@m.mail..nosp@m.com

Functions
template<bool updateVelocities, bool computeVirial>
auto	gmx::settleKernel (sycl::handler &cgh, const int numSettles, const WaterMolecule *__restrict__ gm_settles, SettleParameters pars, const Float3 *__restrict__ gm_x, Float3 *__restrict__ gm_xp, float invdt, Float3 *__restrict__ gm_v, float *__restrict__ gm_virialScaled, PbcAiuc pbcAiuc)
	Function returning the SETTLE kernel lambda.
template<bool updateVelocities, bool computeVirial, class... Args>
static void	gmx::launchSettleKernel (const DeviceStream &deviceStream, int numSettles, Args &&...args)
	SETTLE SYCL kernel launch code.
template<class... Args>
static void	gmx::launchSettleKernel (bool updateVelocities, bool computeVirial, Args &&...args)
	Select templated kernel and launch it.
void	gmx::launchSettleGpuKernel (int numSettles, const DeviceBuffer< WaterMolecule > &d_atomIds, const SettleParameters &settleParameters, const DeviceBuffer< Float3 > &d_x, DeviceBuffer< Float3 > d_xp, bool updateVelocities, DeviceBuffer< Float3 > d_v, real invdt, bool computeVirial, DeviceBuffer< float > d_virialScaled, const PbcAiuc &pbcAiuc, const DeviceStream &deviceStream)
	Apply SETTLE. More...

Variables
static constexpr int	gmx::sc_workGroupSize = 256
	Number of work-items in a work-group.