P3M electrostatics on GPU. More...

#include <config/config.hpp>
#include "electrostatics/p3m_gpu_cuda.cuh"
#include "cuda/utils.cuh"
#include "p3m/math.hpp"
#include "system/System.hpp"
#include <utils/math/bspline.hpp>
#include <utils/math/int_pow.hpp>
#include <utils/math/sqr.hpp>
#include <cuda.h>
#include <cufft.h>
#include <algorithm>
#include <cassert>
#include <cstddef>
#include <limits>
#include <numbers>
#include <stdexcept>

Include dependency graph for p3m_gpu_cuda.cu:

Go to the source code of this file.

Classes
struct	P3MGpuData

struct	P3MGpuFftPlan

struct	P3MGpuParams

Namespaces
namespace	anonymous_namespace{p3m_gpu_cuda.cu}

Macros
#define	ESPRESSO_P3M_GPU_FLOAT

#define	REAL_TYPE float

#define	FFT_TYPE_COMPLEX cufftComplex

#define	FFT_FORW_FFT cufftExecR2C

#define	FFT_BACK_FFT cufftExecC2R

#define	FFT_PLAN_FORW_FLAG CUFFT_R2C

#define	FFT_PLAN_BACK_FLAG CUFFT_C2R

Functions
static auto	p3m_calc_blocks (unsigned int cao, std::size_t n_part)

dim3	p3m_make_grid (unsigned int n_blocks)

template<int cao>
__device__ static void	Aliasing_sums_ik (const P3MGpuData p, int NX, int NY, int NZ, REAL_TYPE Zaehler, REAL_TYPE Nenner)

template<int cao>
__global__ void	calculate_influence_function_device (const P3MGpuData p)

__device__ auto	anonymous_namespace{p3m_gpu_cuda.cu}::linear_index_r (P3MGpuData const &p, int i, int j, int k)

__device__ auto	anonymous_namespace{p3m_gpu_cuda.cu}::linear_index_k (P3MGpuData const &p, int i, int j, int k)

__global__ void	apply_diff_op (const P3MGpuData p)

__device__ int	wrap_index (const int ind, const int mesh)

__global__ void	apply_influence_function (const P3MGpuData p)

template<int cao, bool shared>
__global__ void	assign_charge_kernel (P3MGpuData const params, float const const __restrict__ part_pos, float const const __restrict__ part_q, unsigned int const parts_per_block)

void	assign_charges (P3MGpuData const &params, float const const __restrict__ part_pos, float const const __restrict__ part_q)

template<int cao, bool shared>
__global__ void	assign_forces_kernel (P3MGpuData const params, float const const __restrict__ part_pos, float const const __restrict__ part_q, float *const __restrict__ part_f, REAL_TYPE prefactor, unsigned int const parts_per_block)

void	assign_forces (P3MGpuData const &params, float const const __restrict__ part_pos, float const const __restrict__ part_q, float *const __restrict__ part_f, REAL_TYPE const prefactor)

void	p3m_gpu_init (std::shared_ptr< P3MGpuParams > &data, int cao, Utils::Vector3i const &mesh, double alpha, Utils::Vector3d const &box_l, std::size_t n_part)
	Initialize the internal data structure of the P3M GPU.

void	p3m_gpu_add_farfield_force (P3MGpuParams &data, GpuParticleData &gpu, double prefactor, std::size_t n_part)
	The long-range part of the P3M algorithm.

Detailed Description

P3M electrostatics on GPU.

The corresponding header file is p3m_gpu_cuda.cuh.

Definition in file p3m_gpu_cuda.cu.

Macro Definition Documentation

◆ ESPRESSO_P3M_GPU_FLOAT

#define ESPRESSO_P3M_GPU_FLOAT

Definition at line 32 of file p3m_gpu_cuda.cu.

◆ FFT_BACK_FFT

#define FFT_BACK_FFT cufftExecC2R

Definition at line 39 of file p3m_gpu_cuda.cu.

◆ FFT_FORW_FFT

#define FFT_FORW_FFT cufftExecR2C

Definition at line 38 of file p3m_gpu_cuda.cu.

