nbnxm_hip_kernel_body.h File Reference

#include <memory>
#include <type_traits>
#include <utility>
#include "gromacs/gpu_utils/devicebuffer.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/nbnxm/nbnxm_kernel_utils.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_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

Typedefs
template<typename T , bool Condition>
using	gmx::internal::MaybeConstT = std::conditional_t< Condition, const T, T >

Functions
template<ElecType elecType, VdwType vdwType, bool doPrune, bool doCalcEnergies, PairlistType pairlistType>
static const std::string	gmx::internal::getKernelName ()
	Lookup kernel name based on launch configuration.
template<bool doPruneNBL, bool doCalcEnergies, enum ElecType elecType, enum VdwType vdwType, int nthreadZ, int minBlocksPerMp, PairlistType pairlistType>
	gmx::internal::__launch_bounds__ (c_clSizeSq< pairlistType > *nthreadZ, minBlocksPerMp) __global__ static void nbnxmKernel(NBAtomDataGpu atdat
	Main kernel for NBNXM. More...
template<PairlistType pairlistType, bool doPruneNBL, bool doCalcEnergies, class... Args>
void	gmx::chooseAndLaunchNbnxmKernel (ElecType elecType, VdwType vdwType, const DeviceStream &deviceStream, const int numSci, const DeviceInformation &deviceInfo, Args *...args)
	Select templated kernel and launch it.
template<bool doPruneNBL, bool doCalcEnergies>
void	gmx::launchNbnxmKernel (NbnxmGpu *nb, const StepWorkload &stepWork, InteractionLocality iloc)

Variables
NBParamGpu	gmx::internal::nbparam
NBParamGpu GpuPairlist	gmx::internal::plist
NBParamGpu GpuPairlist MaybeConstT< nbnxm_cj_packed_t,!doPruneNBL > *__restrict__	gmx::internal::gm_plistCJPacked
NBParamGpu GpuPairlist MaybeConstT< nbnxm_cj_packed_t,!doPruneNBL > *__restrict__ const nbnxm_sci_t *__restrict__	gmx::internal::gm_sci