listed_forces_gpu_internal_sycl.cpp File Reference

#include "gmxpre.h"
#include "gromacs/gpu_utils/devicebuffer_sycl.h"
#include "gromacs/gpu_utils/gmxsycl.h"
#include "gromacs/gpu_utils/sycl_kernel_utils.h"
#include "gromacs/listed_forces/listed_forces_gpu.h"
#include "gromacs/math/units.h"
#include "gromacs/mdtypes/interaction_const.h"
#include "gromacs/mdtypes/simulation_workload.h"
#include "gromacs/pbcutil/pbc_aiuc_sycl.h"
#include "gromacs/timing/wallcycle.h"
#include "gromacs/utility/gmxassert.h"
#include "listed_forces_gpu_impl.h"

Include dependency graph for listed_forces_gpu_internal_sycl.cpp:

Description

Implements SYCL bonded functionality.

Author

Andrey Alekseenko al42a.nosp@m.nd@g.nosp@m.mail..nosp@m.com

Jon Vincent jvinc.nosp@m.ent@.nosp@m.nvidi.nosp@m.a.co.nosp@m.m

Magnus Lundborg lundb.nosp@m.org..nosp@m.magnu.nosp@m.s@gm.nosp@m.ail.c.nosp@m.om

Berk Hess hess@.nosp@m.kth..nosp@m.se

Szilárd Páll pall..nosp@m.szil.nosp@m.ard@g.nosp@m.mail.nosp@m..com

Alan Gray alang.nosp@m.@nvi.nosp@m.dia.c.nosp@m.om

Mark Abraham mark..nosp@m.j.ab.nosp@m.raham.nosp@m.@gma.nosp@m.il.co.nosp@m.m

Typedefs
using	gmx::mode = sycl::access_mode

Functions
static void	staggeredAtomicAddForce (sycl::global_ptr< Float3 > gm_f, Float3 f, const int localId)
template<bool calcEner>
static void	harmonic_gpu (const float kA, const float xA, const float x, sycl::private_ptr< float > V, sycl::private_ptr< float > F)
template<bool calcVir, bool calcEner>
static void	bonds_gpu (const int i, sycl::private_ptr< float > vtot_loc, const sycl::global_ptr< const t_iatom > gm_forceatoms, const sycl::global_ptr< const t_iparams > gm_forceparams, const sycl::global_ptr< const sycl::float4 > gm_xq, sycl::global_ptr< Float3 > gm_f, sycl::local_ptr< Float3 > sm_fShiftLoc, const PbcAiuc &pbcAiuc, const int localId)
static float	cos_angle (const Float3 &a, const Float3 &b)
	Cosine of an angle between two vectors. More...
static float	angle (const Float3 &a, const Float3 &b)
	Compute the angle between two vectors. More...