◆ FFT_PLAN_BACK_FLAG

#define FFT_PLAN_BACK_FLAG CUFFT_C2R

Definition at line 41 of file p3m_gpu_cuda.cu.

◆ FFT_PLAN_FORW_FLAG

#define FFT_PLAN_FORW_FLAG CUFFT_R2C

Definition at line 40 of file p3m_gpu_cuda.cu.

◆ FFT_TYPE_COMPLEX

#define FFT_TYPE_COMPLEX cufftComplex

Definition at line 37 of file p3m_gpu_cuda.cu.

◆ REAL_TYPE

#define REAL_TYPE float

Definition at line 36 of file p3m_gpu_cuda.cu.

Function Documentation

◆ Aliasing_sums_ik()

template<int cao>

__device__ static void Aliasing_sums_ik	(	const P3MGpuData	p,
		int	NX,
		int	NY,
		int	NZ,
		REAL_TYPE *	Zaehler,
		REAL_TYPE *	Nenner
	)

static

Definition at line 167 of file p3m_gpu_cuda.cu.

References P3MGpuData::alpha, P3MGpuData::box, P3MGpuData::mesh, P3M_BRILLOUIN, REAL_TYPE, math::sinc(), and stream.

◆ apply_diff_op()

__global__ void apply_diff_op ( const P3MGpuData p )

Definition at line 265 of file p3m_gpu_cuda.cu.

References P3MGpuData::box, P3MGpuData::charge_mesh, FFT_TYPE_COMPLEX, P3MGpuData::force_mesh_x, P3MGpuData::force_mesh_y, P3MGpuData::force_mesh_z, anonymous_namespace{p3m_gpu_cuda.cu}::linear_index_k(), P3MGpuData::mesh, REAL_TYPE, and stream.

Referenced by p3m_gpu_add_farfield_force().

◆ apply_influence_function()

__global__ void apply_influence_function ( const P3MGpuData p )

Definition at line 305 of file p3m_gpu_cuda.cu.

References P3MGpuData::charge_mesh, P3MGpuData::G_hat, anonymous_namespace{p3m_gpu_cuda.cu}::linear_index_k(), and stream.

Referenced by p3m_gpu_add_farfield_force().

◆ assign_charge_kernel()

template<int cao, bool shared>

__global__ void assign_charge_kernel	(	P3MGpuData const	params,
		float const *const __restrict__	part_pos,
		float const *const __restrict__	part_q,
		unsigned int const	parts_per_block
	)

Definition at line 315 of file p3m_gpu_cuda.cu.

References anonymous_namespace{p3m_gpu_cuda.cu}::linear_index_r(), params, REAL_TYPE, stream, and wrap_index().

◆ assign_charges()

void assign_charges	(	P3MGpuData const &	params,
		float const *const __restrict__	part_pos,
		float const *const __restrict__	part_q
	)

Definition at line 377 of file p3m_gpu_cuda.cu.

References block(), cuda_check_errors_exit(), p3m_calc_blocks(), p3m_make_grid(), params, REAL_TYPE, and stream.

Referenced by p3m_gpu_add_farfield_force().

◆ assign_forces()

void assign_forces	(	P3MGpuData const &	params,
		float const *const __restrict__	part_pos,
		float const *const __restrict__	part_q,
		float *const __restrict__	part_f,
		REAL_TYPE const	prefactor
	)

Definition at line 492 of file p3m_gpu_cuda.cu.

References block(), cuda_check_errors_exit(), p3m_calc_blocks(), p3m_make_grid(), params, and stream.

Referenced by p3m_gpu_add_farfield_force().

◆ assign_forces_kernel()

template<int cao, bool shared>

__global__ void assign_forces_kernel	(	P3MGpuData const	params,
		float const *const __restrict__	part_pos,
		float const *const __restrict__	part_q,
		float *const __restrict__	part_f,
		REAL_TYPE	prefactor,
		unsigned int const	parts_per_block
	)

Definition at line 424 of file p3m_gpu_cuda.cu.

References anonymous_namespace{p3m_gpu_cuda.cu}::linear_index_r(), params, REAL_TYPE, stream, and wrap_index().

