Gromacs  2024.2
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
Typedefs | Functions | Variables
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]aFirst vector.
[in]bSecond 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]aFirst vector.
[in]bSecond vector.
Returns
Cosine between a and b.