pme_solve_hip.cpp File Reference

#include "gmxpre.h"
#include <cassert>
#include <hip/hip_math_constants.h>
#include "gromacs/gpu_utils/gputraits.h"
#include "gromacs/gpu_utils/hip_kernel_utils.h"
#include "gromacs/gpu_utils/hip_sycl_kernel_utils.h"
#include "gromacs/gpu_utils/vectype_ops_hip.h"
#include "pme_gpu_constants.h"
#include "pme_gpu_internal.h"
#include "pme_gpu_types.h"

Include dependency graph for pme_solve_hip.cpp:

Description

Implements PME GPU Fourier grid solving in HIP.

Author

Aleksei Iupinov a.yup.nosp@m.inov.nosp@m.@gmai.nosp@m.l.co.nosp@m.m

Macros
#define	INSTANTIATE(parallelExecutionWidth)
	Kernel class instantiations. More...

Functions
template<GridOrdering gridOrdering, bool computeEnergyAndVirial, const int gridIndex, int parallelExecutionWidth>
	__attribute__ ((amdgpu_flat_work_group_size(sc_solveMaxThreadsPerBlock< parallelExecutionWidth >, sc_solveMaxThreadsPerBlock< parallelExecutionWidth >))) __global__ void pmeSolveKernel(const struct PmeGpuKernelParams kernelParams)
	PME complex grid solver kernel function. More...
template __global__ void	pmeSolveKernel< GridOrdering::XYZ, false, 0, 32 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::XYZ, true, 0, 32 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::YZX, false, 0, 32 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::YZX, true, 0, 32 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::XYZ, false, 1, 32 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::XYZ, true, 1, 32 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::YZX, false, 1, 32 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::YZX, true, 1, 32 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::XYZ, false, 0, 64 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::XYZ, true, 0, 64 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::YZX, false, 0, 64 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::YZX, true, 0, 64 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::XYZ, false, 1, 64 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::XYZ, true, 1, 64 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::YZX, false, 1, 64 > (PmeGpuKernelParams kernelParams)
template __global__ void	pmeSolveKernel< GridOrdering::YZX, true, 1, 64 > (PmeGpuKernelParams kernelParams)

Variables
template<int parallelExecutionWidth>
static constexpr int	sc_solveMaxThreadsPerBlock = c_solveMaxWarpsPerBlock * parallelExecutionWidth

Macro Definition Documentation

#define INSTANTIATE

(

parallelExecutionWidth

)

Value:

template __global__ void pmeSolveKernel<GridOrdering::XYZ, false, 0, parallelExecutionWidth>( \
            PmeGpuKernelParams kernelParams);                                                     \
    template __global__ void pmeSolveKernel<GridOrdering::XYZ, true, 0, parallelExecutionWidth>(  \
            PmeGpuKernelParams kernelParams);                                                     \
    template __global__ void pmeSolveKernel<GridOrdering::YZX, false, 0, parallelExecutionWidth>( \
            PmeGpuKernelParams kernelParams);                                                     \
    template __global__ void pmeSolveKernel<GridOrdering::YZX, true, 0, parallelExecutionWidth>(  \
            PmeGpuKernelParams kernelParams);                                                     \
    template __global__ void pmeSolveKernel<GridOrdering::XYZ, false, 1, parallelExecutionWidth>( \
            PmeGpuKernelParams kernelParams);                                                     \
    template __global__ void pmeSolveKernel<GridOrdering::XYZ, true, 1, parallelExecutionWidth>(  \
            PmeGpuKernelParams kernelParams);                                                     \
    template __global__ void pmeSolveKernel<GridOrdering::YZX, false, 1, parallelExecutionWidth>( \
            PmeGpuKernelParams kernelParams);                                                     \
    template __global__ void pmeSolveKernel<GridOrdering::YZX, true, 1, parallelExecutionWidth>(  \
            PmeGpuKernelParams kernelParams);

Kernel class instantiations.

Function Documentation

template<GridOrdering gridOrdering, bool computeEnergyAndVirial, const int gridIndex, int parallelExecutionWidth>

__attribute__

(

(amdgpu_flat_work_group_size(sc_solveMaxThreadsPerBlock< parallelExecutionWidth >,sc_solveMaxThreadsPerBlock< parallelExecutionWidth >))

)

const

PME complex grid solver kernel function.

Template Parameters

gridOrdering	Specifies the dimension ordering of the complex grid.
computeEnergyAndVirial	Tells if the reciprocal energy and virial should be computed.
gridIndex	The index of the grid to use in the kernel.
parallelExecutionWidth	The device parallel execution size to use

param[in] kernelParams Input PME HIP data in constant memory.