50#include "communication.hpp"
54#include "system/System.hpp"
62#include <boost/mpi/collectives/all_reduce.hpp>
63#include <boost/mpi/collectives/reduce.hpp>
81#error "The FFTW3 library shouldn't be visible in this translation unit"
84template <
typename FloatType, Arch Architecture>
87 double local_mu2 = 0.;
89 for (
auto const &p : get_system().
cell_structure->local_particles()) {
91 local_mu2 += p.calc_dip().norm2();
96 boost::mpi::all_reduce(
comm_cart, local_mu2, dp3m.sum_mu2, std::plus<>());
97 boost::mpi::all_reduce(
comm_cart, local_n, dp3m.sum_dip_part, std::plus<>());
101 int n_c_part,
double sum_q2,
double alpha_L);
104 int n_c_part,
double sum_q2,
111 double sum_q2,
double x1,
double x2,
double xacc,
112 double tuned_accuracy);
114template <
typename FloatType, Arch Architecture>
118 auto const &box_geo = *get_system().
box_geo;
119 auto const node_phi =
120 grid_influence_function_self_energy<FloatType, P3M_BRILLOUIN>(
121 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop, dp3m.g_energy);
124 boost::mpi::reduce(
comm_cart, node_phi, phi, std::plus<>(), 0);
125 phi /= 3. * box_geo.length()[0] * Utils::int_pow<3>(dp3m.params.mesh[0]);
126 return phi * std::numbers::pi;
129template <
typename FloatType, Arch Architecture>
132 assert(dp3m.params.cao >= 1 and dp3m.params.cao <= 7);
133 assert(dp3m.params.alpha > 0.);
135 auto const &system = get_system();
136 auto const &box_geo = *system.box_geo;
137 auto const &local_geo = *system.local_geo;
138 auto const verlet_skin = system.cell_structure->get_verlet_skin();
140 dp3m.params.cao3 = Utils::int_pow<3>(dp3m.params.cao);
141 dp3m.params.recalc_a_ai_cao_cut(box_geo.length());
144 dp3m.local_mesh.calc_local_ca_mesh(dp3m.params, local_geo, verlet_skin, 0.);
145 dp3m.fft_buffers->init_halo();
146 dp3m.fft->init(dp3m.params);
147 dp3m.mesh.ks_pnum = dp3m.fft->get_ks_pnum();
148 dp3m.fft_buffers->init_meshes(dp3m.fft->get_ca_mesh_size());
149 dp3m.update_mesh_views();
150 dp3m.calc_differential_operator();
163 typename std::remove_reference_t<
decltype(dp3m)>::value_type;
164 auto const weights = p3m_calculate_interpolation_weights<cao>(
165 real_pos, dp3m.params.ai, dp3m.local_mesh);
166 p3m_interpolate<cao>(
167 dp3m.local_mesh, weights, [&dip, &dp3m](
int ind,
double w) {
168 dp3m.mesh.rs_fields[0u][ind] += value_type(w * dip[0u]);
169 dp3m.mesh.rs_fields[1u][ind] += value_type(w * dip[1u]);
170 dp3m.mesh.rs_fields[2u][ind] += value_type(w * dip[2u]);
173 dp3m.inter_weights.template store<cao>(weights);
178template <
typename FloatType, Arch Architecture>
181 dp3m.inter_weights.reset(dp3m.params.cao);
184 for (
auto &rs_mesh_field : dp3m.mesh.rs_fields)
185 for (
int j = 0; j < dp3m.local_mesh.size; j++)
186 rs_mesh_field[j] = 0.;
188 for (
auto const &p : particles) {
189 if (p.dipm() != 0.) {
190 Utils::integral_parameter<int, AssignDipole, 1, 7>(dp3m.params.cao, dp3m,
191 p.pos(), p.calc_dip());
202 auto p_index = std::size_t{0ul};
204 for (
auto &p : particles) {
205 if (p.dipm() != 0.) {
206 auto const weights = dp3m.inter_weights.template load<cao>(p_index);
210 [&E, &dp3m, d_rs](
int ind,
double w) {
211 E[d_rs] += w * double(dp3m.mesh.rs_scalar[ind]);
226 auto p_index = std::size_t{0ul};
228 for (
auto &p : particles) {
229 if (p.dipm() != 0.) {
230 auto const weights = dp3m.inter_weights.template load<cao>(p_index);
234 [&E, &dp3m](
int ind,
double w) {
235 E[0u] += w * double(dp3m.mesh.rs_fields[0u][ind]);
236 E[1u] += w * double(dp3m.mesh.rs_fields[1u][ind]);
237 E[2u] += w * double(dp3m.mesh.rs_fields[2u][ind]);
240 p.force()[d_rs] += p.calc_dip() * prefac * E;
248template <
typename FloatType, Arch Architecture>
250 bool force_flag,
bool energy_flag,
ParticleRange const &particles) {
253 auto const &system = get_system();
254 auto const &box_geo = *system.box_geo;
255 auto const dipole_prefac = prefactor / Utils::int_pow<3>(dp3m.params.mesh[0]);
257 auto const npt_flag =
260 auto constexpr npt_flag =
false;
263 if (dp3m.sum_mu2 > 0.) {
264 dipole_assign(particles);
265 dp3m.fft_buffers->perform_vector_halo_gather();
266 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
267 dp3m.fft->forward_fft(rs_mesh);
269 dp3m.update_mesh_views();
273 if (energy_flag or npt_flag) {
277 if (dp3m.sum_mu2 > 0.) {
282 auto const &offset = dp3m.mesh.start;
283 auto const &d_op = dp3m.d_op[0u];
284 auto const &mesh_dip = dp3m.mesh.rs_fields;
285 auto const [KX, KY, KZ] = dp3m.fft->get_permutations();
287 auto index = std::size_t(0u);
288 auto it_energy = dp3m.g_energy.begin();
289 auto node_energy = 0.;
290 for_each_3d(mesh_start, dp3m.mesh.size, indices, [&]() {
291 auto const shift = indices + offset;
293 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
294 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
295 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
298 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
299 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
300 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
302 node_energy += *it_energy * (Utils::sqr(re) + Utils::sqr(im));
303 std::advance(it_energy, 1);
306 node_energy *= dipole_prefac * std::numbers::pi * box_geo.length_inv()[0];
307 boost::mpi::reduce(
comm_cart, node_energy, energy, std::plus<>(), 0);
309 if (dp3m.energy_correction == 0.)
310 calc_energy_correction();
314 energy -= prefactor * dp3m.sum_mu2 * std::numbers::inv_sqrtpi *
315 (2. / 3.) * Utils::int_pow<3>(dp3m.params.alpha);
318 energy += prefactor * dp3m.energy_correction / box_geo.volume();
328 if (dp3m.sum_mu2 > 0.) {
329 auto const wavenumber = 2. * std::numbers::pi * box_geo.length_inv()[0u];
331 auto const &mesh_stop = dp3m.mesh.size;
332 auto const offset = dp3m.mesh.start;
333 auto const &d_op = dp3m.d_op[0u];
334 auto const [KX, KY, KZ] = dp3m.fft->get_permutations();
335 auto &mesh_dip = dp3m.mesh.rs_fields;
337 auto index = std::size_t(0u);
338 dp3m.ks_scalar.resize(dp3m.mesh.rs_scalar.size());
341 auto it_energy = dp3m.g_energy.begin();
343 for_each_3d(mesh_start, mesh_stop, indices, [&]() {
344 auto const shift = indices + offset;
346 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
347 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
348 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
349 dp3m.ks_scalar[index] = *it_energy * re;
352 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
353 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
354 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
355 dp3m.ks_scalar[index] = *it_energy * im;
357 std::advance(it_energy, 1);
361 for (
int d = 0; d < 3; d++) {
363 for_each_3d(mesh_start, mesh_stop, indices, [&]() {
364 auto const d_op_val = FloatType(d_op[indices[d] + offset[d]]);
365 dp3m.mesh.rs_scalar[index] = d_op_val * dp3m.ks_scalar[index];
367 dp3m.mesh.rs_scalar[index] = d_op_val * dp3m.ks_scalar[index];
370 dp3m.fft->backward_fft(dp3m.fft_buffers->get_scalar_mesh());
371 dp3m.fft_buffers->perform_scalar_halo_spread();
373 auto const d_rs = (d + dp3m.mesh.ks_pnum) % 3;
374 Utils::integral_parameter<int, AssignTorques, 1, 7>(
375 dp3m.params.cao, dp3m, dipole_prefac * wavenumber, d_rs, particles);
386 auto it_force = dp3m.g_force.begin();
388 for_each_3d(mesh_start, mesh_stop, indices, [&]() {
389 auto const shift = indices + offset;
391 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
392 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
393 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
396 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
397 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
398 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
400 dp3m.ks_scalar[index - 2] = *it_force * im;
401 dp3m.ks_scalar[index - 1] = *it_force * (-re);
402 std::advance(it_force, 1);
406 for (
int d = 0; d < 3; d++) {
408 for_each_3d(mesh_start, mesh_stop, indices, [&]() {
409 auto const d_op_val = FloatType(d_op[indices[d] + offset[d]]);
410 auto const shift = indices + offset;
411 auto const f1 = d_op_val * dp3m.ks_scalar[index];
412 mesh_dip[0u][index] = FloatType(d_op[shift[KX]]) * f1;
413 mesh_dip[1u][index] = FloatType(d_op[shift[KY]]) * f1;
414 mesh_dip[2u][index] = FloatType(d_op[shift[KZ]]) * f1;
416 auto const f2 = d_op_val * dp3m.ks_scalar[index];
417 mesh_dip[0u][index] = FloatType(d_op[shift[KX]]) * f2;
418 mesh_dip[1u][index] = FloatType(d_op[shift[KY]]) * f2;
419 mesh_dip[2u][index] = FloatType(d_op[shift[KZ]]) * f2;
422 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
423 dp3m.fft->backward_fft(rs_mesh);
425 dp3m.fft_buffers->perform_vector_halo_spread();
427 auto const d_rs = (d + dp3m.mesh.ks_pnum) % 3;
428 Utils::integral_parameter<int, AssignForces, 1, 7>(
429 dp3m.params.cao, dp3m, dipole_prefac *
Utils::sqr(wavenumber), d_rs,
436 auto const surface_term =
437 calc_surface_term(force_flag, energy_flag or npt_flag, particles);
439 energy += surface_term;
444 fprintf(stderr,
"dipolar_P3M at this moment is added to p_vir[0]\n");
446 if (not energy_flag) {
453template <
typename FloatType, Arch Architecture>
455 bool force_flag,
bool energy_flag,
ParticleRange const &particles) {
456 auto const &box_geo = *get_system().
box_geo;
457 auto const pref = prefactor * 4. * std::numbers::pi / box_geo.volume() /
458 (2. * dp3m.params.epsilon + 1.);
459 auto const n_local_part = particles.
size();
463 std::vector<double> mx(n_local_part);
464 std::vector<double> my(n_local_part);
465 std::vector<double> mz(n_local_part);
468 for (
auto const &p : particles) {
469 auto const dip = p.calc_dip();
478 for (std::size_t i = 0u; i < n_local_part; i++) {
479 local_dip[0u] += mx[i];
480 local_dip[1u] += my[i];
481 local_dip[2u] += mz[i];
484 boost::mpi::all_reduce(
comm_cart, local_dip, std::plus<>());
489 for (std::size_t i = 0u; i < n_local_part; i++) {
490 sum_e += mx[i] * box_dip[0] + my[i] * box_dip[1] + mz[i] * box_dip[2];
493 0.5 * pref * boost::mpi::all_reduce(
comm_cart, sum_e, std::plus<>());
498 std::vector<double> sumix(n_local_part);
499 std::vector<double> sumiy(n_local_part);
500 std::vector<double> sumiz(n_local_part);
502 for (std::size_t i = 0u; i < n_local_part; i++) {
503 sumix[i] = my[i] * box_dip[2u] - mz[i] * box_dip[1u];
504 sumiy[i] = mz[i] * box_dip[0u] - mx[i] * box_dip[2u];
505 sumiz[i] = mx[i] * box_dip[1u] - my[i] * box_dip[0u];
509 for (
auto &p : particles) {
510 auto &torque = p.torque();
511 torque[0u] -= pref * sumix[ip];
512 torque[1u] -= pref * sumiy[ip];
513 torque[2u] -= pref * sumiz[ip];
521template <
typename FloatType, Arch Architecture>
523 dp3m.g_force = grid_influence_function<FloatType, 3, P3M_BRILLOUIN>(
524 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop,
525 get_system().box_geo->length_inv());
528template <
typename FloatType, Arch Architecture>
530 dp3m.g_energy = grid_influence_function<FloatType, 2, P3M_BRILLOUIN>(
531 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop,
532 get_system().box_geo->length_inv());
535template <
typename FloatType, Arch Architecture>
538 int m_mesh_max = -1, m_mesh_min = -1;
539 std::pair<std::optional<int>, std::optional<int>> m_tune_limits;
543 double prefactor,
int timings,
544 decltype(m_tune_limits) tune_limits)
546 m_tune_limits{std::move(tune_limits)} {}
552 std::optional<std::string>
559 m_logger = std::make_unique<TuningLogger>(
567 std::tuple<double, double, double, double>
569 double r_cut_iL)
const override {
571 double alpha_L, rs_err, ks_err;
583 0.0001 * box_geo.length()[0], 5. * box_geo.length()[0], 0.0001,
606 auto const expo = std::log(std::cbrt(dp3m.
sum_dip_part)) / std::log(2.);
608 m_mesh_min =
static_cast<int>(std::round(std::pow(2., std::floor(expo))));
611 if (m_tune_limits.first) {
612 m_mesh_min = *m_tune_limits.first;
614 if (m_tune_limits.second) {
615 m_mesh_max = *m_tune_limits.second;
618 m_mesh_min = m_mesh_max = dp3m.
params.
mesh[0];
626 for (
auto tmp_mesh = m_mesh_min; tmp_mesh <= m_mesh_max; tmp_mesh += 2) {
631 auto const trial_time =
632 get_m_time(trial_params.mesh, trial_params.cao, trial_params.r_cut_iL,
633 trial_params.alpha_L, trial_params.accuracy);
643 if (trial_time < time_best) {
646 tuned_params = trial_params;
647 time_best = tuned_params.time = trial_time;
658template <
typename FloatType, Arch Architecture>
660 auto &system = get_system();
661 auto const &box_geo = *system.
box_geo;
668 if (not is_tuned()) {
671 throw std::runtime_error(
672 "DipolarP3M: no dipolar particles in the system");
676 system, dp3m, prefactor, tune_timings, tune_limits);
696 double mesh_i,
int cao,
double alpha_L_i) {
700 auto const factor1 =
Utils::sqr(std::numbers::pi * alpha_L_i);
708 mesh_start, mesh_stop, indices,
710 auto const norm_sq = nm.norm2();
711 auto const ex = std::exp(-factor1 * norm_sq);
714 alias2 += U2 * ex * std::pow(shift * nm, 3) / norm_sq;
716 [&](
unsigned dim,
int n) {
717 nm[dim] = shift[dim] + n * mesh;
721 return std::make_pair(alias1, alias2);
726 int n_c_part,
double sum_q2,
double alpha_L) {
729 auto const mesh_i = 1. /
static_cast<double>(mesh);
730 auto const alpha_L_i = 1. / alpha_L;
732 auto const mesh_start = -mesh_stop;
738 mesh_start, mesh_stop, indices,
740 if ((indices[0] != 0) or (indices[1] != 0) or (indices[2] != 0)) {
741 auto const n2 = indices.norm2();
743 auto const [alias1, alias2] =
747 Utils::int_pow<3>(
static_cast<double>(n2));
754 [&values, &mesh_i, cotangent_sum](
unsigned dim,
int n) {
755 values[dim] = cotangent_sum(n, mesh_i);
758 return 8. *
Utils::sqr(std::numbers::pi) / 3. * sum_q2 *
759 sqrt(he_q / n_c_part) / Utils::int_pow<4>(box_size);
769 int n_c_part,
double sum_q2,
771 double d_error_f, d_cc, d_dc, d_con;
773 auto const d_rcut = r_cut_iL * box_size;
779 auto const d_c = sum_q2 * exp(-d_a2 * d_rcut2);
783 d_dc = 8. * Utils::int_pow<3>(d_a2) * Utils::int_pow<3>(d_rcut2) +
784 20. *
Utils::sqr(d_a2) * d_rcut4 + 30. * d_a2 * d_rcut2 + 15.;
786 d_con = 1. / sqrt(Utils::int_pow<3>(box_size) *
Utils::sqr(d_a2) * d_rcut *
787 Utils::sqr(d_rcut4) *
static_cast<double>(n_c_part));
789 d_error_f = d_c * d_con *
791 (2. / 15.) *
Utils::sqr(d_dc) - (13. / 15.) * d_cc * d_dc);
801 double sum_q2,
double x1,
double x2,
double xacc,
802 double tuned_accuracy) {
803 constexpr int JJ_RTBIS_MAX = 40;
805 auto const constant = tuned_accuracy / std::numbers::sqrt2;
813 if (f1 * f2 >= 0.0) {
814 throw std::runtime_error(
815 "Root must be bracketed for bisection in dp3m_rtbisection");
819 double rtb = f1 < 0.0 ? (dx = x2 - x1, x1) : (dx = x1 - x2, x2);
820 for (
int j = 1; j <= JJ_RTBIS_MAX; j++) {
821 auto const xmid = rtb + (dx *= 0.5);
827 if (fabs(dx) < xacc || fmid == 0.0)
830 throw std::runtime_error(
"Too many bisections in dp3m_rtbisection");
835 auto const &box_geo = *system.
box_geo;
836 auto const &local_geo = *system.
local_geo;
837 for (
auto i = 0u; i < 3u; i++) {
840 std::stringstream msg;
842 <<
" is larger than half of box dimension " << box_geo.length()[i];
843 throw std::runtime_error(msg.str());
846 std::stringstream msg;
848 <<
" is larger than local box dimension " << local_geo.length()[i];
849 throw std::runtime_error(msg.str());
853 if ((box_geo.length()[0] != box_geo.length()[1]) or
854 (box_geo.length()[1] != box_geo.length()[2])) {
855 throw std::runtime_error(
"DipolarP3M: requires a cubic box");
861 if (!box_geo.periodic(0) or !box_geo.periodic(1) or !box_geo.periodic(2)) {
862 throw std::runtime_error(
863 "DipolarP3M: requires periodicity (True, True, True)");
868 auto const &local_geo = *
get_system().local_geo;
871 throw std::runtime_error(
872 "DipolarP3M: requires the regular or hybrid decomposition cell system");
876 throw std::runtime_error(
877 "DipolarP3M: does not work with the hybrid decomposition cell system, "
878 "if using more than one MPI node");
884 if (node_grid[0] < node_grid[1] or node_grid[1] < node_grid[2]) {
885 throw std::runtime_error(
886 "DipolarP3M: node grid must be sorted, largest first");
890template <
typename FloatType, Arch Architecture>
892 auto const &box_geo = *get_system().
box_geo;
897 sanity_checks_boxl();
898 calc_influence_function_force();
899 calc_influence_function_energy();
903template <
typename FloatType, Arch Architecture>
905 auto const &box_geo = *get_system().
box_geo;
906 auto const Ukp3m = calc_average_self_energy_k_space() * box_geo.volume();
907 auto const Ewald_volume = Utils::int_pow<3>(dp3m.
params.
alpha_L);
908 auto const Eself = -2. * Ewald_volume * std::numbers::inv_sqrtpi / 3.;
910 -dp3m.
sum_mu2 * (Ukp3m + Eself + 2. * std::numbers::pi / 3.);
@ HYBRID
Hybrid decomposition.
@ REGULAR
Regular decomposition.
Vector implementation and trait types for boost qvm interoperability.
void on_solver_change() const override
void setup_logger(bool verbose) override
TuningAlgorithm::Parameters get_time() override
P3MParameters & get_params() override
DipolarTuningAlgorithm(System::System &system, decltype(dp3m) &input_dp3m, double prefactor, int timings, decltype(m_tune_limits) tune_limits)
std::tuple< double, double, double, double > calculate_accuracy(Utils::Vector3i const &mesh, int cao, double r_cut_iL) const override
std::optional< std::string > layer_correction_veto_r_cut(double) const override
void determine_mesh_limits() override
base_type::size_type size() const
std::shared_ptr< LocalBox > local_geo
std::shared_ptr< CellStructure > cell_structure
std::shared_ptr< BoxGeometry > box_geo
Tuning algorithm for P3M.
double get_m_time(Utils::Vector3i const &mesh, int &tuned_cao, double &tuned_r_cut_iL, double &tuned_alpha_L, double &tuned_accuracy)
Get the optimal alpha and the corresponding computation time for a fixed mesh.
static auto constexpr time_sentinel
Value for invalid time measurements.
static auto constexpr max_n_consecutive_trials
Maximal number of consecutive trials that don't improve runtime.
System::System & m_system
void determine_cao_limits(int initial_cao)
Determine a sensible range for the charge assignment order.
void determine_r_cut_limits()
Determine a sensible range for the real-space cutoff.
std::unique_ptr< TuningLogger > m_logger
static auto constexpr time_granularity
Granularity of the time measurement (milliseconds).
static DEVICE_QUALIFIER constexpr Vector< T, N > broadcast(typename Base::value_type const &value)
Create a vector that has all entries set to the same value.
Communicator communicator
boost::mpi::communicator comm_cart
The communicator.
int this_node
The number of this node.
This file contains the defaults for ESPResSo.
#define ROUND_ERROR_PREC
Precision for capture of round off errors.
#define P3M_BRILLOUIN
P3M: Number of Brillouin zones taken into account in the calculation of the optimal influence functio...
__device__ void vector_product(float const *a, float const *b, float *out)
static int count_magnetic_particles(System::System const &system)
static auto dp3m_tune_aliasing_sums(Utils::Vector3i const &shift, int mesh, double mesh_i, int cao, double alpha_L_i)
Tuning dipolar-P3M.
static double dp3m_k_space_error(double box_size, int mesh, int cao, int n_c_part, double sum_q2, double alpha_L)
Calculate the k-space error of dipolar-P3M.
void npt_add_virial_magnetic_contribution(double energy)
Update the NpT virial.
double dp3m_rtbisection(double box_size, double r_cut_iL, int n_c_part, double sum_q2, double x1, double x2, double xacc, double tuned_accuracy)
Compute the value of alpha through a bisection method.
static double dp3m_real_space_error(double box_size, double r_cut_iL, int n_c_part, double sum_q2, double alpha_L)
Calculate the value of the errors for the REAL part of the force in terms of the splitting parameter ...
P3M algorithm for long-range magnetic dipole-dipole interaction.
This file contains the errorhandling code for severe errors, like a broken bond or illegal parameter ...
Routines, row decomposition, data structures and communication for the 3D-FFT.
and std::invocable< Projector, unsigned, int > void for_each_3d(detail::IndexVectorConcept auto &&start, detail::IndexVectorConcept auto &&stop, detail::IndexVectorConcept auto &&counters, Kernel &&kernel, Projector &&projector=detail::noop_projector)
Repeat an operation on every element of a 3D grid.
void p3m_interpolate(P3MLocalMesh const &local_mesh, InterpolationWeights< cao > const &weights, Kernel kernel)
P3M grid interpolation.
T product(Vector< T, N > const &v)
DEVICE_QUALIFIER constexpr T sqr(T x)
Calculates the SQuaRe of x.
DEVICE_QUALIFIER auto sinc(T x)
Calculate the function .
auto get_analytic_cotangent_sum_kernel(int cao)
void npt_add_virial_contribution(double energy)
Exports for the NpT code.
Common functions for dipolar and charge P3M.
auto constexpr P3M_EPSILON_METALLIC
This value indicates metallic boundary conditions.
Utils::Vector3i node_grid
void calc_energy_correction() override
double long_range_kernel(bool force_flag, bool energy_flag, ParticleRange const &particles)
Compute the k-space part of forces and energies.
double calc_surface_term(bool force_flag, bool energy_flag, ParticleRange const &particles) override
void dipole_assign(ParticleRange const &particles) override
void calc_influence_function_energy() override
void count_magnetic_particles() override
void calc_influence_function_force() override
void scaleby_box_l() override
double calc_average_self_energy_k_space() const override
void sanity_checks_boxl() const
Checks for correctness of the k-space cutoff.
void sanity_checks_cell_structure() const
P3MParameters const & dp3m_params
void sanity_checks_periodicity() const
void sanity_checks_node_grid() const
void recalc_ld_pos(P3MParameters const ¶ms)
Recalculate quantities derived from the mesh and box length: ld_pos (position of the left down mesh).
Structure to hold P3M parameters and some dependent variables.
Utils::Vector3d cao_cut
cutoff for charge assignment.
double alpha
unscaled alpha_L for use with fast inline functions only
double r_cut_iL
cutoff radius for real space electrostatics (>0), rescaled to r_cut_iL = r_cut * box_l_i.
double accuracy
accuracy of the actual parameter set.
double alpha_L
Ewald splitting parameter (0.
double r_cut
unscaled r_cut_iL for use with fast inline functions only
void recalc_a_ai_cao_cut(Utils::Vector3d const &box_l)
Recalculate quantities derived from the mesh and box length: a, ai and cao_cut.
bool tuning
tuning or production?
Utils::Vector3i mesh
number of mesh points per coordinate direction (>0), in real space.
void operator()(auto &dp3m, Utils::Vector3d const &real_pos, Utils::Vector3d const &dip) const
void operator()(auto &dp3m, double prefac, int d_rs, ParticleRange const &particles) const
void operator()(auto &dp3m, double prefac, int d_rs, ParticleRange const &particles) const
double energy_correction
cached k-space self-energy correction
int sum_dip_part
number of dipolar particles (only on head node).
double sum_mu2
Sum of square of magnetic dipoles (only on head node).
P3MLocalMesh local_mesh
Local mesh properties.
P3MParameters params
P3M base parameters.