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();
278#ifdef ESPRESSO_DIPOLE_FIELD_TRACKING
282 "AssignTorques", std::size_t{0
u}, n_part, [&](std::size_t
p_index) {
283 auto const &p = *unique_particles.at(
p_index);
284 if (p.dipm() != 0.) {
299 assert(cao == dp3m.inter_weights.cao());
300 using execution_space = Kokkos::DefaultHostExecutionSpace;
302 auto const kernel = [
d_rs, pref, &dp3m](
auto const &dip,
auto &
p_force,
309 E[0u] += w * double(dp3m.mesh.rs_fields[0u][ind]);
310 E[1u] += w * double(dp3m.mesh.rs_fields[1u][ind]);
311 E[2u] += w * double(dp3m.mesh.rs_fields[2u][ind]);
314 auto access =
p_force.access();
318 auto const n_part = dp3m.inter_weights.size();
322 "AssignForcesDip", std::size_t{0
u}, n_part, [&](std::size_t
p_index) {
323 auto const &p = *unique_particles.at(
p_index);
324 if (p.dipm() != 0.) {
332#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
333template <
typename FloatType,
class FFTConfig>
338 static_cast<std::size_t
>(
Utils::product(this->local_mesh.dim_no_halo));
340 static_cast<std::size_t
>(
Utils::product(heffte.fft->ks_local_size()));
341 for (
auto d : {0
u, 1u, 2u}) {
389template <
typename FloatType, Arch Architecture,
class FFTConfig>
392 auto const &
system = get_system();
393 auto const &box_geo = *
system.box_geo;
397 if (dp3m.sum_mu2 > 0.) {
399 dp3m.fft_buffers->perform_vector_halo_gather();
400 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
401 dp3m.fft->forward_fft(rs_mesh);
403 dp3m.update_mesh_views();
407 auto const wavevector = 2. * std::numbers::pi * box_geo.length_inv()[0];
410 auto index = std::size_t(0
u);
413 auto constexpr KX = 2, KY = 0, KZ = 1;
414 auto const shift = local_index + dp3m.mesh.start;
415 auto const &d_op = dp3m.d_op[0u];
416 auto const &mesh_dip = dp3m.mesh.rs_fields;
417 auto const d_op_x = static_cast<FloatType>(d_op[shift[KX]]);
418 auto const d_op_y = static_cast<FloatType>(d_op[shift[KY]]);
419 auto const d_op_z = static_cast<FloatType>(d_op[shift[KZ]]);
422 auto const Mx_re = mesh_dip[0u][index];
423 auto const My_re = mesh_dip[1u][index];
424 auto const Mz_re = mesh_dip[2u][index];
425 auto const Q_re = Mx_re * d_op_x + My_re * d_op_y + Mz_re * d_op_z;
428 auto const Mx_im = mesh_dip[0u][index];
429 auto const My_im = mesh_dip[1u][index];
430 auto const Mz_im = mesh_dip[2u][index];
431 auto const Q_im = Mx_im * d_op_x + My_im * d_op_y + Mz_im * d_op_z;
434 auto const nx = static_cast<double>(d_op[shift[KX]]);
435 auto const ny = static_cast<double>(d_op[shift[KY]]);
436 auto const nz = static_cast<double>(d_op[shift[KZ]]);
437 auto const kx = nx * wavevector;
438 auto const ky = ny * wavevector;
439 auto const kz = nz * wavevector;
440 auto const norm_sq = Utils::sqr(kx) + Utils::sqr(ky) + Utils::sqr(kz);
442 auto const g = static_cast<double>(*it_energy);
443 auto const cell_energy =
444 g * static_cast<double>(Utils::sqr(Q_re) + Utils::sqr(Q_im));
445 auto const vterm = -2. * (1. / norm_sq + half_alpha_inv_sq);
448 auto const Rx = g * static_cast<double>(Mx_re * Q_re + Mx_im * Q_im);
449 auto const Ry = g * static_cast<double>(My_re * Q_re + My_im * Q_im);
450 auto const Rz = g * static_cast<double>(Mz_re * Q_re + Mz_im * Q_im);
461 node_k_space_pressure_tensor[0u] +=
462 cell_energy * (1. + vterm * kx * kx) + 2. * nx * Rx;
463 node_k_space_pressure_tensor[1u] +=
464 cell_energy * vterm * kx * ky + 2. * ny * Rx;
465 node_k_space_pressure_tensor[2u] +=
466 cell_energy * vterm * kx * kz + 2. * nz * Rx;
467 node_k_space_pressure_tensor[3u] +=
468 cell_energy * vterm * ky * kx + 2. * nx * Ry;
469 node_k_space_pressure_tensor[4u] +=
470 cell_energy * (1. + vterm * ky * ky) + 2. * ny * Ry;
471 node_k_space_pressure_tensor[5u] +=
472 cell_energy * vterm * ky * kz + 2. * nz * Ry;
473 node_k_space_pressure_tensor[6u] +=
474 cell_energy * vterm * kz * kx + 2. * nx * Rz;
475 node_k_space_pressure_tensor[7u] +=
476 cell_energy * vterm * kz * ky + 2. * ny * Rz;
477 node_k_space_pressure_tensor[8u] +=
478 cell_energy * (1. + vterm * kz * kz) + 2. * nz * Rz;
485 box_geo.length_inv()[0];
488template <
typename FloatType, Arch Architecture,
class FFTConfig>
493 auto const &
system = get_system();
494 auto const &box_geo = *
system.box_geo;
504#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
505 auto constexpr r2c_dir = FFTConfig::r2c_dir;
506 auto const rs_local_size = dp3m.heffte.fft->rs_local_size();
507 auto const local_size = dp3m.heffte.fft->ks_local_size();
509 if constexpr (FFTConfig::use_r2c) {
513 auto const local_origin = dp3m.heffte.fft->ks_local_ld_index();
522 dp3m.resize_heffte_buffers();
525 if (dp3m.sum_mu2 > 0.) {
527 dp3m.fft_buffers->perform_vector_halo_gather();
528 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
529 dp3m.fft->forward_fft(rs_mesh);
531 dp3m.update_mesh_views();
533#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
534 if (dp3m.heffte.world_size == 1) {
536 std::array<FloatType *, 3u> rs_fields = {
537 {dp3m.heffte.rs_dipole_density[0
u].data(),
538 dp3m.heffte.rs_dipole_density[1u].data(),
539 dp3m.heffte.rs_dipole_density[2u].data()}};
540 dp3m.heffte.halo_comm.gather_grid(
::comm_cart, rs_fields,
541 dp3m.local_mesh.dim);
543 for (
auto dir : {0
u, 1u, 2u}) {
546 FFTConfig::r_space_order>(
547 dp3m.rs_field_no_halo_kokkos.data(),
548 dp3m.heffte.rs_dipole_density[dir], dp3m.local_mesh.dim,
549 dp3m.local_mesh.n_halo_ld,
550 dp3m.local_mesh.dim - dp3m.local_mesh.n_halo_ur);
555 auto constexpr KX = 1,
KY = 2,
KZ = 0;
560 dp3m.rs_field_no_halo_kokkos(index);
563 dp3m.heffte.fft->forward(dp3m.rs_field_no_halo_reorder_kokkos.data(),
564 dp3m.heffte.ks_dipole_density[dir].data());
566 if (
not dp3m.params.tuning) {
571 auto constexpr KX = 2,
KY = 0,
KZ = 1;
575 auto const old_value = std::complex<FloatType>{
596 if (dp3m.sum_mu2 > 0.) {
600 auto index = std::size_t(0
u);
604 auto constexpr KX = 2, KY = 0, KZ = 1;
605 auto const shift = local_index + dp3m.mesh.start;
606 auto const &d_op = dp3m.d_op[0u];
607 auto const &mesh_dip = dp3m.mesh.rs_fields;
609 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
610 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
611 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
614 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
615 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
616 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
618 node_energy += *it_energy * (Utils::sqr(re) + Utils::sqr(im));
619 std::advance(it_energy, 1);
621#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
622 if (dp3m.heffte.world_size == 1) {
629 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
660 if (dp3m.energy_correction == 0.)
661 calc_energy_correction();
665 energy -= prefactor * dp3m.sum_mu2 * std::numbers::inv_sqrtpi *
666 (2. / 3.) * Utils::int_pow<3>(dp3m.params.alpha);
669 energy += prefactor * dp3m.energy_correction / box_geo.volume();
679 if (dp3m.sum_mu2 > 0.) {
680 auto const wavenumber = 2. * std::numbers::pi * box_geo.length_inv()[0
u];
681 dp3m.ks_scalar.resize(dp3m.local_mesh.size);
684 auto index{std::size_t(0
u)};
688 auto constexpr KX = 2, KY = 0, KZ = 1;
689 auto const shift = local_index + dp3m.mesh.start;
690 auto const &d_op = dp3m.d_op[0u];
691 auto const &mesh_dip = dp3m.mesh.rs_fields;
693 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
694 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
695 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
698 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
699 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
700 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
702 *it_ks_scalar = *it_energy * std::complex<FloatType>{re, im};
707#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
708 if (dp3m.heffte.world_size == 1) {
714 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
724 if (
not dp3m.params.tuning) {
725 auto constexpr KX = 2,
KY = 0,
KZ = 1;
741 for (
int d = 0; d < 3; d++) {
745 auto const &offset = dp3m.mesh.start;
746 auto const &d_op = dp3m.d_op[0u];
747 auto const d_op_val = FloatType(d_op[local_index[d] + offset[d]]);
748 auto const &value = *it_ks_scalar;
749 dp3m.mesh.rs_scalar[index] = d_op_val * value.real();
751 dp3m.mesh.rs_scalar[index] = d_op_val * value.imag();
753 std::advance(it_ks_scalar, 1);
755#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
756 if (dp3m.heffte.world_size == 1) {
757 unsigned int constexpr d_ks[3] = {2u, 0
u, 1u};
768 if (
not dp3m.params.tuning) {
769 auto constexpr KX = 2,
KY = 0,
KZ = 1;
773 auto const old_value = std::complex<FloatType>{
784 dp3m.heffte.fft->backward(dp3m.heffte.ks_B_field_storage.data(),
785 dp3m.heffte.rs_B_fields_no_halo[d].data());
790 dp3m.heffte.rs_B_fields[d].data(),
791 std::span(dp3m.heffte.rs_B_fields_no_halo[d]),
792 dp3m.local_mesh.dim_no_halo, dp3m.local_mesh.n_halo_ld,
793 dp3m.local_mesh.n_halo_ur);
796 dp3m.heffte.rs_B_fields[d].data(),
797 dp3m.local_mesh.dim);
800 dp3m.fft->backward_fft(dp3m.fft_buffers->get_scalar_mesh());
802 dp3m.fft_buffers->perform_scalar_halo_spread();
804 auto const d_rs = (d + dp3m.mesh.ks_pnum) % 3;
805 Utils::integral_parameter<int, AssignTorques, p3m_min_cao, p3m_max_cao>(
819 auto it_force = dp3m.g_force.begin();
821 std::size_t index = 0
u;
823 auto constexpr KX = 2, KY = 0, KZ = 1;
824 auto const shift = local_index + dp3m.mesh.start;
825 auto const &d_op = dp3m.d_op[0u];
826 auto const &mesh_dip = dp3m.mesh.rs_fields;
828 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
829 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
830 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
833 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
834 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
835 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
837 *it_ks_scalar = {*it_force * im, *it_force * (-re)};
843#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
844 if (dp3m.heffte.world_size == 1) {
850 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
862 if (
not dp3m.params.tuning) {
863 auto constexpr KX = 2,
KY = 0,
KZ = 1;
879 for (
int d = 0; d < 3; d++) {
880 std::size_t index = 0
u;
883 auto constexpr KX = 2, KY = 0, KZ = 1;
884 auto const shift = local_index + dp3m.mesh.start;
885 auto const &d_op = dp3m.d_op[0u];
886 auto const &mesh_dip = dp3m.mesh.rs_fields;
887 auto const d_op_val = FloatType(d_op[shift[d]]);
888 auto const f = *it_ks_scalar * d_op_val;
889 mesh_dip[0u][index] = FloatType(d_op[shift[KX]]) * f.real();
890 mesh_dip[1u][index] = FloatType(d_op[shift[KY]]) * f.real();
891 mesh_dip[2u][index] = FloatType(d_op[shift[KZ]]) * f.real();
893 mesh_dip[0u][index] = FloatType(d_op[shift[KX]]) * f.imag();
894 mesh_dip[1u][index] = FloatType(d_op[shift[KY]]) * f.imag();
895 mesh_dip[2u][index] = FloatType(d_op[shift[KZ]]) * f.imag();
897 std::advance(it_ks_scalar, 1);
900#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
901 if (dp3m.heffte.world_size == 1) {
906 auto constexpr KX = 1,
KY = 2,
KZ = 0;
913 auto &
mesh_dip = dp3m.heffte.ks_dipole_density;
924 if (
not FFTConfig::use_r2c
and not dp3m.params.tuning) {
928 for (
int j = 0;
j < 3; ++
j) {
929 auto const old_value = std::complex<FloatType>{
940 for (
int dir = 0
u; dir < 3u; ++dir) {
941 dp3m.heffte.fft->backward(
942 dp3m.heffte.ks_dipole_density[dir].data(),
943 dp3m.heffte.rs_B_fields_no_halo[dir].data());
948 dp3m.heffte.rs_B_fields[d].data(),
949 std::span(dp3m.heffte.rs_B_fields_no_halo[dir]),
950 dp3m.local_mesh.dim_no_halo, dp3m.local_mesh.n_halo_ld,
951 dp3m.local_mesh.n_halo_ur);
955 std::array<FloatType *, 3u>{{dp3m.heffte.rs_B_fields[0
u].data(),
956 dp3m.heffte.rs_B_fields[1u].data(),
957 dp3m.heffte.rs_B_fields[2u].data()}};
958 dp3m.heffte.halo_comm.spread_grid(
::comm_cart, rs_fields,
959 dp3m.local_mesh.dim);
962 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
963 dp3m.fft->backward_fft(rs_mesh);
966 dp3m.fft_buffers->perform_vector_halo_spread();
968 auto const d_rs = (d + dp3m.mesh.ks_pnum) % 3;
1002template <
typename FloatType, Arch Architecture,
class FFTConfig>
1005 auto const &
system = get_system();
1006 auto const &box_geo = *
system.box_geo;
1007 auto const particles =
system.cell_structure->local_particles();
1008 auto const pref = prefactor * 4. * std::numbers::pi / box_geo.volume() /
1009 (2. * dp3m.params.epsilon + 1.);
1018 std::size_t
ip = 0
u;
1019 for (
auto const &p : particles) {
1020 auto const dip = p.calc_dip();
1044 0.5 * pref * boost::mpi::all_reduce(
comm_cart,
sum_e, std::plus<>());
1060 for (
auto &p : particles) {
1061 auto &torque = p.torque();
1063 torque[1u] -= pref *
sumiy[
ip];
1064 torque[2u] -= pref *
sumiz[
ip];
1065#ifdef ESPRESSO_DIPOLE_FIELD_TRACKING
1066 p.dip_fld() -= pref *
box_dip;
1075template <
typename FloatType, Arch Architecture,
class FFTConfig>
1079 FFTConfig::k_space_order>(
1080 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop,
1081 get_system().
box_geo->length_inv());
1082#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
1083 if (dp3m.heffte.world_size == 1) {
1084 dp3m.heffte.g_force =
1086 FFTConfig::k_space_order>(
1087 dp3m.params, dp3m.heffte.fft->ks_local_ld_index(),
1088 dp3m.heffte.fft->ks_local_ur_index(),
1089 get_system().
box_geo->length_inv());
1090 if constexpr (FFTConfig::use_r2c) {
1092 dp3m.heffte.g_force, dp3m.params.mesh,
1093 dp3m.heffte.fft->ks_local_size(),
1094 dp3m.heffte.fft->ks_local_ld_index());
1100template <
typename FloatType, Arch Architecture,
class FFTConfig>
1104 FFTConfig::k_space_order>(
1105 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop,
1106 get_system().
box_geo->length_inv());
1107#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
1108 if (dp3m.heffte.world_size == 1) {
1109 dp3m.heffte.g_energy =
1111 FFTConfig::k_space_order>(
1112 dp3m.params, dp3m.heffte.fft->ks_local_ld_index(),
1113 dp3m.heffte.fft->ks_local_ur_index(),
1114 get_system().
box_geo->length_inv());
1115 if constexpr (FFTConfig::use_r2c) {
1117 dp3m.heffte.g_energy, dp3m.params.mesh,
1118 dp3m.heffte.fft->ks_local_size(),
1119 dp3m.heffte.fft->ks_local_ld_index());
1125template <
typename FloatType, Arch Architecture,
class FFTConfig>
1128 int m_mesh_max = -1, m_mesh_min = -1;
1129 std::pair<std::optional<int>, std::optional<int>> m_tune_limits;
1133 double prefactor,
int timings,
1142 std::optional<std::string>
1149 m_logger = std::make_unique<TuningLogger>(
1157 std::tuple<double, double, double, double>
1159 double r_cut_iL)
const override {
1173 0.0001 * box_geo.length()[0], 5. * box_geo.length()[0], 0.0001,
1198 m_mesh_min =
static_cast<int>(std::round(std::pow(2., std::floor(
expo))));
1201 if (m_tune_limits.first) {
1202 m_mesh_min = *m_tune_limits.first;
1204 if (m_tune_limits.second) {
1205 m_mesh_max = *m_tune_limits.second;
1208 m_mesh_min = m_mesh_max = dp3m.
params.
mesh[0];
1248template <
typename FloatType, Arch Architecture,
class FFTConfig>
1250 auto &
system = get_system();
1251 auto const &box_geo = *
system.box_geo;
1258 if (
not is_tuned()) {
1261 throw std::runtime_error(
1262 "DipolarP3M: no dipolar particles in the system");
1266 system, dp3m, prefactor, tuning.timings, tuning.limits);
1275 system.on_dipoles_change();
1306 [&](
unsigned dim,
int n) {
1307 nm[dim] = shift[dim] + n * mesh;
1316 std::size_t
n_c_part,
double sum_q2,
1320 auto const mesh_i = 1. /
static_cast<double>(mesh);
1338 Utils::int_pow<3>(
static_cast<double>(
n2));
1349 return 8. *
Utils::sqr(std::numbers::pi) / 3. * sum_q2 *
1361 std::size_t
n_c_part,
double sum_q2,
1363 auto constexpr exp_min = -708.4;
1395 double sum_q2,
double x1,
double x2,
double xacc,
1407 if (
f1 *
f2 >= 0.0) {
1408 throw std::runtime_error(
1409 "Root must be bracketed for bisection in dp3m_rtbisection");
1424 throw std::runtime_error(
"Too many bisections in dp3m_rtbisection");
1429 auto const &box_geo = *
system.box_geo;
1430 auto const &local_geo = *
system.local_geo;
1431 for (
auto i = 0
u; i < 3u; i++) {
1434 std::stringstream
msg;
1436 <<
" is larger than half of box dimension " << box_geo.length()[i];
1437 throw std::runtime_error(
msg.str());
1440 std::stringstream
msg;
1442 <<
" is larger than local box dimension " << local_geo.length()[i];
1443 throw std::runtime_error(
msg.str());
1447 if ((box_geo.length()[0] != box_geo.length()[1])
or
1448 (box_geo.length()[1] != box_geo.length()[2])) {
1449 throw std::runtime_error(
"DipolarP3M: requires a cubic box");
1455 if (!box_geo.periodic(0)
or !box_geo.periodic(1)
or !box_geo.periodic(2)) {
1456 throw std::runtime_error(
1457 "DipolarP3M: requires periodicity (True, True, True)");
1462 auto const &local_geo = *
get_system().local_geo;
1465 throw std::runtime_error(
1466 "DipolarP3M: requires the regular or hybrid decomposition cell system");
1470 throw std::runtime_error(
1471 "DipolarP3M: does not work with the hybrid decomposition cell system, "
1472 "if using more than one MPI node");
1478 if (node_grid[0] < node_grid[1]
or node_grid[1] < node_grid[2]) {
1479 throw std::runtime_error(
1480 "DipolarP3M: node grid must be sorted, largest first");
1484template <
typename FloatType, Arch Architecture,
class FFTConfig>
1486 auto const &box_geo = *get_system().
box_geo;
1491 sanity_checks_boxl();
1492 calc_influence_function_force();
1493 calc_influence_function_energy();
1495#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
1502template <
typename FloatType, Arch Architecture,
class FFTConfig>
1505 auto const &box_geo = *get_system().
box_geo;
1506 auto const Ukp3m = calc_average_self_energy_k_space() * box_geo.volume();
1514template <
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()
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