49#include "communication.hpp"
54#include "system/System.hpp"
62#include <boost/mpi/collectives/all_reduce.hpp>
63#include <boost/mpi/collectives/reduce.hpp>
65#include <Kokkos_Core.hpp>
66#include <Kokkos_ScatterView.hpp>
85#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
89 auto const diff = std::abs(value - reference);
90 using FT = std::remove_cvref_t<
decltype(
diff)>;
91 auto constexpr atol = std::is_same_v<FT, float> ?
FT{2
e-4} :
FT{1
e-6};
92 auto constexpr rtol = std::is_same_v<FT, float> ?
FT{5
e-5} :
FT{1
e-5};
93 auto const non_zero = std::abs(reference) !=
FT{0};
99template <
typename FloatType, Arch Architecture,
class FFTConfig>
105 for (
auto const &p : get_system().
cell_structure->local_particles()) {
106 if (p.dipm() != 0.) {
113 boost::mpi::all_reduce(
comm_cart,
local_n, dp3m.sum_dip_part, std::plus<>());
117 std::size_t
n_c_part,
double sum_q2,
121 std::size_t
n_c_part,
double sum_q2,
128 double sum_q2,
double x1,
double x2,
double xacc,
131template <
typename FloatType, Arch Architecture,
class FFTConfig>
134 auto const &box_geo = *get_system().
box_geo;
137 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop, dp3m.g_energy);
141 phi /= 3. * box_geo.length()[0] *
142 Utils::int_pow<3>(
static_cast<double>(dp3m.params.mesh[0]));
143 return phi * std::numbers::pi;
146template <
typename FloatType, Arch Architecture,
class FFTConfig>
150 assert(dp3m.params.alpha > 0.);
152 auto const &
system = get_system();
153 auto const &box_geo = *
system.box_geo;
154 auto const &local_geo = *
system.local_geo;
157 dp3m.params.cao3 = Utils::int_pow<3>(dp3m.params.cao);
158 dp3m.params.recalc_a_ai_cao_cut(box_geo.length());
161 dp3m.local_mesh.calc_local_ca_mesh(dp3m.params, local_geo,
verlet_skin, 0.);
162 dp3m.fft_buffers->init_halo();
163 dp3m.fft->init(dp3m.params);
164 dp3m.mesh.ks_pnum = dp3m.fft->get_ks_pnum();
165 dp3m.fft_buffers->init_meshes(dp3m.fft->get_ca_mesh_size());
166 dp3m.update_mesh_views();
167#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
168 dp3m.heffte.world_size =
comm_cart.size();
170 std::make_shared<P3MFFT<FloatType, Architecture, FFTConfig>>(
171 nullptr,
::comm_cart, dp3m.params.mesh, dp3m.local_mesh.ld_no_halo,
173 dp3m.resize_heffte_buffers();
175 dp3m.calc_differential_operator();
188 using execution_space = Kokkos::DefaultHostExecutionSpace;
189 auto const &aosoa = cell_structure.get_aosoa();
190 auto const &unique_particles = cell_structure.get_unique_particles();
191 auto const n_part = cell_structure.count_local_particles();
194 "InterpolateDipoles", std::size_t{0
u}, n_part, [&](
auto p_index) {
197 auto const p_pos = aosoa.get_span_at(aosoa.position,
p_index);
198 auto const dip = unique_particles.at(
p_index)->calc_dip();
201 p_pos, dp3m.params.ai, dp3m.local_mesh);
202 dp3m.inter_weights.store_at(
p_index, weights);
204 dp3m.local_mesh, weights, [&dip,
tid, &dp3m](
int ind,
double w) {
205 dp3m.rs_fields_kokkos(tid, 0u, ind) += value_type(w * dip[0u]);
206 dp3m.rs_fields_kokkos(tid, 1u, ind) += value_type(w * dip[1u]);
207 dp3m.rs_fields_kokkos(tid, 2u, ind) += value_type(w * dip[2u]);
213 "ReduceInterpolatedDipoles", std::size_t{0}, dp3m.local_mesh.size,
215 for (
int dir = 0; dir < 3; ++dir) {
218 acc += dp3m.rs_fields_kokkos(
tid, dir, i);
220 dp3m.mesh.rs_fields[dir][i] += acc;
221#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
222 dp3m.heffte.rs_dipole_density[dir][i] += acc;
231template <
typename FloatType, Arch Architecture,
class FFTConfig>
235 Utils::integral_parameter<int, AssignDipole, p3m_min_cao, p3m_max_cao>(
236 dp3m.params.cao, dp3m, *get_system().cell_structure);
244 assert(cao == dp3m.inter_weights.cao());
245 using execution_space = Kokkos::DefaultHostExecutionSpace;
247 auto const kernel = [
d_rs, pref, &dp3m](
auto const &dip,
248#ifdef ESPRESSO_DIPOLE_FIELD_TRACKING
256 [&
E, &dp3m,
d_rs](
int ind,
double w) {
258 E[d_rs] += w * double(dp3m.mesh.rs_scalar[ind]);
263 access(
p_index, 0) -= torque[0];
264 access(
p_index, 1) -= torque[1];
265 access(
p_index, 2) -= torque[2];
266#ifdef ESPRESSO_DIPOLE_FIELD_TRACKING
267 auto const dip_fld = pref *
E;
275 auto const n_part = dp3m.inter_weights.size();
279#ifdef ESPRESSO_DIPOLE_FIELD_TRACKING
283 "AssignTorques", std::size_t{0
u}, n_part, [&](std::size_t
p_index) {
284 auto const &p = *unique_particles.at(
p_index);
285 if (p.dipm() != 0.) {
300 assert(cao == dp3m.inter_weights.cao());
301 using execution_space = Kokkos::DefaultHostExecutionSpace;
303 auto const kernel = [
d_rs, pref, &dp3m](
auto const &dip,
auto &
p_force,
310 E[0u] += w * double(dp3m.mesh.rs_fields[0u][ind]);
311 E[1u] += w * double(dp3m.mesh.rs_fields[1u][ind]);
312 E[2u] += w * double(dp3m.mesh.rs_fields[2u][ind]);
315 auto access =
p_force.access();
319 auto const n_part = dp3m.inter_weights.size();
323 "AssignForcesDip", std::size_t{0
u}, n_part, [&](std::size_t
p_index) {
324 auto const &p = *unique_particles.at(
p_index);
325 if (p.dipm() != 0.) {
333#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
334template <
typename FloatType,
class FFTConfig>
339 static_cast<std::size_t
>(
Utils::product(this->local_mesh.dim_no_halo));
341 static_cast<std::size_t
>(
Utils::product(heffte.fft->ks_local_size()));
342 for (
auto d : {0
u, 1u, 2u}) {
390template <
typename FloatType, Arch Architecture,
class FFTConfig>
393 auto const &
system = get_system();
394 auto const &box_geo = *
system.box_geo;
398 if (dp3m.sum_mu2 > 0.) {
400 dp3m.fft_buffers->perform_vector_halo_gather();
401 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
402 dp3m.fft->forward_fft(rs_mesh);
404 dp3m.update_mesh_views();
408 auto const wavevector = 2. * std::numbers::pi * box_geo.length_inv()[0];
411 auto index = std::size_t(0
u);
414 auto constexpr KX = 2, KY = 0, KZ = 1;
415 auto const shift = local_index + dp3m.mesh.start;
416 auto const &d_op = dp3m.d_op[0u];
417 auto const &mesh_dip = dp3m.mesh.rs_fields;
418 auto const d_op_x = static_cast<FloatType>(d_op[shift[KX]]);
419 auto const d_op_y = static_cast<FloatType>(d_op[shift[KY]]);
420 auto const d_op_z = static_cast<FloatType>(d_op[shift[KZ]]);
423 auto const Mx_re = mesh_dip[0u][index];
424 auto const My_re = mesh_dip[1u][index];
425 auto const Mz_re = mesh_dip[2u][index];
426 auto const Q_re = Mx_re * d_op_x + My_re * d_op_y + Mz_re * d_op_z;
429 auto const Mx_im = mesh_dip[0u][index];
430 auto const My_im = mesh_dip[1u][index];
431 auto const Mz_im = mesh_dip[2u][index];
432 auto const Q_im = Mx_im * d_op_x + My_im * d_op_y + Mz_im * d_op_z;
435 auto const nx = static_cast<double>(d_op[shift[KX]]);
436 auto const ny = static_cast<double>(d_op[shift[KY]]);
437 auto const nz = static_cast<double>(d_op[shift[KZ]]);
438 auto const kx = nx * wavevector;
439 auto const ky = ny * wavevector;
440 auto const kz = nz * wavevector;
441 auto const norm_sq = Utils::sqr(kx) + Utils::sqr(ky) + Utils::sqr(kz);
443 auto const g = static_cast<double>(*it_energy);
444 auto const cell_energy =
445 g * static_cast<double>(Utils::sqr(Q_re) + Utils::sqr(Q_im));
446 auto const vterm = -2. * (1. / norm_sq + half_alpha_inv_sq);
449 auto const Rx = g * static_cast<double>(Mx_re * Q_re + Mx_im * Q_im);
450 auto const Ry = g * static_cast<double>(My_re * Q_re + My_im * Q_im);
451 auto const Rz = g * static_cast<double>(Mz_re * Q_re + Mz_im * Q_im);
462 node_k_space_pressure_tensor[0u] +=
463 cell_energy * (1. + vterm * kx * kx) + 2. * nx * Rx;
464 node_k_space_pressure_tensor[1u] +=
465 cell_energy * vterm * kx * ky + 2. * ny * Rx;
466 node_k_space_pressure_tensor[2u] +=
467 cell_energy * vterm * kx * kz + 2. * nz * Rx;
468 node_k_space_pressure_tensor[3u] +=
469 cell_energy * vterm * ky * kx + 2. * nx * Ry;
470 node_k_space_pressure_tensor[4u] +=
471 cell_energy * (1. + vterm * ky * ky) + 2. * ny * Ry;
472 node_k_space_pressure_tensor[5u] +=
473 cell_energy * vterm * ky * kz + 2. * nz * Ry;
474 node_k_space_pressure_tensor[6u] +=
475 cell_energy * vterm * kz * kx + 2. * nx * Rz;
476 node_k_space_pressure_tensor[7u] +=
477 cell_energy * vterm * kz * ky + 2. * ny * Rz;
478 node_k_space_pressure_tensor[8u] +=
479 cell_energy * (1. + vterm * kz * kz) + 2. * nz * Rz;
486 box_geo.length_inv()[0];
489template <
typename FloatType, Arch Architecture,
class FFTConfig>
494 auto const &
system = get_system();
495 auto const &box_geo = *
system.box_geo;
500#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
501 auto constexpr r2c_dir = FFTConfig::r2c_dir;
502 auto const rs_local_size = dp3m.heffte.fft->rs_local_size();
503 auto const local_size = dp3m.heffte.fft->ks_local_size();
505 if constexpr (FFTConfig::use_r2c) {
509 auto const local_origin = dp3m.heffte.fft->ks_local_ld_index();
518 dp3m.resize_heffte_buffers();
521 if (dp3m.sum_mu2 > 0.) {
523 dp3m.fft_buffers->perform_vector_halo_gather();
524 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
525 dp3m.fft->forward_fft(rs_mesh);
527 dp3m.update_mesh_views();
529#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
530 if (dp3m.heffte.world_size == 1) {
532 std::array<FloatType *, 3u> rs_fields = {
533 {dp3m.heffte.rs_dipole_density[0
u].data(),
534 dp3m.heffte.rs_dipole_density[1u].data(),
535 dp3m.heffte.rs_dipole_density[2u].data()}};
536 dp3m.heffte.halo_comm.gather_grid(
::comm_cart, rs_fields,
537 dp3m.local_mesh.dim);
539 for (
auto dir : {0
u, 1u, 2u}) {
542 FFTConfig::r_space_order>(
543 dp3m.rs_field_no_halo_kokkos.data(),
544 dp3m.heffte.rs_dipole_density[dir], dp3m.local_mesh.dim,
545 dp3m.local_mesh.n_halo_ld,
546 dp3m.local_mesh.dim - dp3m.local_mesh.n_halo_ur);
551 auto constexpr KX = 1,
KY = 2,
KZ = 0;
556 dp3m.rs_field_no_halo_kokkos(index);
559 dp3m.heffte.fft->forward(dp3m.rs_field_no_halo_reorder_kokkos.data(),
560 dp3m.heffte.ks_dipole_density[dir].data());
562 if (
not dp3m.params.tuning) {
567 auto constexpr KX = 2,
KY = 0,
KZ = 1;
571 auto const old_value = std::complex<FloatType>{
592 if (dp3m.sum_mu2 > 0.) {
596 auto index = std::size_t(0
u);
600 auto constexpr KX = 2, KY = 0, KZ = 1;
601 auto const shift = local_index + dp3m.mesh.start;
602 auto const &d_op = dp3m.d_op[0u];
603 auto const &mesh_dip = dp3m.mesh.rs_fields;
605 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
606 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
607 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
610 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
611 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
612 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
614 node_energy += *it_energy * (Utils::sqr(re) + Utils::sqr(im));
615 std::advance(it_energy, 1);
617#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
618 if (dp3m.heffte.world_size == 1) {
625 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
656 if (dp3m.energy_correction == 0.)
657 calc_energy_correction();
661 energy -= prefactor * dp3m.sum_mu2 * std::numbers::inv_sqrtpi *
662 (2. / 3.) * Utils::int_pow<3>(dp3m.params.alpha);
665 energy += prefactor * dp3m.energy_correction / box_geo.volume();
675 if (dp3m.sum_mu2 > 0.) {
676 auto const wavenumber = 2. * std::numbers::pi * box_geo.length_inv()[0
u];
677 dp3m.ks_scalar.resize(dp3m.local_mesh.size);
680 auto index{std::size_t(0
u)};
684 auto constexpr KX = 2, KY = 0, KZ = 1;
685 auto const shift = local_index + dp3m.mesh.start;
686 auto const &d_op = dp3m.d_op[0u];
687 auto const &mesh_dip = dp3m.mesh.rs_fields;
689 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
690 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
691 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
694 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
695 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
696 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
698 *it_ks_scalar = *it_energy * std::complex<FloatType>{re, im};
703#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
704 if (dp3m.heffte.world_size == 1) {
710 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
720 if (
not dp3m.params.tuning) {
721 auto constexpr KX = 2,
KY = 0,
KZ = 1;
737 for (
int d = 0; d < 3; d++) {
741 auto const &offset = dp3m.mesh.start;
742 auto const &d_op = dp3m.d_op[0u];
743 auto const d_op_val = FloatType(d_op[local_index[d] + offset[d]]);
744 auto const &value = *it_ks_scalar;
745 dp3m.mesh.rs_scalar[index] = d_op_val * value.real();
747 dp3m.mesh.rs_scalar[index] = d_op_val * value.imag();
749 std::advance(it_ks_scalar, 1);
751#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
752 if (dp3m.heffte.world_size == 1) {
753 unsigned int constexpr d_ks[3] = {2u, 0
u, 1u};
764 if (
not dp3m.params.tuning) {
765 auto constexpr KX = 2,
KY = 0,
KZ = 1;
769 auto const old_value = std::complex<FloatType>{
780 dp3m.heffte.fft->backward(dp3m.heffte.ks_B_field_storage.data(),
781 dp3m.heffte.rs_B_fields_no_halo[d].data());
786 dp3m.heffte.rs_B_fields[d].data(),
787 std::span(dp3m.heffte.rs_B_fields_no_halo[d]),
788 dp3m.local_mesh.dim_no_halo, dp3m.local_mesh.n_halo_ld,
789 dp3m.local_mesh.n_halo_ur);
792 dp3m.heffte.rs_B_fields[d].data(),
793 dp3m.local_mesh.dim);
796 dp3m.fft->backward_fft(dp3m.fft_buffers->get_scalar_mesh());
798 dp3m.fft_buffers->perform_scalar_halo_spread();
800 auto const d_rs = (d + dp3m.mesh.ks_pnum) % 3;
801 Utils::integral_parameter<int, AssignTorques, p3m_min_cao, p3m_max_cao>(
815 auto it_force = dp3m.g_force.begin();
817 std::size_t index = 0
u;
819 auto constexpr KX = 2, KY = 0, KZ = 1;
820 auto const shift = local_index + dp3m.mesh.start;
821 auto const &d_op = dp3m.d_op[0u];
822 auto const &mesh_dip = dp3m.mesh.rs_fields;
824 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
825 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
826 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
829 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
830 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
831 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
833 *it_ks_scalar = {*it_force * im, *it_force * (-re)};
839#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
840 if (dp3m.heffte.world_size == 1) {
846 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
858 if (
not dp3m.params.tuning) {
859 auto constexpr KX = 2,
KY = 0,
KZ = 1;
875 for (
int d = 0; d < 3; d++) {
876 std::size_t index = 0
u;
879 auto constexpr KX = 2, KY = 0, KZ = 1;
880 auto const shift = local_index + dp3m.mesh.start;
881 auto const &d_op = dp3m.d_op[0u];
882 auto const &mesh_dip = dp3m.mesh.rs_fields;
883 auto const d_op_val = FloatType(d_op[shift[d]]);
884 auto const f = *it_ks_scalar * d_op_val;
885 mesh_dip[0u][index] = FloatType(d_op[shift[KX]]) * f.real();
886 mesh_dip[1u][index] = FloatType(d_op[shift[KY]]) * f.real();
887 mesh_dip[2u][index] = FloatType(d_op[shift[KZ]]) * f.real();
889 mesh_dip[0u][index] = FloatType(d_op[shift[KX]]) * f.imag();
890 mesh_dip[1u][index] = FloatType(d_op[shift[KY]]) * f.imag();
891 mesh_dip[2u][index] = FloatType(d_op[shift[KZ]]) * f.imag();
893 std::advance(it_ks_scalar, 1);
896#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
897 if (dp3m.heffte.world_size == 1) {
902 auto constexpr KX = 1,
KY = 2,
KZ = 0;
909 auto &
mesh_dip = dp3m.heffte.ks_dipole_density;
920 if (
not FFTConfig::use_r2c
and not dp3m.params.tuning) {
924 for (
int j = 0;
j < 3; ++
j) {
925 auto const old_value = std::complex<FloatType>{
936 for (
int dir = 0
u; dir < 3u; ++dir) {
937 dp3m.heffte.fft->backward(
938 dp3m.heffte.ks_dipole_density[dir].data(),
939 dp3m.heffte.rs_B_fields_no_halo[dir].data());
944 dp3m.heffte.rs_B_fields[d].data(),
945 std::span(dp3m.heffte.rs_B_fields_no_halo[dir]),
946 dp3m.local_mesh.dim_no_halo, dp3m.local_mesh.n_halo_ld,
947 dp3m.local_mesh.n_halo_ur);
951 std::array<FloatType *, 3u>{{dp3m.heffte.rs_B_fields[0
u].data(),
952 dp3m.heffte.rs_B_fields[1u].data(),
953 dp3m.heffte.rs_B_fields[2u].data()}};
954 dp3m.heffte.halo_comm.spread_grid(
::comm_cart, rs_fields,
955 dp3m.local_mesh.dim);
958 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
959 dp3m.fft->backward_fft(rs_mesh);
962 dp3m.fft_buffers->perform_vector_halo_spread();
964 auto const d_rs = (d + dp3m.mesh.ks_pnum) % 3;
998template <
typename FloatType, Arch Architecture,
class FFTConfig>
1001 auto const &
system = get_system();
1002 auto const &box_geo = *
system.box_geo;
1003 auto const particles =
system.cell_structure->local_particles();
1004 auto const pref = prefactor * 4. * std::numbers::pi / box_geo.volume() /
1005 (2. * dp3m.params.epsilon + 1.);
1014 std::size_t
ip = 0
u;
1015 for (
auto const &p : particles) {
1016 auto const dip = p.calc_dip();
1040 0.5 * pref * boost::mpi::all_reduce(
comm_cart,
sum_e, std::plus<>());
1056 for (
auto &p : particles) {
1057 auto &torque = p.torque();
1059 torque[1u] -= pref *
sumiy[
ip];
1060 torque[2u] -= pref *
sumiz[
ip];
1061#ifdef ESPRESSO_DIPOLE_FIELD_TRACKING
1062 p.dip_fld() -= pref *
box_dip;
1071template <
typename FloatType, Arch Architecture,
class FFTConfig>
1075 FFTConfig::k_space_order>(
1076 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop,
1077 get_system().
box_geo->length_inv());
1078#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
1079 if (dp3m.heffte.world_size == 1) {
1080 dp3m.heffte.g_force =
1082 FFTConfig::k_space_order>(
1083 dp3m.params, dp3m.heffte.fft->ks_local_ld_index(),
1084 dp3m.heffte.fft->ks_local_ur_index(),
1085 get_system().
box_geo->length_inv());
1086 if constexpr (FFTConfig::use_r2c) {
1088 dp3m.heffte.g_force, dp3m.params.mesh,
1089 dp3m.heffte.fft->ks_local_size(),
1090 dp3m.heffte.fft->ks_local_ld_index());
1096template <
typename FloatType, Arch Architecture,
class FFTConfig>
1100 FFTConfig::k_space_order>(
1101 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop,
1102 get_system().
box_geo->length_inv());
1103#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
1104 if (dp3m.heffte.world_size == 1) {
1105 dp3m.heffte.g_energy =
1107 FFTConfig::k_space_order>(
1108 dp3m.params, dp3m.heffte.fft->ks_local_ld_index(),
1109 dp3m.heffte.fft->ks_local_ur_index(),
1110 get_system().
box_geo->length_inv());
1111 if constexpr (FFTConfig::use_r2c) {
1113 dp3m.heffte.g_energy, dp3m.params.mesh,
1114 dp3m.heffte.fft->ks_local_size(),
1115 dp3m.heffte.fft->ks_local_ld_index());
1121template <
typename FloatType, Arch Architecture,
class FFTConfig>
1124 int m_mesh_max = -1, m_mesh_min = -1;
1125 std::pair<std::optional<int>, std::optional<int>> m_tune_limits;
1129 double prefactor,
int timings,
1138 std::optional<std::string>
1145 m_logger = std::make_unique<TuningLogger>(
1153 std::tuple<double, double, double, double>
1155 double r_cut_iL)
const override {
1169 0.0001 * box_geo.length()[0], 5. * box_geo.length()[0], 0.0001,
1194 m_mesh_min =
static_cast<int>(std::round(std::pow(2., std::floor(
expo))));
1197 if (m_tune_limits.first) {
1198 m_mesh_min = *m_tune_limits.first;
1200 if (m_tune_limits.second) {
1201 m_mesh_max = *m_tune_limits.second;
1204 m_mesh_min = m_mesh_max = dp3m.
params.
mesh[0];
1244template <
typename FloatType, Arch Architecture,
class FFTConfig>
1246 auto &
system = get_system();
1247 auto const &box_geo = *
system.box_geo;
1254 if (
not is_tuned()) {
1257 throw std::runtime_error(
1258 "DipolarP3M: no dipolar particles in the system");
1262 system, dp3m, prefactor, tuning.timings, tuning.limits);
1271 system.on_dipoles_change();
1302 [&](
unsigned dim,
int n) {
1303 nm[dim] = shift[dim] + n * mesh;
1312 std::size_t
n_c_part,
double sum_q2,
1316 auto const mesh_i = 1. /
static_cast<double>(mesh);
1334 Utils::int_pow<3>(
static_cast<double>(
n2));
1345 return 8. *
Utils::sqr(std::numbers::pi) / 3. * sum_q2 *
1357 std::size_t
n_c_part,
double sum_q2,
1359 auto constexpr exp_min = -708.4;
1391 double sum_q2,
double x1,
double x2,
double xacc,
1403 if (
f1 *
f2 >= 0.0) {
1404 throw std::runtime_error(
1405 "Root must be bracketed for bisection in dp3m_rtbisection");
1420 throw std::runtime_error(
"Too many bisections in dp3m_rtbisection");
1425 auto const &box_geo = *
system.box_geo;
1426 auto const &local_geo = *
system.local_geo;
1427 for (
auto i = 0
u; i < 3u; i++) {
1430 std::stringstream
msg;
1432 <<
" is larger than half of box dimension " << box_geo.length()[i];
1433 throw std::runtime_error(
msg.str());
1436 std::stringstream
msg;
1438 <<
" is larger than local box dimension " << local_geo.length()[i];
1439 throw std::runtime_error(
msg.str());
1443 if ((box_geo.length()[0] != box_geo.length()[1])
or
1444 (box_geo.length()[1] != box_geo.length()[2])) {
1445 throw std::runtime_error(
"DipolarP3M: requires a cubic box");
1451 if (!box_geo.periodic(0)
or !box_geo.periodic(1)
or !box_geo.periodic(2)) {
1452 throw std::runtime_error(
1453 "DipolarP3M: requires periodicity (True, True, True)");
1458 auto const &local_geo = *
get_system().local_geo;
1461 throw std::runtime_error(
1462 "DipolarP3M: requires the regular or hybrid decomposition cell system");
1466 throw std::runtime_error(
1467 "DipolarP3M: does not work with the hybrid decomposition cell system, "
1468 "if using more than one MPI node");
1474 if (node_grid[0] < node_grid[1]
or node_grid[1] < node_grid[2]) {
1475 throw std::runtime_error(
1476 "DipolarP3M: node grid must be sorted, largest first");
1480template <
typename FloatType, Arch Architecture,
class FFTConfig>
1482 auto const &box_geo = *get_system().
box_geo;
1487 sanity_checks_boxl();
1488 calc_influence_function_force();
1489 calc_influence_function_energy();
1491#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
1498template <
typename FloatType, Arch Architecture,
class FFTConfig>
1501 auto const &box_geo = *get_system().
box_geo;
1502 auto const Ukp3m = calc_average_self_energy_k_space() * box_geo.volume();
1510template <
typename FloatType, Arch Architecture,
class FFTConfig>
@ HYBRID
Hybrid decomposition.
@ REGULAR
Regular decomposition.
Vector implementation and trait types for boost qvm interoperability.
Describes a cell structure / cell system.
auto const & get_unique_particles() const
auto get_scatter_dip_fld()
auto get_scatter_torque()
void mark_torque_replicas_dirty()
Declare that a kernel scattering into the torque view is about to run.
std::tuple< double, double, double, double > calculate_accuracy(Utils::Vector3i const &mesh, int cao, double r_cut_iL) const override
TuningAlgorithm::Parameters get_time() override
DipolarTuningAlgorithm(System::System &system, decltype(dp3m) &input_dp3m, double prefactor, int timings, decltype(m_tune_limits) tune_limits)
void on_solver_change() const override
void determine_mesh_limits() override
P3MParameters & get_params() override
std::optional< std::string > layer_correction_veto_r_cut(double) const override
void setup_logger(bool verbose) override
base_type::size_type size() const
void npt_add_virial_contribution(double energy)
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
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) noexcept
Create a vector that has all entries set to the same value.
void zfill(std::size_t size)
Fill cache with zero-initialized data.
cudaStream_t stream[1]
CUDA streams for parallel computing on CPU and GPU.
Communicator communicator
boost::mpi::communicator comm_cart
The communicator.
int this_node
The number of this node.
constexpr auto round_error_prec
Precision below which a double-precision float is assumed to be zero.
__device__ void vector_product(float const *a, float const *b, float *out)
static std::size_t count_magnetic_particles(ParticleRange const &particles)
P3M algorithm for long-range magnetic dipole-dipole interaction.
double dp3m_real_space_error(double box_size, double r_cut_iL, std::size_t 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 ...
double dp3m_rtbisection(double box_size, double r_cut_iL, std::size_t n_c_part, double sum_q2, double x1, double x2, double xacc, double tuned_accuracy)
Compute the value of alpha through a bisection method.
auto dp3m_tune_aliasing_sums(Utils::Vector3i const &shift, int mesh, double mesh_i, int cao, double alpha_L_i)
Tuning dipolar-P3M.
double dp3m_k_space_error(double box_size, int mesh, int cao, std::size_t n_c_part, double sum_q2, double alpha_L)
Calculate the k-space error of dipolar-P3M.
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.
void extract_block_into(OutValue *out, Container const &in_array, Utils::Vector3i const &dimensions, Utils::Vector3i const &start, Utils::Vector3i const &stop)
Extract a 3D block from the halo field into a caller-provided buffer.
void pad_with_zeros_discard_imag_into(OutValue *out, std::span< T > cropped_array, Utils::Vector3i const &cropped_dim, Utils::Vector3i const &pad_left, Utils::Vector3i const &pad_right)
Pad a 3D matrix with zeros to restore halo regions, writing into a caller-provided buffer of product(...
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.
std::vector< FloatType > grid_influence_function_dipolar(P3MParameters const ¶ms, Utils::Vector3i const &n_start, Utils::Vector3i const &n_stop, Utils::Vector3d const &inv_box_l)
Map influence function over a grid.
void p3m_interpolate(P3MLocalMesh const &local_mesh, WeightsStorage< cao > const &weights, Kernel kernel)
P3M grid interpolation.
constexpr int p3m_min_cao
Minimal charge assignment order.
constexpr int p3m_max_cao
Maximal charge assignment order.
#define P3M_BRILLOUIN
P3M: Number of Brillouin zones taken into account in the calculation of the optimal influence functio...
T product(Vector< T, N > const &v)
DEVICE_QUALIFIER constexpr T sqr(T x)
Calculates the SQuaRe of x.
decltype(auto) integral_parameter(T i, Args &&...args)
Generate a call table for an integral non-type template parameter.
DEVICE_QUALIFIER auto sinc(T x)
Calculate the function .
auto get_analytic_cotangent_sum_kernel(int cao)
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
double calc_surface_term(bool force_flag, bool energy_flag) override
void dipole_assign() override
Utils::Vector9d long_range_pressure() override
Reciprocal-space virial for the dipolar Ewald/P3M sum.
void scaleby_box_l() override
double long_range_kernel(bool force_flag, bool energy_flag)
Compute the k-space part of forces and energies.
Base class for the magnetostatics P3M algorithm.
double sum_mu2
Sum of square of magnetic dipoles.
p3m_interpolation_cache inter_weights
std::size_t sum_dip_part
number of dipolar particles.
p3m_send_mesh< FloatType > halo_comm
double energy_correction
cached k-space self-energy correction
void resize_heffte_buffers()
struct DipolarP3MState::@1 heffte
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.
P3MLocalMesh local_mesh
Local mesh geometry information for this MPI rank.
P3MParameters params
P3M base parameters.
void operator()(auto &dp3m, auto &cell_structure)
void operator()(auto &dp3m, double pref, int d_rs, CellStructure &cell_structure) const
void operator()(auto &dp3m, double pref, int d_rs, CellStructure &cell_structure) const