pme_gpu_grid_hip.cpp File Reference

#include "gmxpre.h"
#include "config.h"
#include <cstdlib>
#include "gromacs/fft/parallel_3dfft.h"
#include "gromacs/gpu_utils/devicebuffer_hip.h"
#include "gromacs/gpu_utils/hiputils.h"
#include "gromacs/hardware/device_information.h"
#include "gromacs/math/functions.h"
#include "gromacs/timing/wallcycle.h"
#include "gromacs/utility/template_mp.h"
#include "gromacs/utility/vec.h"
#include "pme_gpu_grid.h"
#include "pme_gpu_internal.h"
#include "pme_gpu_types.h"
#include "pme_gpu_types_host.h"
#include "pme_gpu_types_host_impl.h"

Include dependency graph for pme_gpu_grid_hip.cpp:

Description

Implements PME GPU halo exchange and PME GPU - Host FFT grid conversion functions. These functions are used for PME decomposition in mixed-mode.

Author

Gaurav Garg gauga.nosp@m.rg@n.nosp@m.vidia.nosp@m..com

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

Functions
bool	anonymous_namespace{pme_gpu_grid_hip.cpp}::deviceHas64ParallelExecutionSize (const DeviceInformation &deviceInfo)
template<bool is64ExecutionWidth>
__global__ void	anonymous_namespace{pme_gpu_grid_hip.cpp}::pmeGpuPackHaloExternal (const float *__restrict__ gm_realGrid, float *__restrict__ gm_transferGridUp, float *__restrict__ gm_transferGridDown, float *__restrict__ gm_transferGridLeft, float *__restrict__ gm_transferGridRight, float *__restrict__ gm_transferGridUpLeft, float *__restrict__ gm_transferGridDownLeft, float *__restrict__ gm_transferGridUpRight, float *__restrict__ gm_transferGridDownRight, int overlapSizeUp, int overlapSizeDown, int overlapSizeLeft, int overlapSizeRight, int myGridX, int myGridY, int3 pmeSize)
	A HIP kernel which packs non-contiguous overlap data in all 8 neighboring directions. More...
template<bool is64ExecutionWidth>
__global__ void	anonymous_namespace{pme_gpu_grid_hip.cpp}::pmeGpuUnpackHaloExternal (float *__restrict__ gm_realGrid, const float *__restrict__ gm_transferGridUp, const float *__restrict__ gm_transferGridDown, const float *__restrict__ gm_transferGridLeft, const float *__restrict__ gm_transferGridRight, const float *__restrict__ gm_transferGridUpLeft, const float *__restrict__ gm_transferGridDownLeft, const float *__restrict__ gm_transferGridUpRight, const float *__restrict__ gm_transferGridDownRight, int overlapSizeUp, int overlapSizeDown, int overlapSizeLeft, int overlapSizeRight, int myGridX, int myGridY, int3 pmeSize)
	A HIP kernel which assigns data in halo region in all 8 neighboring directions. More...
template<bool is64ExecutionWidth>
__global__ void	anonymous_namespace{pme_gpu_grid_hip.cpp}::pmeGpuUnpackAndAddHaloInternal (float *__restrict__ gm_realGrid, const float *__restrict__ gm_transferGridUp, const float *__restrict__ gm_transferGridDown, const float *__restrict__ gm_transferGridLeft, const float *__restrict__ gm_transferGridRight, const float *__restrict__ gm_transferGridUpLeft, const float *__restrict__ gm_transferGridDownLeft, const float *__restrict__ gm_transferGridUpRight, const float *__restrict__ gm_transferGridDownRight, int overlapSizeX, int overlapSizeY, int overlapUp, int overlapLeft, int myGridX, int myGridY, int3 pmeSize)
	A HIP kernel which adds grid overlap data received from neighboring ranks. More...
template<bool is64ExecutionWidth>
__global__ void	anonymous_namespace{pme_gpu_grid_hip.cpp}::pmeGpuPackHaloInternal (const float *__restrict__ gm_realGrid, float *__restrict__ gm_transferGridUp, float *__restrict__ gm_transferGridDown, float *__restrict__ gm_transferGridLeft, float *__restrict__ gm_transferGridRight, float *__restrict__ gm_transferGridUpLeft, float *__restrict__ gm_transferGridDownLeft, float *__restrict__ gm_transferGridUpRight, float *__restrict__ gm_transferGridDownRight, int overlapSizeX, int overlapSizeY, int overlapUp, int overlapLeft, int myGridX, int myGridY, int3 pmeSize)
	A HIP kernel which packs non-contiguous overlap data in all 8 neighboring directions. More...
