ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
|
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>
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 ¶ms, 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 ¶ms, 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. | |
P3M electrostatics on GPU.
The corresponding header file is p3m_gpu_cuda.cuh.
Definition in file p3m_gpu_cuda.cu.
#define FFT_BACK_FFT cufftExecC2R |
Definition at line 39 of file p3m_gpu_cuda.cu.
#define FFT_FORW_FFT cufftExecR2C |
Definition at line 38 of file p3m_gpu_cuda.cu.
#define FFT_PLAN_BACK_FLAG CUFFT_C2R |
Definition at line 41 of file p3m_gpu_cuda.cu.
#define FFT_PLAN_FORW_FLAG CUFFT_R2C |
Definition at line 40 of file p3m_gpu_cuda.cu.
#define FFT_TYPE_COMPLEX cufftComplex |
Definition at line 37 of file p3m_gpu_cuda.cu.
#define P3M_GPU_FLOAT |
Definition at line 32 of file p3m_gpu_cuda.cu.
#define REAL_TYPE float |
Definition at line 36 of file p3m_gpu_cuda.cu.
|
static |
Definition at line 167 of file p3m_gpu_cuda.cu.
References P3MGpuData::alpha, P3MGpuData::box, P3MGpuData::mesh, P3M_BRILLOUIN, REAL_TYPE, and math::sinc().
__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, and REAL_TYPE.
Referenced by p3m_gpu_add_farfield_force().
__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, and anonymous_namespace{p3m_gpu_cuda.cu}::linear_index_k().
Referenced by p3m_gpu_add_farfield_force().
__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, and wrap_index().
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, and REAL_TYPE.
Referenced by p3m_gpu_add_farfield_force().
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(), and params.
Referenced by p3m_gpu_add_farfield_force().
__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, and wrap_index().
__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.
|
static |
Definition at line 141 of file p3m_gpu_cuda.cu.
Referenced by assign_charges(), and assign_forces().
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 672 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(), and REAL_TYPE.
Referenced by CoulombP3MImpl< FloatType, Architecture >::add_long_range_forces_gpu().
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().
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().
|
inline |
Definition at line 297 of file p3m_gpu_cuda.cu.
Referenced by assign_charge_kernel(), and assign_forces_kernel().