nbnxm_sycl_kernel_body.h File Reference

#include "gromacs/gpu_utils/devicebuffer.h"
#include "gromacs/gpu_utils/gmxsycl.h"
#include "gromacs/math/functions.h"
#include "gromacs/mdtypes/simulation_workload.h"
#include "gromacs/pbcutil/ishift.h"
#include "gromacs/utility/template_mp.h"
#include "nbnxm_sycl_kernel.h"
#include "nbnxm_sycl_kernel_utils.h"
#include "nbnxm_sycl_types.h"

Include dependency graph for nbnxm_sycl_kernel_body.h:

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

Description

NBNXM SYCL kernels.

Classes
class	NbnxmKernel< doPruneNBL, doCalcEnergies, elecType, vdwType, subGroupSize >
	Class name for NBNXM kernel. More...
struct	Nbnxm::EnergyFunctionProperties< elecType, vdwType >
	Set of boolean constants mimicking preprocessor macros. More...

Macros
#define	GMX_NBNXM_ENABLE_PACKED_FLOAT3   1
	Macro to control the enablement of manually-packed Float3 structure. More...

Typedefs
using	FCiFloat3 = Float3
using	Nbnxm::mode = sycl::access_mode

Functions
static Float2	Nbnxm::convertSigmaEpsilonToC6C12 (const float sigma, const float epsilon)
	Convert sigma and epsilon VdW parameters to c6,c12 pair.
template<bool doCalcEnergies>
static void	Nbnxm::ljForceSwitch (const shift_consts_t dispersionShift, const shift_consts_t repulsionShift, const float rVdwSwitch, const float c6, const float c12, const float rInv, const float r2, sycl::private_ptr< float > fInvR, sycl::private_ptr< float > eLJ)
	Calculate force and energy for a pair of atoms, VdW force-switch flavor.
