lincs_gpu_internal_sycl.cpp File Reference

#include "gmxpre.h"
#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"

Include dependency graph for lincs_gpu_internal_sycl.cpp:

Description

Implements LINCS kernels using SYCL.

This file contains SYCL kernels of LINCS constraints algorithm.

Author

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

Functions
template<bool updateVelocities, bool computeVirial, bool haveCoupledConstraints>
auto	gmx::lincsKernel (sycl::handler &cgh, const int numConstraintsThreads, const AtomPair *__restrict__ gm_constraints, const float *__restrict__ gm_constraintsTargetLengths, const int *__restrict__ gm_coupledConstraintsCounts, const int *__restrict__ gm_coupledConstraintsIndices, const float *__restrict__ gm_massFactors, float *__restrict__ gm_matrixA, const float *__restrict__ gm_inverseMasses, const int numIterations, const int expansionOrder, const Float3 *__restrict__ gm_x, Float3 *__restrict__ gm_xp, const float invdt, Float3 *__restrict__ gm_v, float *__restrict__ gm_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...