template<bool pmeToFft, bool is64ExecutionWidth>
__global__ void	anonymous_namespace{pme_gpu_grid_hip.cpp}::pmegrid_to_fftgrid (float *__restrict__ gm_realGrid, float *__restrict__ gm_fftGrid, int3 fftNData, int3 fftSize, int3 pmeSize)
	A HIP kernel which copies data from pme grid to FFT grid and back. More...
template<bool is64ExecutionWidth>
void	anonymous_namespace{pme_gpu_grid_hip.cpp}::packHaloDataExternal (const PmeGpu *pmeGpu, int overlapUp, int overlapDown, int overlapLeft, int overlapRight, int myGridX, int myGridY, const ivec &pmeSize, DeviceBuffer< float > realGrid, DeviceBuffer< float > packedGridUp, DeviceBuffer< float > packedGridDown, DeviceBuffer< float > packedGridLeft, DeviceBuffer< float > packedGridRight, DeviceBuffer< float > packedGridUpLeft, DeviceBuffer< float > packedGridDownLeft, DeviceBuffer< float > packedGridUpRight, DeviceBuffer< float > packedGridDownRight)
	Launches HIP kernel to pack non-contiguous external halo data.
template<bool is64ExecutionWidth>
void	anonymous_namespace{pme_gpu_grid_hip.cpp}::packHaloDataInternal (const PmeGpu *pmeGpu, int overlapSizeX, int overlapSizeY, int overlapUp, int overlapLeft, int myGridX, int myGridY, const ivec &pmeSize, DeviceBuffer< float > realGrid, DeviceBuffer< float > packedGridUp, DeviceBuffer< float > packedGridDown, DeviceBuffer< float > packedGridLeft, DeviceBuffer< float > packedGridRight, DeviceBuffer< float > packedGridUpLeft, DeviceBuffer< float > packedGridDownLeft, DeviceBuffer< float > packedGridUpRight, DeviceBuffer< float > packedGridDownRight)
	Launches HIP kernel to pack non-contiguous internal halo data.
template<bool is64ExecutionWidth>
void	anonymous_namespace{pme_gpu_grid_hip.cpp}::unpackAndAddHaloDataInternal (const PmeGpu *pmeGpu, int overlapSizeX, int overlapSizeY, int overlapUp, int overlapLeft, int myGridX, int myGridY, const ivec &pmeSize, DeviceBuffer< float > realGrid, DeviceBuffer< float > packedGridUp, DeviceBuffer< float > packedGridDown, DeviceBuffer< float > packedGridLeft, DeviceBuffer< float > packedGridRight, DeviceBuffer< float > packedGridUpLeft, DeviceBuffer< float > packedGridDownLeft, DeviceBuffer< float > packedGridUpRight, DeviceBuffer< float > packedGridDownRight)
	Launches HIP kernel to unpack and reduce overlap data.
template<bool is64ExecutionWidth>
void	anonymous_namespace{pme_gpu_grid_hip.cpp}::unpackHaloDataExternal (const PmeGpu *pmeGpu, int overlapUp, int overlapDown, int overlapLeft, int overlapRight, int myGridX, int myGridY, const ivec &pmeSize, DeviceBuffer< float > realGrid, DeviceBuffer< float > packedGridUp, DeviceBuffer< float > packedGridDown, DeviceBuffer< float > packedGridLeft, DeviceBuffer< float > packedGridRight, DeviceBuffer< float > packedGridUpLeft, DeviceBuffer< float > packedGridDownLeft, DeviceBuffer< float > packedGridUpRight, DeviceBuffer< float > packedGridDownRight)
	Launches HIP kernel to initialize overlap data.
void	anonymous_namespace{pme_gpu_grid_hip.cpp}::receiveAndSend (DeviceBuffer< float > sendBuf, int sendCount, int dest, MPI_Request *sendRequest, DeviceBuffer< float > recvBuf, int recvCount, int src, MPI_Request *recvRequest, int tag, MPI_Comm comm)
	utility function to send and recv halo data from neighboring ranks
void	pmeGpuGridHaloExchange (const PmeGpu *pmeGpu, gmx_wallcycle *wcycle)
	Grid Halo exchange after PME spread ToDo: Current implementation transfers halo region from/to only immediate neighbours And, expects that overlapSize <= local grid width. Implement exchange with multiple neighbors to remove this limitation ToDo: Current implementation synchronizes pmeStream to make sure data is ready on GPU after spread. Consider using events for this synchnozation. More...