template<enum VdwType vdwType>
static float	Nbnxm::calculateLJEwaldC6Grid (const sycl::global_ptr< const Float2 > a_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 void	Nbnxm::ljEwaldComb (const sycl::global_ptr< const Float2 > a_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 int_bit, sycl::private_ptr< float > fInvR, sycl::private_ptr< float > eLJ)
	Calculate LJ-PME grid force contribution with geometric or LB combination rule.
template<bool doCalcEnergies>
static void	Nbnxm::ljPotentialSwitch (const switch_consts_t vdwSwitch, const float rVdwSwitch, const float rInv, const float r2, sycl::private_ptr< float > fInvR, sycl::private_ptr< float > eLJ)
	Apply potential switch.
static float	Nbnxm::pmeCorrF (const float z2)
	Calculate analytical Ewald correction term.
template<typename T >
static T	Nbnxm::lerp (T d0, T d1, T t)
	Linear interpolation using exactly two FMA operations. More...
static float	Nbnxm::interpolateCoulombForceR (const sycl::global_ptr< const float > a_coulombTab, const float coulombTabScale, const float r)
	Interpolate Ewald coulomb force correction using the F*r table.
static void	Nbnxm::reduceForceJShuffle (Float3 f, const sycl::nd_item< 3 > itemIdx, const int tidxi, const int aidx, sycl::global_ptr< Float3 > a_f)
	Reduce c_clSize j-force components using shifts and atomically accumulate into a_f. More...
template<int subGroupSize, int groupSize>
static float	Nbnxm::groupReduce (const sycl::nd_item< 3 > itemIdx, const unsigned int tidxi, sycl::local_ptr< float > sm_buf, float valueToReduce)
	Do workgroup-level reduction of a single float. More...
static void	Nbnxm::reduceForceJGeneric (sycl::local_ptr< float > sm_buf, Float3 f, const sycl::nd_item< 3 > itemIdx, const int tidxi, const int tidxj, const int aidx, sycl::global_ptr< Float3 > a_f)
	Reduce c_clSize j-force components using local memory and atomically accumulate into a_f. More...
template<bool useShuffleReduction>
static void	Nbnxm::reduceForceJ (sycl::local_ptr< float > sm_buf, Float3 f, const sycl::nd_item< 3 > itemIdx, const int tidxi, const int tidxj, const int aidx, sycl::global_ptr< Float3 > a_f)
	Reduce c_clSize j-force components using either shifts or local memory and atomically accumulate into a_f.
template<typename FCiBufferWrapperX , typename FCiBufferWrapperY , typename FCiBufferWrapperZ >
static void	Nbnxm::reduceForceIAndFShiftGeneric (sycl::local_ptr< float > sm_buf, const FCiBufferWrapperX &fCiBufX, const FCiBufferWrapperY &fCiBufY, const FCiBufferWrapperZ &fCiBufZ, const bool calcFShift, const sycl::nd_item< 3 > itemIdx, const int tidxi, const int tidxj, const int sci, const int shift, sycl::global_ptr< Float3 > a_f, sycl::global_ptr< Float3 > a_fShift)
	Local memory-based i-force reduction. More...
template<int numShuffleReductionSteps, typename FCiBufferWrapperX , typename FCiBufferWrapperY , typename FCiBufferWrapperZ >
static std::enable_if_t < numShuffleReductionSteps!=1, void >	Nbnxm::reduceForceIAndFShiftShuffles (const FCiBufferWrapperX &fCiBufX, const FCiBufferWrapperY &fCiBufY, const FCiBufferWrapperZ &fCiBufZ, const bool calcFShift, const sycl::nd_item< 3 > itemIdx, const int tidxi, const int tidxj, const int sci, const int shift, sycl::global_ptr< Float3 > a_f, sycl::global_ptr< Float3 > a_fShift)
	Shuffle-based i-force reduction. More...
template<int numShuffleReductionSteps, typename FCiBufferWrapperX , typename FCiBufferWrapperY , typename FCiBufferWrapperZ >
static std::enable_if_t < numShuffleReductionSteps==1, void >	Nbnxm::reduceForceIAndFShiftShuffles (const FCiBufferWrapperX &fCiBufX, const FCiBufferWrapperY &fCiBufY, const FCiBufferWrapperZ &fCiBufZ, const bool calcFShift, const sycl::nd_item< 3 > itemIdx, const int tidxi, const int tidxj, const int sci, const int shift, sycl::global_ptr< Float3 > a_f, sycl::global_ptr< Float3 > a_fShift)
	reduceForceIAndFShiftShuffles specialization for single-step reduction (e.g., Intel iGPUs). More...
template<bool useShuffleReduction, int subGroupSize, typename FCiBufferWrapperX , typename FCiBufferWrapperY , typename FCiBufferWrapperZ >
static void	Nbnxm::reduceForceIAndFShift (sycl::local_ptr< float > sm_buf, const FCiBufferWrapperX &fCiBufX, const FCiBufferWrapperY &fCiBufY, const FCiBufferWrapperZ &fCiBufZ, const bool calcFShift, const sycl::nd_item< 3 > itemIdx, const int tidxi, const int tidxj, const int sci, const int shift, sycl::global_ptr< Float3 > a_f, sycl::global_ptr< Float3 > a_fShift)
	Final i-force reduction. More...
template<int subGroupSize, bool doPruneNBL, bool doCalcEnergies, enum ElecType elecType, enum VdwType vdwType>
static auto	Nbnxm::nbnxmKernel (sycl::handler &cgh, DeviceAccessor< Float4, mode::read > a_xq, DeviceAccessor< Float3, mode::read_write > a_f, DeviceAccessor< Float3, mode::read > a_shiftVec, DeviceAccessor< Float3, mode::read_write > a_fShift, OptionalAccessor< float, mode::read_write, doCalcEnergies > a_energyElec, OptionalAccessor< float, mode::read_write, doCalcEnergies > a_energyVdw, DeviceAccessor< nbnxn_cj_packed_t, doPruneNBL?mode::read_write:mode::read > a_plistCJPacked, DeviceAccessor< nbnxn_sci_t, mode::read > a_plistSci, DeviceAccessor< nbnxn_excl_t, mode::read > a_plistExcl, OptionalAccessor< Float2, mode::read, ljComb< vdwType >> a_ljComb, OptionalAccessor< int, mode::read,!ljComb< vdwType >> a_atomTypes, OptionalAccessor< Float2, mode::read,!ljComb< vdwType >> a_nbfp, OptionalAccessor< Float2, mode::read, ljEwald< vdwType >> a_nbfpComb, OptionalAccessor< float, mode::read, elecEwaldTab< elecType >> a_coulombTab, const int numTypes, const float rCoulombSq, const float rVdwSq, const float twoKRf, const float ewaldBeta, const float rlistOuterSq, const float ewaldShift, const float epsFac, const float ewaldCoeffLJ_2, const float cRF, const shift_consts_t dispersionShift, const shift_consts_t repulsionShift, const switch_consts_t vdwSwitch, const float rVdwSwitch, const float ljEwaldShift, const float coulombTabScale, const bool calcShift)
	Main kernel for NBNXM. More...
template<int subGroupSize, bool doPruneNBL, bool doCalcEnergies, enum ElecType elecType, enum VdwType vdwType, class... Args>
static void	Nbnxm::launchNbnxmKernel (const DeviceStream &deviceStream, const int numSci, Args &&...args)
	NBNXM kernel launch code.
template<int subGroupSize, bool doPruneNBL, bool doCalcEnergies, class... Args>
void	Nbnxm::chooseAndLaunchNbnxmKernel (enum ElecType elecType, enum VdwType vdwType, Args &&...args)
	Select templated kernel and launch it.
template<int subGroupSize, bool doPruneNBL, bool doCalcEnergies>
void	Nbnxm::launchNbnxmKernelHelper (NbnxmGpu *nb, const gmx::StepWorkload &stepWork, const InteractionLocality iloc)

Variables
constexpr bool	Nbnxm::c_avoidFloatingPointAtomics = (c_clSize == 4)
	Should we avoid FP atomics to the same location from the same work-group? More...
template<enum VdwType vdwType>
constexpr bool	Nbnxm::ljComb = EnergyFunctionProperties<ElecType::Count, vdwType>().vdwComb
	Templated constants to shorten kernel function declaration.
template<enum ElecType elecType>
constexpr bool	Nbnxm::elecEwald = EnergyFunctionProperties<elecType, VdwType::Count>().elecEwald
template<enum ElecType elecType>
constexpr bool	Nbnxm::elecEwaldTab = EnergyFunctionProperties<elecType, VdwType::Count>().elecEwaldTab
template<enum VdwType vdwType>
constexpr bool	Nbnxm::ljEwald = EnergyFunctionProperties<ElecType::Count, vdwType>().vdwEwald

Macro Definition Documentation

#define GMX_NBNXM_ENABLE_PACKED_FLOAT3   1

Macro to control the enablement of manually-packed Float3 structure.

If enabled (default), the explicit packed math will be used on devices where it is known to be beneficial (currently, AMD MI250X / gfx90a). If disabled, packed math will never be explicitly used.

This only controls the use of AmdPackedFloat3 datastructure, not the layout of fCi buffer.

See issue #4854