49#include "communication.hpp"
53#include "system/System.hpp"
61#include <boost/mpi/collectives/all_reduce.hpp>
62#include <boost/mpi/collectives/reduce.hpp>
64#include <Kokkos_Core.hpp>
65#include <Kokkos_ScatterView.hpp>
84#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
88 auto const diff = std::abs(value - reference);
89 using FT = std::remove_cvref_t<
decltype(
diff)>;
90 auto constexpr atol = std::is_same_v<FT, float> ?
FT{2
e-4} :
FT{1
e-6};
91 auto constexpr rtol = std::is_same_v<FT, float> ?
FT{5
e-5} :
FT{1
e-5};
92 auto const non_zero = std::abs(reference) !=
FT{0};
98template <
typename FloatType, Arch Architecture,
class FFTConfig>
104 for (
auto const &p : get_system().
cell_structure->local_particles()) {
105 if (p.dipm() != 0.) {
112 boost::mpi::all_reduce(
comm_cart,
local_n, dp3m.sum_dip_part, std::plus<>());
116 std::size_t
n_c_part,
double sum_q2,
120 std::size_t
n_c_part,
double sum_q2,
127 double sum_q2,
double x1,
double x2,
double xacc,
130template <
typename FloatType, Arch Architecture,
class FFTConfig>
133 auto const &box_geo = *get_system().
box_geo;
136 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop, dp3m.g_energy);
140 phi /= 3. * box_geo.length()[0] * Utils::int_pow<3>(dp3m.params.mesh[0]);
141 return phi * std::numbers::pi;
144template <
typename FloatType, Arch Architecture,
class FFTConfig>
148 assert(dp3m.params.alpha > 0.);
150 auto const &
system = get_system();
151 auto const &box_geo = *
system.box_geo;
152 auto const &local_geo = *
system.local_geo;
155 dp3m.params.cao3 = Utils::int_pow<3>(dp3m.params.cao);
156 dp3m.params.recalc_a_ai_cao_cut(box_geo.length());
159 dp3m.local_mesh.calc_local_ca_mesh(dp3m.params, local_geo,
verlet_skin, 0.);
160 dp3m.fft_buffers->init_halo();
161 dp3m.fft->init(dp3m.params);
162 dp3m.mesh.ks_pnum = dp3m.fft->get_ks_pnum();
163 dp3m.fft_buffers->init_meshes(dp3m.fft->get_ca_mesh_size());
164 dp3m.update_mesh_views();
165#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
166 dp3m.heffte.world_size =
comm_cart.size();
168 std::make_shared<P3MFFT<FloatType, Architecture, FFTConfig>>(
169 nullptr,
::comm_cart, dp3m.params.mesh, dp3m.local_mesh.ld_no_halo,
171 dp3m.resize_heffte_buffers();
173 dp3m.calc_differential_operator();
186 auto const &aosoa = cell_structure.get_aosoa();
187 auto const &unique_particles = cell_structure.get_unique_particles();
188 auto const n_part = cell_structure.count_local_particles();
191 "InterpolateDipoles", std::size_t{0
u}, n_part, [&](
auto p_index) {
194 auto const p_pos = aosoa.get_span_at(aosoa.position,
p_index);
195 auto const dip = unique_particles.at(
p_index)->calc_dip();
198 p_pos, dp3m.params.ai, dp3m.local_mesh);
199 dp3m.inter_weights.store_at(
p_index, weights);
201 dp3m.local_mesh, weights, [&dip,
tid, &dp3m](
int ind,
double w) {
202 dp3m.rs_fields_kokkos(tid, 0u, ind) += value_type(w * dip[0u]);
203 dp3m.rs_fields_kokkos(tid, 1u, ind) += value_type(w * dip[1u]);
204 dp3m.rs_fields_kokkos(tid, 2u, ind) += value_type(w * dip[2u]);
208 using execution_space = Kokkos::DefaultHostExecutionSpace;
210 Kokkos::RangePolicy<execution_space> policy(std::size_t{0},
211 dp3m.local_mesh.size);
212 Kokkos::parallel_for(
"ReduceInterpolatedDipoles", policy,
214 for (
int dir = 0; dir < 3; ++dir) {
217 acc += dp3m.rs_fields_kokkos(
tid, dir, i);
219 dp3m.mesh.rs_fields[dir][i] += acc;
220#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
221 dp3m.heffte.rs_dipole_density[dir][i] += acc;
230template <
typename FloatType, Arch Architecture,
class FFTConfig>
234 Utils::integral_parameter<int, AssignDipole, p3m_min_cao, p3m_max_cao>(
235 dp3m.params.cao, dp3m, *get_system().cell_structure);
243 assert(cao == dp3m.inter_weights.cao());
245 auto const kernel = [
d_rs, &dp3m](
auto const &pref,
auto &
p_torque,
250 [&
E, &dp3m,
d_rs](
int ind,
double w) {
252 E[d_rs] += w * double(dp3m.mesh.rs_scalar[ind]);
257 access(
p_index, 0) -= torque[0];
258 access(
p_index, 1) -= torque[1];
259 access(
p_index, 2) -= torque[2];
262 auto const n_part = dp3m.inter_weights.size();
266 "AssignTorques", std::size_t{0
u}, n_part, [&](std::size_t
p_index) {
267 auto const &p = *unique_particles.at(
p_index);
268 if (p.dipm() != 0.) {
279 assert(cao == dp3m.inter_weights.cao());
281 auto const kernel = [
d_rs, &dp3m](
auto const &pref,
auto &
p_force,
288 E[0u] += w * double(dp3m.mesh.rs_fields[0u][ind]);
289 E[1u] += w * double(dp3m.mesh.rs_fields[1u][ind]);
290 E[2u] += w * double(dp3m.mesh.rs_fields[2u][ind]);
293 auto access =
p_force.access();
297 auto const n_part = dp3m.inter_weights.size();
301 "AssignForcesDip", std::size_t{0
u}, n_part, [&](std::size_t
p_index) {
302 auto const &p = *unique_particles.at(
p_index);
303 if (p.dipm() != 0.) {
311#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
312template <
typename FloatType,
class FFTConfig>
317 static_cast<std::size_t
>(
Utils::product(this->local_mesh.dim_no_halo));
319 static_cast<std::size_t
>(
Utils::product(heffte.fft->ks_local_size()));
320 for (
auto d : {0
u, 1u, 2u}) {
331template <
typename FloatType, Arch Architecture,
class FFTConfig>
336 auto const &
system = get_system();
337 auto const &box_geo = *
system.box_geo;
347#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
348 auto constexpr r2c_dir = FFTConfig::r2c_dir;
349 auto const rs_local_size = dp3m.heffte.fft->rs_local_size();
350 auto const local_size = dp3m.heffte.fft->ks_local_size();
352 if constexpr (FFTConfig::use_r2c) {
356 auto const local_origin = dp3m.heffte.fft->ks_local_ld_index();
365 dp3m.resize_heffte_buffers();
368 if (dp3m.sum_mu2 > 0.) {
370 dp3m.fft_buffers->perform_vector_halo_gather();
371 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
372 dp3m.fft->forward_fft(rs_mesh);
374 dp3m.update_mesh_views();
376#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
377 if (dp3m.heffte.world_size == 1) {
379 std::array<FloatType *, 3u> rs_fields = {
380 {dp3m.heffte.rs_dipole_density[0
u].data(),
381 dp3m.heffte.rs_dipole_density[1u].data(),
382 dp3m.heffte.rs_dipole_density[2u].data()}};
383 dp3m.heffte.halo_comm.gather_grid(
::comm_cart, rs_fields,
384 dp3m.local_mesh.dim);
386 for (
auto dir : {0
u, 1u, 2u}) {
389 FFTConfig::r_space_order>(
390 dp3m.heffte.rs_dipole_density[dir], dp3m.local_mesh.dim,
391 dp3m.local_mesh.n_halo_ld,
392 dp3m.local_mesh.dim - dp3m.local_mesh.n_halo_ur);
399 auto constexpr KX = 1,
KY = 2,
KZ = 0;
408 dp3m.heffte.ks_dipole_density[dir].data());
410 if (
not dp3m.params.tuning) {
415 auto constexpr KX = 2,
KY = 0,
KZ = 1;
419 auto const old_value = std::complex<FloatType>{
440 if (dp3m.sum_mu2 > 0.) {
444 auto index = std::size_t(0
u);
448 auto constexpr KX = 2, KY = 0, KZ = 1;
449 auto const shift = local_index + dp3m.mesh.start;
450 auto const &d_op = dp3m.d_op[0u];
451 auto const &mesh_dip = dp3m.mesh.rs_fields;
453 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
454 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
455 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
458 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
459 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
460 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
462 node_energy += *it_energy * (Utils::sqr(re) + Utils::sqr(im));
463 std::advance(it_energy, 1);
465#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
466 if (dp3m.heffte.world_size == 1) {
473 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
504 if (dp3m.energy_correction == 0.)
505 calc_energy_correction();
509 energy -= prefactor * dp3m.sum_mu2 * std::numbers::inv_sqrtpi *
510 (2. / 3.) * Utils::int_pow<3>(dp3m.params.alpha);
513 energy += prefactor * dp3m.energy_correction / box_geo.volume();
523 if (dp3m.sum_mu2 > 0.) {
524 auto const wavenumber = 2. * std::numbers::pi * box_geo.length_inv()[0
u];
525 dp3m.ks_scalar.resize(dp3m.local_mesh.size);
528 auto index{std::size_t(0
u)};
532 auto constexpr KX = 2, KY = 0, KZ = 1;
533 auto const shift = local_index + dp3m.mesh.start;
534 auto const &d_op = dp3m.d_op[0u];
535 auto const &mesh_dip = dp3m.mesh.rs_fields;
537 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
538 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
539 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
542 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
543 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
544 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
546 *it_ks_scalar = *it_energy * std::complex<FloatType>{re, im};
551#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
552 if (dp3m.heffte.world_size == 1) {
558 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
568 if (
not dp3m.params.tuning) {
569 auto constexpr KX = 2,
KY = 0,
KZ = 1;
585 for (
int d = 0; d < 3; d++) {
589 auto const &offset = dp3m.mesh.start;
590 auto const &d_op = dp3m.d_op[0u];
591 auto const d_op_val = FloatType(d_op[local_index[d] + offset[d]]);
592 auto const &value = *it_ks_scalar;
593 dp3m.mesh.rs_scalar[index] = d_op_val * value.real();
595 dp3m.mesh.rs_scalar[index] = d_op_val * value.imag();
597 std::advance(it_ks_scalar, 1);
599#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
600 if (dp3m.heffte.world_size == 1) {
601 unsigned int constexpr d_ks[3] = {2u, 0
u, 1u};
612 if (
not dp3m.params.tuning) {
613 auto constexpr KX = 2,
KY = 0,
KZ = 1;
617 auto const old_value = std::complex<FloatType>{
628 dp3m.heffte.fft->backward(dp3m.heffte.ks_B_field_storage.data(),
629 dp3m.heffte.rs_B_fields_no_halo[d].data());
631 dp3m.heffte.rs_B_fields[d] =
634 std::span(dp3m.heffte.rs_B_fields_no_halo[d]),
635 dp3m.local_mesh.dim_no_halo, dp3m.local_mesh.n_halo_ld,
636 dp3m.local_mesh.n_halo_ur);
639 dp3m.heffte.rs_B_fields[d].data(),
640 dp3m.local_mesh.dim);
643 dp3m.fft->backward_fft(dp3m.fft_buffers->get_scalar_mesh());
645 dp3m.fft_buffers->perform_scalar_halo_spread();
647 auto const d_rs = (d + dp3m.mesh.ks_pnum) % 3;
648 Utils::integral_parameter<int, AssignTorques, p3m_min_cao, p3m_max_cao>(
662 auto it_force = dp3m.g_force.begin();
664 std::size_t index = 0
u;
666 auto constexpr KX = 2, KY = 0, KZ = 1;
667 auto const shift = local_index + dp3m.mesh.start;
668 auto const &d_op = dp3m.d_op[0u];
669 auto const &mesh_dip = dp3m.mesh.rs_fields;
671 auto const re = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
672 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
673 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
676 auto const im = mesh_dip[0u][index] * FloatType(d_op[shift[KX]]) +
677 mesh_dip[1u][index] * FloatType(d_op[shift[KY]]) +
678 mesh_dip[2u][index] * FloatType(d_op[shift[KZ]]);
680 *it_ks_scalar = {*it_force * im, *it_force * (-re)};
686#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
687 if (dp3m.heffte.world_size == 1) {
693 auto const &
mesh_dip = dp3m.heffte.ks_dipole_density;
705 if (
not dp3m.params.tuning) {
706 auto constexpr KX = 2,
KY = 0,
KZ = 1;
722 for (
int d = 0; d < 3; d++) {
723 std::size_t index = 0
u;
726 auto constexpr KX = 2, KY = 0, KZ = 1;
727 auto const shift = local_index + dp3m.mesh.start;
728 auto const &d_op = dp3m.d_op[0u];
729 auto const &mesh_dip = dp3m.mesh.rs_fields;
730 auto const d_op_val = FloatType(d_op[shift[d]]);
731 auto const f = *it_ks_scalar * d_op_val;
732 mesh_dip[0u][index] = FloatType(d_op[shift[KX]]) * f.real();
733 mesh_dip[1u][index] = FloatType(d_op[shift[KY]]) * f.real();
734 mesh_dip[2u][index] = FloatType(d_op[shift[KZ]]) * f.real();
736 mesh_dip[0u][index] = FloatType(d_op[shift[KX]]) * f.imag();
737 mesh_dip[1u][index] = FloatType(d_op[shift[KY]]) * f.imag();
738 mesh_dip[2u][index] = FloatType(d_op[shift[KZ]]) * f.imag();
740 std::advance(it_ks_scalar, 1);
743#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
744 if (dp3m.heffte.world_size == 1) {
749 auto constexpr KX = 1,
KY = 2,
KZ = 0;
756 auto &
mesh_dip = dp3m.heffte.ks_dipole_density;
767 if (
not FFTConfig::use_r2c
and not dp3m.params.tuning) {
771 for (
int j = 0;
j < 3; ++
j) {
772 auto const old_value = std::complex<FloatType>{
783 for (
int dir = 0
u; dir < 3u; ++dir) {
784 dp3m.heffte.fft->backward(
785 dp3m.heffte.ks_dipole_density[dir].data(),
786 dp3m.heffte.rs_B_fields_no_halo[dir].data());
788 dp3m.heffte.rs_B_fields[d] =
791 std::span(dp3m.heffte.rs_B_fields_no_halo[dir]),
792 dp3m.local_mesh.dim_no_halo, dp3m.local_mesh.n_halo_ld,
793 dp3m.local_mesh.n_halo_ur);
797 std::array<FloatType *, 3u>{{dp3m.heffte.rs_B_fields[0
u].data(),
798 dp3m.heffte.rs_B_fields[1u].data(),
799 dp3m.heffte.rs_B_fields[2u].data()}};
800 dp3m.heffte.halo_comm.spread_grid(
::comm_cart, rs_fields,
801 dp3m.local_mesh.dim);
804 for (
auto &rs_mesh : dp3m.fft_buffers->get_vector_mesh()) {
805 dp3m.fft->backward_fft(rs_mesh);
808 dp3m.fft_buffers->perform_vector_halo_spread();
810 auto const d_rs = (d + dp3m.mesh.ks_pnum) % 3;
838template <
typename FloatType, Arch Architecture,
class FFTConfig>
841 auto const &
system = get_system();
842 auto const &box_geo = *
system.box_geo;
843 auto const particles =
system.cell_structure->local_particles();
844 auto const pref = prefactor * 4. * std::numbers::pi / box_geo.volume() /
845 (2. * dp3m.params.epsilon + 1.);
855 for (
auto const &p : particles) {
856 auto const dip = p.calc_dip();
880 0.5 * pref * boost::mpi::all_reduce(
comm_cart,
sum_e, std::plus<>());
896 for (
auto &p : particles) {
897 auto &torque = p.torque();
899 torque[1u] -= pref *
sumiy[
ip];
900 torque[2u] -= pref *
sumiz[
ip];
908template <
typename FloatType, Arch Architecture,
class FFTConfig>
912 FFTConfig::k_space_order>(
913 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop,
914 get_system().
box_geo->length_inv());
915#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
916 if (dp3m.heffte.world_size == 1) {
917 dp3m.heffte.g_force =
919 FFTConfig::k_space_order>(
920 dp3m.params, dp3m.heffte.fft->ks_local_ld_index(),
921 dp3m.heffte.fft->ks_local_ur_index(),
922 get_system().
box_geo->length_inv());
923 if constexpr (FFTConfig::use_r2c) {
925 dp3m.heffte.g_force, dp3m.params.mesh,
926 dp3m.heffte.fft->ks_local_size(),
927 dp3m.heffte.fft->ks_local_ld_index());
933template <
typename FloatType, Arch Architecture,
class FFTConfig>
937 FFTConfig::k_space_order>(
938 dp3m.params, dp3m.mesh.start, dp3m.mesh.stop,
939 get_system().
box_geo->length_inv());
940#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
941 if (dp3m.heffte.world_size == 1) {
942 dp3m.heffte.g_energy =
944 FFTConfig::k_space_order>(
945 dp3m.params, dp3m.heffte.fft->ks_local_ld_index(),
946 dp3m.heffte.fft->ks_local_ur_index(),
947 get_system().
box_geo->length_inv());
948 if constexpr (FFTConfig::use_r2c) {
950 dp3m.heffte.g_energy, dp3m.params.mesh,
951 dp3m.heffte.fft->ks_local_size(),
952 dp3m.heffte.fft->ks_local_ld_index());
958template <
typename FloatType, Arch Architecture,
class FFTConfig>
961 int m_mesh_max = -1, m_mesh_min = -1;
962 std::pair<std::optional<int>, std::optional<int>> m_tune_limits;
966 double prefactor,
int timings,
975 std::optional<std::string>
982 m_logger = std::make_unique<TuningLogger>(
990 std::tuple<double, double, double, double>
992 double r_cut_iL)
const override {
1006 0.0001 * box_geo.length()[0], 5. * box_geo.length()[0], 0.0001,
1031 m_mesh_min =
static_cast<int>(std::round(std::pow(2., std::floor(
expo))));
1034 if (m_tune_limits.first) {
1035 m_mesh_min = *m_tune_limits.first;
1037 if (m_tune_limits.second) {
1038 m_mesh_max = *m_tune_limits.second;
1041 m_mesh_min = m_mesh_max = dp3m.
params.
mesh[0];
1081template <
typename FloatType, Arch Architecture,
class FFTConfig>
1083 auto &
system = get_system();
1084 auto const &box_geo = *
system.box_geo;
1091 if (
not is_tuned()) {
1094 throw std::runtime_error(
1095 "DipolarP3M: no dipolar particles in the system");
1099 system, dp3m, prefactor, tuning.timings, tuning.limits);
1108 system.on_dipoles_change();
1139 [&](
unsigned dim,
int n) {
1140 nm[dim] = shift[dim] + n * mesh;
1149 std::size_t
n_c_part,
double sum_q2,
1153 auto const mesh_i = 1. /
static_cast<double>(mesh);
1171 Utils::int_pow<3>(
static_cast<double>(
n2));
1182 return 8. *
Utils::sqr(std::numbers::pi) / 3. * sum_q2 *
1194 std::size_t
n_c_part,
double sum_q2,
1196 auto constexpr exp_min = -708.4;
1228 double sum_q2,
double x1,
double x2,
double xacc,
1240 if (
f1 *
f2 >= 0.0) {
1241 throw std::runtime_error(
1242 "Root must be bracketed for bisection in dp3m_rtbisection");
1257 throw std::runtime_error(
"Too many bisections in dp3m_rtbisection");
1262 auto const &box_geo = *
system.box_geo;
1263 auto const &local_geo = *
system.local_geo;
1264 for (
auto i = 0
u; i < 3u; i++) {
1267 std::stringstream
msg;
1269 <<
" is larger than half of box dimension " << box_geo.length()[i];
1270 throw std::runtime_error(
msg.str());
1273 std::stringstream
msg;
1275 <<
" is larger than local box dimension " << local_geo.length()[i];
1276 throw std::runtime_error(
msg.str());
1280 if ((box_geo.length()[0] != box_geo.length()[1])
or
1281 (box_geo.length()[1] != box_geo.length()[2])) {
1282 throw std::runtime_error(
"DipolarP3M: requires a cubic box");
1288 if (!box_geo.periodic(0)
or !box_geo.periodic(1)
or !box_geo.periodic(2)) {
1289 throw std::runtime_error(
1290 "DipolarP3M: requires periodicity (True, True, True)");
1295 auto const &local_geo = *
get_system().local_geo;
1298 throw std::runtime_error(
1299 "DipolarP3M: requires the regular or hybrid decomposition cell system");
1303 throw std::runtime_error(
1304 "DipolarP3M: does not work with the hybrid decomposition cell system, "
1305 "if using more than one MPI node");
1311 if (node_grid[0] < node_grid[1]
or node_grid[1] < node_grid[2]) {
1312 throw std::runtime_error(
1313 "DipolarP3M: node grid must be sorted, largest first");
1317template <
typename FloatType, Arch Architecture,
class FFTConfig>
1319 auto const &box_geo = *get_system().
box_geo;
1324 sanity_checks_boxl();
1325 calc_influence_function_force();
1326 calc_influence_function_energy();
1328#ifdef ESPRESSO_DP3M_HEFFTE_CROSS_CHECKS
1335template <
typename FloatType, Arch Architecture,
class FFTConfig>
1338 auto const &box_geo = *get_system().
box_geo;
1339 auto const Ukp3m = calc_average_self_energy_k_space() * box_geo.volume();
1347template <
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_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.
auto pad_with_zeros_discard_imag(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.
auto extract_block(Container const &in_array, Utils::Vector3i const &dimensions, Utils::Vector3i const &start, Utils::Vector3i const &stop)
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.
ESPRESSO_ATTR_ALWAYS_INLINE void kokkos_parallel_range_for(auto const &name, auto start, auto end, auto const &kernel)
Utils::Vector3i node_grid
double calc_surface_term(bool force_flag, bool energy_flag) override
void dipole_assign() override
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 prefac, int d_rs, CellStructure &cell_structure) const
void operator()(auto &dp3m, double prefac, int d_rs, CellStructure &cell_structure) const