void	pmeGpuGridHaloExchangeReverse (const PmeGpu *pmeGpu, gmx_wallcycle *wcycle)
	Grid reverse Halo exchange before PME gather ToDo: Current implementation transfers halo region from/to only immediate neighbours And, expects that overlapSize <= local grid width. Implement exchange with multiple neighbors to remove this limitation ToDo: Current implementation synchronizes pmeStream to make sure data is ready on GPU after FFT to PME grid conversion. Consider using events for this synchnozation. More...
template<bool pmeToFft>
void	convertPmeGridToFftGrid (const PmeGpu *pmeGpu, float *h_fftRealGrid, gmx_parallel_3dfft *fftSetup, const int gridIndex)
	Copy PME Grid with overlap region to host FFT grid and vice-versa. Used in mixed mode PME decomposition. More...
template<bool pmeToFft>
void	convertPmeGridToFftGrid (const PmeGpu *pmeGpu, DeviceBuffer< float > *d_fftRealGrid, const int gridIndex)
	Copy PME Grid with overlap region to device FFT grid and vice-versa. Used in full GPU PME decomposition. More...
template void	convertPmeGridToFftGrid< true > (const PmeGpu *pmeGpu, float *h_fftRealGrid, gmx_parallel_3dfft *fftSetup, const int gridIndex)
template void	convertPmeGridToFftGrid< false > (const PmeGpu *pmeGpu, float *h_fftRealGrid, gmx_parallel_3dfft *fftSetup, const int gridIndex)
template void	convertPmeGridToFftGrid< true > (const PmeGpu *pmeGpu, DeviceBuffer< float > *d_fftRealGrid, const int gridIndex)
template void	convertPmeGridToFftGrid< false > (const PmeGpu *pmeGpu, DeviceBuffer< float > *d_fftRealGrid, const int gridIndex)

Variables
template<bool is64ExecutionWidth>
constexpr int	anonymous_namespace{pme_gpu_grid_hip.cpp}::sc_subGroupSizeX = is64ExecutionWidth ? 64 : 32
	Sub-group size for conversion kernels. More...
template<bool is64ExecutionWidth>
constexpr int	anonymous_namespace{pme_gpu_grid_hip.cpp}::sc_subGroupSizeY = is64ExecutionWidth ? 2 : 4
constexpr int	anonymous_namespace{pme_gpu_grid_hip.cpp}::sc_subGroupSizeZ = 1

Function Documentation

template<bool pmeToFft>

void convertPmeGridToFftGrid	(	const PmeGpu *	pmeGpu,
		float *	h_fftRealGrid,
		gmx_parallel_3dfft *	fftSetup,
		int	gridIndex
	)

Copy PME Grid with overlap region to host FFT grid and vice-versa. Used in mixed mode PME decomposition.

Parameters

[in]	pmeGpu	The PME GPU structure.
[in]	h_fftRealGrid	FFT grid on host
[in]	fftSetup	Host FFT setup structure
[in]	gridIndex	Grid index which is to be converted

Template Parameters

pmeToFft

A boolean which tells if this is conversion from PME grid to FFT grid or reverse

template<bool pmeToFft>

void convertPmeGridToFftGrid	(	const PmeGpu *	pmeGpu,
		DeviceBuffer< float > *	d_fftRealGrid,
		int	gridIndex
	)

Copy PME Grid with overlap region to device FFT grid and vice-versa. Used in full GPU PME decomposition.

Parameters

[in]	pmeGpu	The PME GPU structure.
[in]	d_fftRealGrid	FFT grid on device
[in]	gridIndex	Grid index which is to be converted

Template Parameters

pmeToFft

A boolean which tells if this is conversion from PME grid to FFT grid or reverse

void pmeGpuGridHaloExchange	(	const PmeGpu *	pmeGpu,
		gmx_wallcycle *	wcycle
	)

Grid Halo exchange after PME spread ToDo: Current implementation transfers halo region from/to only immediate neighbours And, expects that overlapSize <= local grid width. Implement exchange with multiple neighbors to remove this limitation ToDo: Current implementation synchronizes pmeStream to make sure data is ready on GPU after spread. Consider using events for this synchnozation.

Parameters

[in]	pmeGpu	The PME GPU structure.
[in]	wcycle	The wallclock counter.

void pmeGpuGridHaloExchangeReverse	(	const PmeGpu *	pmeGpu,
		gmx_wallcycle *	wcycle
	)

Grid reverse Halo exchange before PME gather ToDo: Current implementation transfers halo region from/to only immediate neighbours And, expects that overlapSize <= local grid width. Implement exchange with multiple neighbors to remove this limitation ToDo: Current implementation synchronizes pmeStream to make sure data is ready on GPU after FFT to PME grid conversion. Consider using events for this synchnozation.

Parameters

[in]	pmeGpu	The PME GPU structure.
[in]	wcycle	The wallclock counter.