listed_forces_gpu_internal_hip.cpp File Reference

#include "gmxpre.h"
#include "hip/hip_math_constants.h"
#include "gromacs/gpu_utils/hiputils.h"
#include "gromacs/gpu_utils/typecasts_cuda_hip.h"
#include "gromacs/gpu_utils/vectype_ops_hip.h"
#include "gromacs/listed_forces/listed_forces_gpu.h"
#include "gromacs/listed_forces/listed_forces_gpu_internal_shared.h"
#include "gromacs/math/units.h"
#include "gromacs/mdlib/force_flags.h"
#include "gromacs/mdtypes/interaction_const.h"
#include "gromacs/mdtypes/simulation_workload.h"
#include "gromacs/pbcutil/pbc_aiuc_hip.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_hip.cpp:

Description

Implements HIP bonded functionality.

Author

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

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

Functions
template<bool calcVir, bool calcEner>
	gmx::__launch_bounds__ (c_threadsBondedPerBlock) __global__ void bonded_kernel_gpu(BondedGpuKernelParameters kernelParams
if	gmx::constexpr (calcVir)
	gmx::for (int j=0;j< numFTypesOnGpu;j++)
	gmx::if (threadComputedPotential)

Variables
BondedGpuKernelBuffers	gmx::kernelBuffers
BondedGpuKernelBuffers float4 *	gmx::gm_xq_in
BondedGpuKernelBuffers float4 float3 *	gmx::gm_f
BondedGpuKernelBuffers float4 float3 float3 *	gmx::gm_fShift
const int	gmx::tid = blockIdx.x * blockDim.x + threadIdx.x
float	gmx::vtot_loc = 0.0F
float	gmx::vtotElec_loc = 0.0F
__shared__ float3	gmx::sm_dynamicShmem []
float3 *	gmx::sm_fShiftLoc = sm_dynamicShmem
DeviceFloat4 *	gmx::gm_xq = reinterpret_cast<DeviceFloat4*>(gm_xq_in)
InteractionFunction	gmx::fType
bool	gmx::threadComputedPotential = false