Gromacs  2026.3
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
Macros | Functions | Variables
pme_spread_hip.cpp File Reference
#include "gmxpre.h"
#include <cassert>
#include "gromacs/gpu_utils/hip_kernel_utils.h"
#include "gromacs/gpu_utils/typecasts_cuda_hip.h"
#include "pme_gpu_calculate_splines_hip.h"
#include "pme_gpu_constants.h"
#include "pme_gpu_types.h"
+ Include dependency graph for pme_spread_hip.cpp:

Description

Implements PME GPU spline calculation and charge spreading in HIP. TODO: consider always pre-sorting particles (as in DD case).

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

Macros

#define INSTANTIATE_3(order, computeSplines, spreadCharges, numGrids, writeGlobal, threadsPerAtom, parallelExecutionWidth)
 Kernel instantiations. More...
 
#define INSTANTIATE_2(order, numGrids, threadsPerAtom, parallelExecutionWidth)
 
#define INSTANTIATE(order, parallelExecutionWidth)
 

Functions

template<int order, bool wrapX, bool wrapY, int gridIndex, ThreadsPerAtom threadsPerAtom, int parallelExecutionWidth>
__device__ __forceinline__ void spread_charges (const PmeGpuKernelParams kernelParams, const float *atomCharge, const int *__restrict__ sm_gridlineIndices, const float *__restrict__ sm_theta)
 Charge spreading onto the grid. This corresponds to the CPU function spread_coefficients_bsplines_thread(). Optional second stage of the spline_and_spread_kernel. More...
 
template<int order, bool computeSplines, bool spreadCharges, bool wrapX, bool wrapY, int numGrids, bool writeGlobal, ThreadsPerAtom threadsPerAtom, int parallelExecutionWidth>
 __attribute__ ((amdgpu_flat_work_group_size(sc_spreadMaxThreadsPerBlock< parallelExecutionWidth >, sc_spreadMaxThreadsPerBlock< parallelExecutionWidth >))) __global__ void pmeSplineAndSpreadKernel(const PmeGpuKernelParams kernelParams)
 A spline computation and charge spreading kernel function. More...
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 1, true, ThreadsPerAtom::Order, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, false, true, true, 1, true, ThreadsPerAtom::Order, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, false, true, true, true, 1, true, ThreadsPerAtom::Order, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 1, false, ThreadsPerAtom::Order, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 1, true, ThreadsPerAtom::OrderSquared, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, false, true, true, 1, true, ThreadsPerAtom::OrderSquared, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, false, true, true, true, 1, true, ThreadsPerAtom::OrderSquared, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 1, false, ThreadsPerAtom::OrderSquared, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 2, true, ThreadsPerAtom::Order, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, false, true, true, 2, true, ThreadsPerAtom::Order, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, false, true, true, true, 2, true, ThreadsPerAtom::Order, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 2, false, ThreadsPerAtom::Order, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 2, true, ThreadsPerAtom::OrderSquared, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, false, true, true, 2, true, ThreadsPerAtom::OrderSquared, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, false, true, true, true, 2, true, ThreadsPerAtom::OrderSquared, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 2, false, ThreadsPerAtom::OrderSquared, 32 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 1, true, ThreadsPerAtom::Order, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, false, true, true, 1, true, ThreadsPerAtom::Order, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, false, true, true, true, 1, true, ThreadsPerAtom::Order, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 1, false, ThreadsPerAtom::Order, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 1, true, ThreadsPerAtom::OrderSquared, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, false, true, true, 1, true, ThreadsPerAtom::OrderSquared, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, false, true, true, true, 1, true, ThreadsPerAtom::OrderSquared, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 1, false, ThreadsPerAtom::OrderSquared, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 2, true, ThreadsPerAtom::Order, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, false, true, true, 2, true, ThreadsPerAtom::Order, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, false, true, true, true, 2, true, ThreadsPerAtom::Order, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 2, false, ThreadsPerAtom::Order, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 2, true, ThreadsPerAtom::OrderSquared, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, false, true, true, 2, true, ThreadsPerAtom::OrderSquared, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, false, true, true, true, 2, true, ThreadsPerAtom::OrderSquared, 64 > (PmeGpuKernelParams kernelParams)
 
template __global__ void pmeSplineAndSpreadKernel< 4, true, true, true, true, 2, false, ThreadsPerAtom::OrderSquared, 64 > (PmeGpuKernelParams kernelParams)
 

Variables

template<int parallelExecutionWidth>
static constexpr int sc_spreadHipMaxWarpsPerBlock = (parallelExecutionWidth == 64) ? 8 : 4
 
template<int parallelExecutionWidth>
static constexpr int sc_spreadMaxThreadsPerBlock
 

Macro Definition Documentation

