nbnxm_hip_kernel_body.h File Reference

#include <memory>
#include <type_traits>
#include <utility>
#include "gromacs/gpu_utils/devicebuffer.h"
#include "gromacs/gpu_utils/gpu_kernel_utils.h"
#include "gromacs/gpu_utils/typecasts_cuda_hip.h"
#include "gromacs/hardware/device_information.h"
#include "gromacs/math/functions.h"
#include "gromacs/mdtypes/simulation_workload.h"
#include "gromacs/nbnxm/gpu_types_common.h"
#include "gromacs/nbnxm/nbnxm_enums.h"
#include "gromacs/pbcutil/ishift.h"
#include "gromacs/utility/arrayref.h"
#include "gromacs/utility/enumerationhelpers.h"
#include "gromacs/utility/stringutil.h"
#include "gromacs/utility/template_mp.h"
#include "nbnxm_hip_kernel.h"
#include "nbnxm_hip_kernel_utils.h"
#include "nbnxm_hip_types.h"

Include dependency graph for nbnxm_hip_kernel_body.h:

This graph shows which files directly or indirectly include this file:

Description

NBNXM HIP kernels.

Author

Paul Bauer paul..nosp@m.baue.nosp@m.r.q@g.nosp@m.mail.nosp@m..com

Functions
template<ElecType elecType, VdwType vdwType, bool doPrune, bool doCalcEnergies, PairlistType pairlistType>
static const std::string	gmx::getKernelName ()
	Lookup kernel name based on launch configuration.
static __device__ float2	gmx::fetchNbfpC6C12 (const float2 *nbfpComb, int type)
static __device__ float2	gmx::convertSigmaEpsilonToC6C12 (const float sigma, const float epsilon)
	Convert sigma and epsilon VdW parameters to c6,c12 pair.
template<bool doCalcEnergies>
static __device__ void	gmx::ljForceSwitch (const shift_consts_t dispersionShift, const shift_consts_t repulsionShift, const float2 c6c12, const float rVdwSwitch, const float rInv, const float r2, float *fInvR, float *eLJ)
	Calculate force and energy for a pair of atoms, VdW force-switch flavor.
template<enum VdwType vdwType>
static __device__ float	gmx::calculateLJEwaldC6Grid (const Float2 *nbfpComb, const int typeI, const int typeJ)
	Fetch C6 grid contribution coefficients and return the product of these.
template<bool doCalcEnergies, enum VdwType vdwType>
static __device__ void	gmx::ljEwaldComb (const Float2 *nbfpComb, const float sh_lj_ewald, const int typeI, const int typeJ, const float r2, const float r2Inv, const float lje_coeff2, const float lje_coeff6_6, const float pairExclMask, float *fInvR, float *eLJ)
	Calculate LJ-PME grid force contribution with geometric or LB combination rule.
template<bool doCalcEnergies>
static __device__ void	gmx::ljPotentialSwitch (const switch_consts_t vdwSwitch, const float rVdwSwitch, const float rInv, const float r2, float *fInvR, float *eLJ)
	Apply potential switch.
static __device__ float	gmx::pmeCorrF (const float z2)
	Calculate analytical Ewald correction term.
static __device__ float2	gmx::fetchCoulombForceR (const float *coulombTable, int index)
static __device__ float	gmx::interpolateCoulombForceR (const float *coulombTable, const float coulombTabScale, const float r)
	Interpolate Ewald coulomb force correction using the F*r table.
template<PairlistType pairlistType>
static __device__ float	gmx::reduceForceJWarpShuffle (AmdPackedFloat3 f, const int tidxi)
	Reduce c_clSize j-force components using AMD DPP instruction. More...
template<PairlistType pairlistType>
static __device__ float	gmx::reduceForceIWarpShuffle (AmdPackedFloat3 f, const int tidxi, const int tidxj)
	Lowest level i force reduction. More...
template<PairlistType pairlistType, bool calculateShift>
static __device__ float3	gmx::reduceForceIAtomics (AmdPackedFloat3 input, float3 *result, const int tidxj, const int aidx)
	Lowest level i force reduction. More...
template<bool calculateShift, PairlistType pairlistType>
static __device__ void	gmx::reduceForceI (AmdPackedFloat3 *input, float3 *result, const int tidxi, const int tidxj, const int tidx, const int sci, float3 *fShift, const int shiftBase)
	Reduce i forces. More...
template<PairlistType pairlistType>
static __device__ void	gmx::reduceEnergyWarpShuffle (float localLJ, float localEl, float *gm_LJ, float *gm_El, int tidx)
	Energy reduction kernel. More...
template<bool doPruneNBL, bool doCalcEnergies, enum ElecType elecType, enum VdwType vdwType, int nthreadZ, int minBlocksPerMp, PairlistType pairlistType>
	gmx::__launch_bounds__ (c_clSizeSq< pairlistType > *nthreadZ, minBlocksPerMp) __global__ static void nbnxmKernel(NBAtomDataGpu atdat
	Main kernel for NBNXM. More...

Variables
NBParamGpu	gmx::nbparam
NBParamGpu GpuPairlist	gmx::plist