pbc_aiuc_hip.h File Reference

#include "gromacs/gpu_utils/gputraits.h"
#include "gromacs/gpu_utils/vectype_ops_hip.h"
#include "gromacs/pbcutil/ishift.h"
#include "gromacs/pbcutil/pbc_aiuc.h"

Include dependency graph for pbc_aiuc_hip.h:

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

Description

Basic routines to handle periodic boundary conditions with HIP.

This file contains GPU implementation of the PBC-aware vector evaluation.

Todo:

CPU, GPU and SIMD routines essentially do the same operations on different data-types. Currently this leads to code duplication, which has to be resolved. For details, see Issue #2863 https://gitlab.com/gromacs/gromacs/-/issues/2863

Author

Mark Abraham mark..nosp@m.j.ab.nosp@m.raham.nosp@m.@gma.nosp@m.il.co.nosp@m.m

Berk Hess hess@.nosp@m.kth..nosp@m.se

Artem Zhmurov zhmur.nosp@m.ov@g.nosp@m.mail..nosp@m.com

Functions
static __device__ int	xyzToShiftIndex (int x, int y, int z)
static __device__ int	int3ToShiftIndex (int3 iv)
template<bool returnShift>
static __forceinline__ __device__ int	pbcDxAiucGpu (const PbcAiuc &pbcAiuc, const float4 r1, const float4 r2, float3 &dr)
	Computes the vector between two points taking PBC into account. More...
static __forceinline__ __host__ __device__ float3	pbcDxAiuc (const PbcAiuc &pbcAiuc, const float3 &r1, const float3 &r2)
	Computes the vector between two points taking PBC into account. More...

Function Documentation

static __forceinline__ __host__ __device__ float3 pbcDxAiuc ( const PbcAiuc &  pbcAiuc, const float3 &  r1, const float3 &  r2  )

static

Computes the vector between two points taking PBC into account.

Computes the vector dr between points r2 and r1, taking into account the periodic boundary conditions, described in pbcAiuc object. Same as above, only takes and returns data in float3 format. Does not return shifts.

Todo:

This routine uses HIP float3 types for both input and returns values. Other than that, it does essentially the same thing as the version above, as well as SIMD and CPU versions. This routine is used in GPU-based constraints. To avoid code duplication, these implementations should be unified. See Issue #2863: https://gitlab.com/gromacs/gromacs/-/issues/2863

Parameters

[in]	pbcAiuc	PBC object.
[in]	r1	Coordinates of the first point.
[in]	r2	Coordinates of the second point.

Returns

dr Resulting distance.

template<bool returnShift>

static __forceinline__ __device__ int pbcDxAiucGpu ( const PbcAiuc &  pbcAiuc, const float4  r1, const float4  r2, float3 &  dr  )

static

Computes the vector between two points taking PBC into account.

Computes the vector dr between points r2 and r1, taking into account the periodic boundary conditions, described in pbcAiuc object. Note that this routine always does the PBC arithmetic for all directions, multiplying the displacements by zeroes if the corresponding direction is not periodic. For triclinic boxes only distances up to half the smallest box diagonal element are guaranteed to be the shortest. This means that distances from 0.5/sqrt(2) times a box vector length (e.g. for a rhombic dodecahedron) can use a more distant periodic image.

Todo:

This routine uses HIP float4 types for input coordinates and returns in float3 data-type. Other than that, it does essentially the same thing as the version below, as well as SIMD and CPU versions. This routine is used in GPU listed forces module. To avoid code duplication, these implementations should be unified. See Issue #2863: https://gitlab.com/gromacs/gromacs/-/issues/2863

Parameters

[in]	pbcAiuc	PBC object.
[in]	r1	Coordinates of the first point.
[in]	r2	Coordinates of the second point.
[out]	dr	Resulting distance.