#define INSTANTIATE (   order,
  parallelExecutionWidth 
)
Value:
INSTANTIATE_2(order, 1, ThreadsPerAtom::Order, parallelExecutionWidth); \
INSTANTIATE_2(order, 1, ThreadsPerAtom::OrderSquared, parallelExecutionWidth); \
INSTANTIATE_2(order, 2, ThreadsPerAtom::Order, parallelExecutionWidth); \
INSTANTIATE_2(order, 2, ThreadsPerAtom::OrderSquared, parallelExecutionWidth);
ThreadsPerAtom::Order
Use a number of threads equal to the PME order (ie. 4)
ThreadsPerAtom::OrderSquared
Use a number of threads equal to the square of the PME order (ie. 16)
#define INSTANTIATE_2 (   order,
  numGrids,
  threadsPerAtom,
  parallelExecutionWidth 
)
Value:
INSTANTIATE_3(order, true, true, numGrids, true, threadsPerAtom, parallelExecutionWidth); \
INSTANTIATE_3(order, true, false, numGrids, true, threadsPerAtom, parallelExecutionWidth); \
INSTANTIATE_3(order, false, true, numGrids, true, threadsPerAtom, parallelExecutionWidth); \
INSTANTIATE_3(order, true, true, numGrids, false, threadsPerAtom, parallelExecutionWidth);
gmx::numGrids
static int numGrids(const GridSet::DomainSetup &domainSetup)
Returns the number of search grids.
Definition: gridset.cpp:67
INSTANTIATE_3
#define INSTANTIATE_3(order, computeSplines, spreadCharges, numGrids, writeGlobal, threadsPerAtom, parallelExecutionWidth)
Kernel instantiations.
Definition: pme_spread_hip.cpp:311
#define INSTANTIATE_3 (   order,
  computeSplines,
  spreadCharges,
  numGrids,
  writeGlobal,
  threadsPerAtom,
  parallelExecutionWidth 
)
Value:
template __global__ void \
pmeSplineAndSpreadKernel<order, computeSplines, spreadCharges, true, true, numGrids, writeGlobal, threadsPerAtom, parallelExecutionWidth>( \
PmeGpuKernelParams kernelParams);
PmeGpuKernelParams
int PmeGpuKernelParams
A dummy typedef for the GPU kernel arguments data placeholder on non-GPU builds.
Definition: pme_gpu_types_host.h:79

Kernel instantiations.

Function Documentation

template<int order, bool computeSplines, bool spreadCharges, bool wrapX, bool wrapY, int numGrids, bool writeGlobal, ThreadsPerAtom threadsPerAtom, int parallelExecutionWidth>
__attribute__ ( (amdgpu_flat_work_group_size(sc_spreadMaxThreadsPerBlock< parallelExecutionWidth >,sc_spreadMaxThreadsPerBlock< parallelExecutionWidth >))  ) const

A spline computation and charge spreading kernel function.

Two tuning parameters can be used for additional performance. For small systems and for debugging writeGlobal should be used removing the need to recalculate the theta values in the gather kernel. Similarly for useOrderThreads large systems order threads per atom gives higher performance than order*order threads

Template Parameters
orderPME interpolation order.
computeSplinesA boolean which tells if the spline parameter and gridline indices' computation should be performed.
spreadChargesA boolean which tells if the charge spreading should be performed.
wrapXA boolean which tells if the grid overlap in dimension X should be wrapped.
wrapYA boolean which tells if the grid overlap in dimension Y should be wrapped.
numGridsThe number of grids to use in the kernel. Can be 1 or 2.
writeGlobalA boolean which tells if the theta values and gridlines should be written to global memory.
threadsPerAtomHow many threads work on each atom param[in] kernelParams Input PME HIP data in constant memory.
template<int order, bool wrapX, bool wrapY, int gridIndex, ThreadsPerAtom threadsPerAtom, int parallelExecutionWidth>
__device__ __forceinline__ void spread_charges ( const PmeGpuKernelParams  kernelParams,
const float *  atomCharge,
const int *__restrict__  sm_gridlineIndices,
const float *__restrict__  sm_theta 
)

Charge spreading onto the grid. This corresponds to the CPU function spread_coefficients_bsplines_thread(). Optional second stage of the spline_and_spread_kernel.

Template Parameters
orderPME interpolation order.
wrapXWhether the grid overlap in dimension X should be wrapped.
wrapYWhether the grid overlap in dimension Y should be wrapped.
gridIndexThe index of the grid to use in the kernel.
threadsPerAtomHow many threads work on each atom
Parameters
[in]kernelParamsInput PME HIP data in constant memory.
[in]atomChargeAtom charge/coefficient of atom processed by thread.
[in]sm_gridlineIndicesAtom gridline indices in the shared memory.
[in]sm_thetaAtom spline values in the shared memory.

Variable Documentation

template<int parallelExecutionWidth>
constexpr int sc_spreadMaxThreadsPerBlock
static
Initial value:
=
sc_spreadHipMaxWarpsPerBlock<parallelExecutionWidth> * parallelExecutionWidth
Generated on Thu Jun 25 2026 13:58:45 for Gromacs by   doxygen 1.8.5