ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
Loading...
Searching...
No Matches
p3m_gpu_cuda.cu File Reference

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 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

◆ 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.

◆ P3M_GPU_FLOAT

#define P3M_GPU_FLOAT

Definition at line 32 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

◆ apply_diff_op()

◆ apply_influence_function()

__global__ void apply_influence_function ( const P3MGpuData  p)

◆ 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 
)

◆ assign_charges()

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

◆ 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 
)

◆ 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 
)

◆ 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, and REAL_TYPE.

◆ 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.

Referenced by assign_charges(), and assign_forces().

◆ p3m_gpu_add_farfield_force()

◆ 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 Utils::Array< T, N >::begin(), block(), Utils::Vector< T, N >::broadcast(), cuda_safe_mem, Utils::Array< T, N >::end(), FFT_PLAN_BACK_FLAG, FFT_PLAN_FORW_FLAG, FFT_TYPE_COMPLEX, KERNELCALL, and REAL_TYPE.

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

◆ p3m_make_grid()

dim3 p3m_make_grid ( unsigned int  n_blocks)

Definition at line 154 of file p3m_gpu_cuda.cu.

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().