◆ calculate_influence_function_device()

template<int cao>

__global__ void calculate_influence_function_device ( const P3MGpuData p )

Definition at line 213 of file p3m_gpu_cuda.cu.

References P3MGpuData::box, P3MGpuData::G_hat, P3MGpuData::mesh, REAL_TYPE, and stream.

◆ p3m_calc_blocks()

static auto p3m_calc_blocks	(	unsigned int	cao,
		std::size_t	n_part
	)

static

Definition at line 141 of file p3m_gpu_cuda.cu.

References stream.

Referenced by assign_charges(), and assign_forces().

◆ p3m_gpu_add_farfield_force()

void p3m_gpu_add_farfield_force	(	P3MGpuParams &	data,
		GpuParticleData &	gpu,
		double	prefactor,
		std::size_t	n_part
	)

The long-range part of the P3M algorithm.

Definition at line 671 of file p3m_gpu_cuda.cu.

References apply_diff_op(), apply_influence_function(), assign_charges(), assign_forces(), P3MGpuFftPlan::back_plan, cuda_safe_mem, FFT_BACK_FFT, FFT_FORW_FFT, P3MGpuFftPlan::forw_plan, GpuParticleData::get_particle_charges_device(), GpuParticleData::get_particle_forces_device(), GpuParticleData::get_particle_positions_device(), KERNELCALL, P3MGpuData::n_part, P3MGpuParams::p3m_fft, P3MGpuParams::p3m_gpu_data, Utils::product(), REAL_TYPE, and stream.

Referenced by CoulombP3MHeffte< FloatType, Architecture, FFTConfig >::add_long_range_forces_gpu().

◆ p3m_gpu_init()

void p3m_gpu_init	(	std::shared_ptr< P3MGpuParams > &	data,
		int	cao,
		Utils::Vector3i const &	mesh,
		double	alpha,
		Utils::Vector3d const &	box_l,
		std::size_t	n_part
	)

Initialize the internal data structure of the P3M GPU.

Mainly allocation on the device and influence function calculation. Be advised: this needs mesh^3*5*sizeof(REAL_TYPE) of device memory. We use real to complex FFTs, so the size of the reciprocal mesh is (cuFFT convention) Nx * Ny * ( Nz /2 + 1 ).

Definition at line 549 of file p3m_gpu_cuda.cu.

References block(), Utils::Vector< T, N >::broadcast(), cuda_safe_mem, FFT_PLAN_BACK_FLAG, FFT_PLAN_FORW_FLAG, FFT_TYPE_COMPLEX, KERNELCALL, REAL_TYPE, and stream.

Referenced by CoulombP3MHeffte< FloatType, Architecture, FFTConfig >::init_gpu_kernels().

◆ p3m_make_grid()

dim3 p3m_make_grid ( unsigned int n_blocks )

Definition at line 154 of file p3m_gpu_cuda.cu.

References stream.

Referenced by assign_charges(), and assign_forces().

◆ wrap_index()

__device__ int wrap_index	(	const int	ind,
		const int	mesh
	)

inline

Definition at line 297 of file p3m_gpu_cuda.cu.

Referenced by assign_charge_kernel(), and assign_forces_kernel().

Classes

Namespaces

Macros

Functions

Detailed Description

Macro Definition Documentation

◆ ESPRESSO_P3M_GPU_FLOAT

◆ FFT_BACK_FFT

◆ FFT_FORW_FFT

◆ FFT_PLAN_BACK_FLAG

◆ FFT_PLAN_FORW_FLAG

◆ FFT_TYPE_COMPLEX

◆ REAL_TYPE

Function Documentation

◆ Aliasing_sums_ik()

◆ apply_diff_op()

◆ apply_influence_function()

◆ assign_charge_kernel()

◆ assign_charges()

◆ assign_forces()

◆ assign_forces_kernel()

◆ calculate_influence_function_device()

◆ p3m_calc_blocks()

◆ p3m_gpu_add_farfield_force()

◆ p3m_gpu_init()

◆ p3m_make_grid()

◆ wrap_index()