template<bool returnShift>
static float	bond_angle_gpu (const sycl::float4 xi, const sycl::float4 xj, const sycl::float4 xk, const PbcAiuc &pbcAiuc, sycl::private_ptr< Float3 > r_ij, sycl::private_ptr< Float3 > r_kj, sycl::private_ptr< float > costh, sycl::private_ptr< int > t1, sycl::private_ptr< int > t2)
template<bool calcVir, bool calcEner>
static void	angles_gpu (const int i, sycl::private_ptr< float > vtot_loc, const sycl::global_ptr< const t_iatom > gm_forceatoms, const sycl::global_ptr< const t_iparams > gm_forceparams, const sycl::global_ptr< const sycl::float4 > gm_xq, sycl::global_ptr< Float3 > gm_f, sycl::local_ptr< Float3 > sm_fShiftLoc, const PbcAiuc &pbcAiuc, const int localId)
template<bool calcVir, bool calcEner>
static void	urey_bradley_gpu (const int i, sycl::private_ptr< float > vtot_loc, const sycl::global_ptr< const t_iatom > gm_forceatoms, const sycl::global_ptr< const t_iparams > gm_forceparams, const sycl::global_ptr< const sycl::float4 > gm_xq, sycl::global_ptr< Float3 > gm_f, sycl::local_ptr< Float3 > sm_fShiftLoc, const PbcAiuc &pbcAiuc, const int localId)
template<bool returnShift, typename T >
static float	dih_angle_gpu (const T xi, const T xj, const T xk, const T xl, const PbcAiuc &pbcAiuc, sycl::private_ptr< Float3 > r_ij, sycl::private_ptr< Float3 > r_kj, sycl::private_ptr< Float3 > r_kl, sycl::private_ptr< Float3 > m, sycl::private_ptr< Float3 > n, sycl::private_ptr< int > t1, sycl::private_ptr< int > t2, sycl::private_ptr< int > t3)
template<bool returnShift, typename T >
static float	dih_angle_gpu_sincos (const T xi, const T xj, const T xk, const T xl, const PbcAiuc &pbcAiuc, sycl::private_ptr< Float3 > r_ij, sycl::private_ptr< Float3 > r_kj, sycl::private_ptr< Float3 > r_kl, sycl::private_ptr< Float3 > m, sycl::private_ptr< Float3 > n, sycl::private_ptr< int > t1, sycl::private_ptr< int > t2, sycl::private_ptr< float > cosval)
static void	dopdihs_gpu (const float cpA, const float phiA, const int mult, const float phi, sycl::private_ptr< float > v, sycl::private_ptr< float > f)
template<bool calcVir>
static void	do_dih_fup_gpu (const int i, const int j, const int k, const int l, const float ddphi, const Float3 r_ij, const Float3 r_kj, const Float3 r_kl, const Float3 m, const Float3 n, sycl::global_ptr< Float3 > gm_f, sycl::local_ptr< Float3 > sm_fShiftLoc, const PbcAiuc &pbcAiuc, const sycl::global_ptr< const Float4 > gm_xq, const int t1, const int t2, const int localId)
template<bool calcVir, bool calcEner>
static void	pdihs_gpu (const int i, sycl::private_ptr< float > vtot_loc, const sycl::global_ptr< const t_iatom > gm_forceatoms, const sycl::global_ptr< const t_iparams > gm_forceparams, const sycl::global_ptr< const sycl::float4 > gm_xq, sycl::global_ptr< Float3 > gm_f, sycl::local_ptr< Float3 > sm_fShiftLoc, const PbcAiuc &pbcAiuc, const int localId)
template<bool calcVir, bool calcEner>
static void	rbdihs_gpu (const int i, sycl::private_ptr< float > vtot_loc, const sycl::global_ptr< const t_iatom > gm_forceatoms, const sycl::global_ptr< const t_iparams > gm_forceparams, const sycl::global_ptr< const sycl::float4 > gm_xq, sycl::global_ptr< Float3 > gm_f, sycl::local_ptr< Float3 > sm_fShiftLoc, const PbcAiuc &pbcAiuc, const int localId)
static constexpr float	wrapAngle (float a)
	Wrap angle from range [-3*pi; 3*pi) to [-pi; pi)
template<bool calcVir, bool calcEner>
static void	idihs_gpu (const int i, sycl::private_ptr< float > vtot_loc, const sycl::global_ptr< const t_iatom > gm_forceatoms, const sycl::global_ptr< const t_iparams > gm_forceparams, const sycl::global_ptr< const sycl::float4 > gm_xq, sycl::global_ptr< Float3 > gm_f, sycl::local_ptr< Float3 > sm_fShiftLoc, const PbcAiuc &pbcAiuc, const int localId)
template<bool calcVir, bool calcEner>
static void	pairs_gpu (const int i, const sycl::global_ptr< const t_iatom > gm_forceatoms, const sycl::global_ptr< const t_iparams > gm_iparams, const sycl::global_ptr< const sycl::float4 > gm_xq, sycl::global_ptr< Float3 > gm_f, sycl::local_ptr< Float3 > sm_fShiftLoc, const PbcAiuc &pbcAiuc, const float scale_factor, sycl::private_ptr< float > vtotVdw_loc, sycl::private_ptr< float > vtotElec_loc, const int localId)
template<bool calcVir, bool calcEner>
auto	gmx::bondedKernel (sycl::handler &cgh, const BondedGpuKernelParameters &kernelParams, const DeviceBuffer< t_iatom > gm_iatoms_[numFTypesOnGpu], float *__restrict__ gm_vTot, const t_iparams *__restrict__ gm_forceParams_, const sycl::float4 *__restrict__ gm_xq_, Float3 *__restrict__ gm_f_, Float3 *__restrict__ gm_fShift_)

Variables
static constexpr float	c_deg2RadF = gmx::c_deg2Rad
constexpr float	c_Pi = M_PI

Function Documentation

static float angle ( const Float3 &  a, const Float3 &  b  )

inlinestatic

Compute the angle between two vectors.

Uses atan( |axb| / a.b ) formula.

Parameters

[in]	a	First vector.
[in]	b	Second vector.

Returns

Angle between vectors in radians.

static float cos_angle ( const Float3 &  a, const Float3 &  b  )

inlinestatic

Cosine of an angle between two vectors.

Computes cosine using the following formula:

             ax*bx + ay*by + az*bz

cos-vec (a,b) = ------------------— ||a|| * ||b||

This function also makes sure that the cosine does not leave the [-1, 1] interval, which can happen due to numerical errors.

Parameters

[in]	a	First vector.
[in]	b	Second vector.

Returns

Cosine between a and b.