nbnxm_hip_kernel_utils.h File Reference

#include <assert.h>
#include <type_traits>
#include "gromacs/gpu_utils/hip_kernel_utils.h"
#include "gromacs/gpu_utils/typecasts_cuda_hip.h"
#include "gromacs/gpu_utils/vectype_ops_hip.h"
#include "gromacs/hardware/device_information.h"
#include "gromacs/nbnxm/nbnxm_enums.h"
#include "nbnxm_hip_types.h"

Include dependency graph for nbnxm_hip_kernel_utils.h:

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

Description

Utility constant and function declaration for the HIP non-bonded kernels. This header should be included once at the top level, just before the kernels are included (has to be preceded by nbnxn_hip_types.h).

Author

Szilárd Páll pall..nosp@m.szil.nosp@m.ard@g.nosp@m.mail.nosp@m..com

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

Functions
template<PairlistType pairlistType>
__device__ __forceinline__ int	gmx::nb_any_internal (int predicate, int widx)
template<PairlistType pairlistType, typename T >
__device__ size_t	gmx::incrementSharedMemorySlotPtr ()
	Increment the pointer into shared memory. More...
size_t	gmx::numberOfKernelBlocksSanityCheck (int numSci, const DeviceInformation &deviceInfo)
template<bool isPruneKernel, int numThreadZ, VdwType vdwType, PairlistType pairlistType>
constexpr size_t	gmx::requiredSharedMemorySize ()
bool	gmx::targetHasLargeRegisterPool (const DeviceInformation &deviceInfo)
	Find out if the target device has a large enough register pool (MI2xx and later)

Variables
template<PairlistType pairlistType>
__device__ constexpr int	gmx::c_subWarp = sc_gpuParallelExecutionWidth(pairlistType)
template<PairlistType pairlistType>
__device__ constexpr int	gmx::c_clSizeLog2 = StaticLog2<sc_gpuClusterSize(pairlistType)>::value
	Log of the i and j cluster size. change this together with c_clSize !
template<PairlistType pairlistType>
__device__ constexpr int	gmx::c_clSizeSq = sc_gpuClusterSize(pairlistType) * sc_gpuClusterSize(pairlistType)
	Square of cluster size.
template<PairlistType pairlistType>
static __device__ constexpr int	gmx::c_fbufStride = c_clSizeSq<pairlistType>