270 bool has_B_field = solver_config.has_field_B_on_rho();
271 bool has_psi_field = solver_config.has_field_psi_on_ch();
272 bool has_epsilon_field = solver_config.dust_config.has_epsilon_field();
273 bool has_deltav_field = solver_config.dust_config.has_deltav_field();
274 bool has_s_j_field = solver_config.dust_config.has_s_j_field();
276 using namespace shamrock::solvergraph;
302 if (has_epsilon_field) {
306 if (has_deltav_field) {
314 u32 ndust = solver_config.dust_config.get_dust_nvar();
321 gpart_mass.data = solver_config.gpart_mass;
331 std::vector<std::shared_ptr<shamrock::solvergraph::INode>> attach_field_sequence;
335 "set_scheduler_patchdata",
338 scheduler().for_each_patchdata_nonempty(
341 scheduler_patchdata.patchdatas.
add_obj(p.id_patch, std::ref(pdat));
346 attach_field_sequence.push_back(set_scheduler_patchdata);
350 auto attach_part_counts
356 attach_field_sequence.push_back(attach_part_counts);
366 attach_field_sequence.push_back(attach_xyz);
376 attach_field_sequence.push_back(attach_vxyz);
386 attach_field_sequence.push_back(attach_axyz);
396 attach_field_sequence.push_back(attach_uint);
406 attach_field_sequence.push_back(attach_duint);
416 attach_field_sequence.push_back(attach_hpart);
426 attach_field_sequence.push_back(attach_B_on_rho);
436 attach_field_sequence.push_back(attach_dB_on_rho);
446 attach_field_sequence.push_back(attach_psi_on_ch);
456 attach_field_sequence.push_back(attach_dpsi_on_ch);
459 if (has_epsilon_field) {
466 attach_field_sequence.push_back(attach_epsilon);
469 if (has_epsilon_field) {
476 attach_field_sequence.push_back(attach_dtepsilon);
479 if (has_deltav_field) {
486 attach_field_sequence.push_back(attach_deltav);
489 if (has_deltav_field) {
496 attach_field_sequence.push_back(attach_dtdeltav);
506 attach_field_sequence.push_back(attach_s_j);
516 attach_field_sequence.push_back(attach_ds_j_dt);
519 "attach fields to scheduler",
529 auto make_half_step_sequence = [&](std::string prefix) {
530 std::vector<std::shared_ptr<shamrock::solvergraph::INode>> half_step_sequence;
541 half_step_sequence.push_back(half_step_vxyz);
553 half_step_sequence.push_back(half_step_uint);
565 half_step_sequence.push_back(half_step_B_on_rho);
577 half_step_sequence.push_back(half_step_psi_on_ch);
580 if (has_epsilon_field) {
589 half_step_sequence.push_back(half_step_epsilon);
592 if (has_deltav_field) {
601 half_step_sequence.push_back(half_step_deltav);
605 u32 ndust = solver_config.dust_config.get_dust_nvar();
607 auto &cfg = solver_config.dust_config.get_monofluid_tva();
609 if (cfg.ensure_s_j_positivity) {
619 half_step_sequence.push_back(half_step_s_j);
629 half_step_sequence.push_back(half_step_s_j);
632 if (cfg.should_clamp_dust_density()) {
633 auto hfactd_edge = IDataEdge<Tscal>::make_shared(
"hfactd",
"hfactd");
634 hfactd_edge->data = Kernel::hfactd;
637 = IDataEdge<Tscal>::make_shared(
"clamp_frac",
"clamp_frac");
638 clamp_frac_edge->data = cfg.get_clamp_dust_frac();
640 auto half_step_s_j_density_clamp = solver_graph.
register_node(
641 prefix +
"_s_j_density_clamp",
651 half_step_sequence.push_back(half_step_s_j_density_clamp);
658 solver_graph.
register_node(
"half_step1", make_half_step_sequence(
"half_step1"));
659 solver_graph.
register_node(
"half_step2", make_half_step_sequence(
"half_step2"));
674 "leapfrog predictor",
676 "leapfrog predictor",
688 bool do_part_killing_step = solver_config.particle_killing.kill_list.size() > 0;
690 if (do_part_killing_step) {
698 std::vector<std::shared_ptr<shamrock::solvergraph::INode>> part_kill_sequence{};
702 auto empty_part_to_remove
705 part_kill_sequence.push_back(empty_part_to_remove);
708 using kill_t =
typename ParticleKillingConfig<Tvec>::kill_t;
712 for (kill_t &kill_obj : solver_config.particle_killing.kill_list) {
713 if (kill_sphere *kill_info = std::get_if<kill_sphere>(&kill_obj)) {
716 kill_info->center, kill_info->radius);
717 node_selector.set_edges(xyz_edge, part_to_remove);
719 part_kill_sequence.push_back(
720 std::make_shared<
decltype(node_selector)>(std::move(node_selector)));
726 node_killer.set_edges(part_to_remove, patchdatas);
728 part_kill_sequence.push_back(
729 std::make_shared<
decltype(node_killer)>(std::move(node_killer)));
734 part_kill_sequence.push_back(solver_graph.
get_node_ptr_base(
"attach fields to scheduler"));
746 half_dt.data = dt.data / 2;
755 = std::make_shared<shamrock::solvergraph::Indexes<u32>>(
"part_counts",
"N_{\\rm part}");
757 storage.part_counts_with_ghost = std::make_shared<shamrock::solvergraph::Indexes<u32>>(
758 "part_counts_with_ghost",
"N_{\\rm part, with ghost}");
760 storage.patch_rank_owner = std::make_shared<shamrock::solvergraph::RankGetter>(
761 [&](
u64 patch_id) ->
u32 {
762 return scheduler().get_patch_rank_owner(patch_id);
768 storage.positions_with_ghosts
769 = std::make_shared<shamrock::solvergraph::FieldRefs<Tvec>>(
"part_pos",
"\\mathbf{r}");
770 storage.hpart_with_ghosts
771 = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"h_part",
"h");
774 = std::make_shared<shammodels::sph::solvergraph::NeighCache>(
"neigh_cache",
"neigh");
776 storage.omega = std::make_shared<shamrock::solvergraph::Field<Tscal>>(1,
"omega",
"\\Omega");
778 if (solver_config.has_field_alphaAV()) {
779 storage.alpha_av_updated = std::make_shared<shamrock::solvergraph::Field<Tscal>>(
780 1,
"alpha_av_updated",
"\\alpha_{\\rm AV}");
783 storage.pressure = std::make_shared<shamrock::solvergraph::Field<Tscal>>(1,
"pressure",
"P");
785 = std::make_shared<shamrock::solvergraph::Field<Tscal>>(1,
"soundspeed",
"c_s");
787 storage.exchange_gz_alpha
788 = std::make_shared<shamrock::solvergraph::ExchangeGhostField<Tscal>>();
789 storage.exchange_gz_node
790 = std::make_shared<shamrock::solvergraph::ExchangeGhostLayer>(storage.ghost_layout);
791 storage.exchange_gz_positions
792 = std::make_shared<shamrock::solvergraph::ExchangeGhostLayer>(storage.xyzh_ghost_layout);
824 auto sink_accelerations
827 "sink_angular_momentum");
831 "sink_accretion_radius");
834 "sink_accretion_table",
Field<u32>(1,
"sink_accretion_table",
"\\mathrm{acc}"));
909 "sink_ext_force_epsilon",
IDataEdge<Tscal>(
"epsilon_grav_sink",
"\\epsilon"));
913 g_edge.data = solver_config.get_constant_G();
919 "set_sink_ext_force_epsilon",
921 epsilon_edge.data = 1e-9;
927 "reset_sink_acc_ext",
930 for (Tvec &a : acc_ext.data) {
949 "sink_ext_force_body",
951 "sink ext force body",
962 "sink ext force",
OperationIf(
"sink ext force", ext_force_body));
974 auto sink_predictor_dt_to_half_dt = solver_graph.
register_node(
975 "sink_predictor_dt_to_half_dt",
978 half_dt.data = dt.data / 2;
1003 "sink_predictor_body",
1005 "sink predictor body",
1010 sink_predictor_dt_to_half_dt,
1011 sink_predictor_vel_update,
1012 sink_predictor_pos_update,
1018 "sink predictor",
OperationIf(
"sink predictor", sink_predictor_body));
1027 auto sink_corrector_vel_update = solver_graph.
register_node(
1039 "sink corrector",
OperationIf(
"sink corrector", sink_corrector_vel_update));
1048 using EF_PointMass =
typename Config::ExtForceConfig::PointMass;
1049 using EF_PN_PW =
typename Config::ExtForceConfig::PN_PW;
1050 using EF_LenseThirring =
typename Config::ExtForceConfig::LenseThirring;
1054 auto collect_accretors = [&]() {
1055 std::vector<Tvec> positions{};
1056 std::vector<Tscal> radii{};
1057 for (
auto &var_force : solver_config.ext_force_config.ext_forces) {
1058 if (EF_PointMass *ext_force = std::get_if<EF_PointMass>(&var_force.val)) {
1059 positions.push_back(ext_force->central_pos);
1060 radii.push_back(ext_force->Racc);
1061 }
else if (EF_PN_PW *ext_force = std::get_if<EF_PN_PW>(&var_force.val)) {
1062 positions.push_back(ext_force->central_pos);
1063 radii.push_back(ext_force->Racc);
1065 EF_LenseThirring *ext_force = std::get_if<EF_LenseThirring>(&var_force.val)) {
1066 positions.push_back(ext_force->central_pos);
1067 radii.push_back(ext_force->Racc);
1070 return std::pair{std::move(positions), std::move(radii)};
1074 "ext_force_accretion_pos",
1075 IDataEdge<std::vector<Tvec>>(
"ext_force_accretion_pos",
"\\mathbf{r}_{\\rm acc, ext}"));
1077 "ext_force_accretion_racc",
1078 IDataEdge<std::vector<Tscal>>(
"ext_force_accretion_racc",
"R_{\\rm acc, ext}"));
1080 "ext_force_accretion_table",
1081 Field<u32>(1,
"ext_force_accretion_table",
"\\mathrm{acc}"));
1083 "has_ext_force_accretion",
1084 IDataEdge<bool>(
"has_ext_force_accretion",
"\\rm has\\_ext\\_force\\_accretion"));
1087 "set_ext_force_accretion_pos",
1089 [collect_accretors](
IDataEdge<std::vector<Tvec>> &accretion_pos) {
1090 accretion_pos.data = std::get<0>(collect_accretors());
1097 "set_ext_force_accretion_racc",
1099 [collect_accretors](
IDataEdge<std::vector<Tscal>> &accretion_racc) {
1100 accretion_racc.data = std::get<1>(collect_accretors());
1104 "ext_force_accretion_racc"));
1107 "set_has_ext_force_accretion",
1111 has_ext_force_accretion.data = !accretion_pos.data.empty();
1138 "ext_force_accretion_body",
1140 "ext force accretion",
1155 "if_has_ext_force_accretion",
1156 OperationIf(
"if_has_ext_force_accretion", accretion_body));
1162 "point mass accretion",
1164 "point mass accretion",
1174 std::vector<std::shared_ptr<shamrock::solvergraph::INode>> seq{};
1183 if (do_part_killing_step) {
1195 if (solver_config.self_grav_config.is_sg_on()) {
1199 auto constant_G = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"",
"");
1203 constant_G.data = solver_config.get_constant_G();
1208 auto field_xyz = shamrock::solvergraph::FieldRefs<Tvec>::make_shared(
"",
"");
1213 scheduler().for_each_patchdata_nonempty(
1215 auto &field = pdat.get_field<Tvec>(ixyz);
1216 field_xyz_refs.
add_obj(p.id_patch, std::ref(field));
1218 field_xyz_edge.set_refs(field_xyz_refs);
1224 auto field_axyz_ext = shamrock::solvergraph::FieldRefs<Tvec>::make_shared(
"",
"");
1230 scheduler().for_each_patchdata_nonempty(
1232 auto &field = pdat.get_field<Tvec>(iaxyz_ext);
1233 field_axyz_ext_refs.
add_obj(p.id_patch, std::ref(field));
1235 field_axyz_ext_edge.set_refs(field_axyz_ext_refs);
1237 set_field_axyz_ext.
set_edges(field_axyz_ext);
1239 auto sizes = shamrock::solvergraph::Indexes<u32>::make_shared(
"",
"");
1244 scheduler().for_each_patchdata_nonempty(
1246 sizes.indexes.add_obj(p.id_patch, pdat.
get_obj_cnt());
1251 auto gpart_mass = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"",
"");
1255 gpart_mass.data = solver_config.gpart_mass;
1260 std::shared_ptr<shamrock::solvergraph::INode> sg_inode = build_self_gravity_node<Tvec>(
1261 solver_config.self_grav_config,
1269 "self gravity sequence",
1273 shambase::to_shared(std::move(set_gpart_mass)),
1274 shambase::to_shared(std::move(set_constant_G)),
1275 shambase::to_shared(std::move(set_field_xyz)),
1276 shambase::to_shared(std::move(set_field_axyz_ext)),
1277 shambase::to_shared(std::move(set_sizes)),
1587 using namespace shamrock::patch;
1592 SPHUtils sph_utils(scheduler());
1600 auto should_set_omega_mask = std::make_shared<shamrock::solvergraph::Field<u32>>(
1601 1,
"should_set_omega_mask",
"should_set_omega_mask");
1604 u32 hstep_max = solver_config.h_max_subcycles_count;
1605 for (; hstep_cnt < hstep_max; hstep_cnt++) {
1607 gen_ghost_handler(time_val + dt);
1615 _h_old = utility.
save_field<Tscal>(ihpart,
"h_old");
1619 if (solver_config.gpart_mass == 0) {
1621 "invalid gpart_mass {}, this configuration can not converge.\n"
1622 "Please set it using either model.set_particle_mass(pmass) or "
1623 "cfg.set_particle_mass(pmass)",
1624 solver_config.gpart_mass));
1628 std::shared_ptr<shamrock::solvergraph::Indexes<u32>> sizes
1629 = std::make_shared<shamrock::solvergraph::Indexes<u32>>(
"",
"");
1631 sizes->indexes.add_obj(p.id_patch, pdat.
get_obj_cnt());
1635 auto &neigh_cache = storage.neigh_cache;
1638 auto &pos_merged = storage.positions_with_ghosts;
1641 std::shared_ptr<shamrock::solvergraph::FieldRefs<Tscal>> hold
1642 = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"",
"");
1645 auto &field = _h_old.get_field(p.id_patch);
1646 hold_refs.add_obj(p.id_patch, std::ref(field));
1648 hold->set_refs(hold_refs);
1651 std::shared_ptr<shamrock::solvergraph::FieldRefs<Tscal>> hnew
1652 = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"",
"");
1655 auto &field = pdat.get_field<Tscal>(ihpart);
1656 hnew_refs.add_obj(p.id_patch, std::ref(field));
1658 hnew->set_refs(hnew_refs);
1661 std::shared_ptr<shamrock::solvergraph::FieldRefs<Tscal>> eps_h
1662 = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"",
"");
1665 auto &field = _epsilon_h.get_field(p.id_patch);
1666 eps_h_refs.add_obj(p.id_patch, std::ref(field));
1668 eps_h->set_refs(eps_h_refs);
1670 std::shared_ptr<shamrock::solvergraph::INode> smth_h_iter_ptr;
1675 if (h_conf_density_based *conf
1676 = std::get_if<h_conf_density_based>(&solver_config.smoothing_length_config.config)) {
1677 std::shared_ptr<shammodels::sph::modules::IterateSmoothingLengthDensity<Tvec, Kernel>>
1678 smth_h_iter = std::make_shared<
1680 solver_config.gpart_mass,
1681 solver_config.htol_up_coarse_cycle,
1682 solver_config.htol_up_fine_cycle,
1683 solver_config.epsilon_h);
1684 smth_h_iter->set_edges(sizes, neigh_cache, pos_merged, hold, hnew, eps_h);
1685 smth_h_iter_ptr = smth_h_iter;
1687 h_conf_neigh_lim *conf
1688 = std::get_if<h_conf_neigh_lim>(&solver_config.smoothing_length_config.config)) {
1691 smth_h_iter_neigh_lim = std::make_shared<
1693 solver_config.gpart_mass,
1694 solver_config.htol_up_coarse_cycle,
1695 solver_config.htol_up_fine_cycle,
1696 conf->max_neigh_count,
1697 solver_config.epsilon_h);
1698 smth_h_iter_neigh_lim->set_edges(
1699 sizes, neigh_cache, pos_merged, hold, hnew, eps_h, should_set_omega_mask);
1700 smth_h_iter_ptr = smth_h_iter_neigh_lim;
1706 std::shared_ptr<shamrock::solvergraph::IDataEdge<bool>> is_converged
1707 = shamrock::solvergraph::IDataEdge<bool>::make_shared(
"",
"");
1710 smth_h_iter_ptr, solver_config.epsilon_h, solver_config.h_iter_per_subcycles,
false);
1711 loop_smth_h_iter.set_edges(eps_h, is_converged);
1715 if (!is_converged->data) {
1717 Tscal largest_h = 0;
1720 largest_h = sham::max(largest_h, pdat.get_field<Tscal>(ihpart).compute_max());
1722 Tscal global_largest_h = shamalgs::collective::allreduce_max(largest_h);
1724 std::string add_info =
"";
1725 u64 cnt_unconverged = 0;
1728 = _epsilon_h.get_field(p.id_patch).get_ids_buf_where([](
auto access,
u32 id) {
1729 return access[id] == -1;
1732 if (hstep_cnt == hstep_max - 1) {
1733 if (std::get<0>(res)) {
1734 add_info +=
"\n patch " + std::to_string(p.id_patch) +
" ";
1735 add_info +=
"errored parts : \n";
1736 sycl::buffer<u32> &idx_err = *std::get<0>(res);
1741 auto pos = xyz.copy_to_stdvec();
1742 auto h = hpart.copy_to_stdvec();
1745 sycl::host_accessor acc{idx_err};
1746 for (
u32 i = 0; i < idx_err.size(); i++) {
1747 add_info += sham::format(
1748 "{} - pos : {}, hpart : {}\n", acc[i], pos[acc[i]], h[acc[i]]);
1754 cnt_unconverged += std::get<1>(res);
1757 u64 global_cnt_unconverged = shamalgs::collective::allreduce_sum(cnt_unconverged);
1762 "smoothing length is not converged, rerunning the iterator ...\n largest h "
1765 "unconverged cnt =",
1766 global_cnt_unconverged,
1770 reset_ghost_handler();
1778 storage.merged_xyzh.reset();
1799 if (hstep_cnt == hstep_max) {
1800 logger::err_ln(
"SPH",
"the h iterator is not converged after", hstep_cnt,
"iterations");
1803 std::shared_ptr<shamrock::solvergraph::FieldRefs<Tscal>> hnew_edge
1804 = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"",
"");
1807 auto &field = pdat.get_field<Tscal>(ihpart);
1808 hnew_refs.add_obj(p.id_patch, std::ref(field));
1810 hnew_edge->set_refs(hnew_refs);
1813 compute_omega.set_edges(
1814 storage.part_counts,
1815 storage.neigh_cache,
1816 storage.positions_with_ghosts,
1821 if (solver_config.smoothing_length_config.is_density_based_neigh_lim()) {
1826 set_omega_mask.set_edges(storage.part_counts, should_set_omega_mask, storage.omega);
1923 timer_interf.
start();
1926 using namespace shamrock::patch;
1928 bool has_alphaAV_field = solver_config.has_field_alphaAV();
1929 bool has_soundspeed_field = solver_config.ghost_has_soundspeed();
1931 bool has_B_field = solver_config.has_field_B_on_rho();
1932 bool has_psi_field = solver_config.has_field_psi_on_ch();
1933 bool has_curlB_field = solver_config.has_field_curlB();
1934 bool has_epsilon_field = solver_config.dust_config.has_epsilon_field();
1935 bool has_deltav_field = solver_config.dust_config.has_deltav_field();
1936 bool has_s_j_field = solver_config.dust_config.has_s_j_field();
1946 const u32 ialpha_AV = (has_alphaAV_field) ? pdl.
get_field_idx<Tscal>(
"alpha_AV") : 0;
1947 const u32 isoundspeed = (has_soundspeed_field) ? pdl.
get_field_idx<Tscal>(
"soundspeed") : 0;
1949 const u32 iB_on_rho = (has_B_field) ? pdl.
get_field_idx<Tvec>(
"B/rho") : 0;
1950 const u32 idB_on_rho = (has_B_field) ? pdl.
get_field_idx<Tvec>(
"dB/rho") : 0;
1951 const u32 ipsi_on_ch = (has_psi_field) ? pdl.
get_field_idx<Tscal>(
"psi/ch") : 0;
1952 const u32 idpsi_on_ch = (has_psi_field) ? pdl.
get_field_idx<Tscal>(
"dpsi/ch") : 0;
1953 const u32 icurlB = (has_curlB_field) ? pdl.
get_field_idx<Tvec>(
"curlB") : 0;
1955 bool do_MHD_debug = solver_config.do_MHD_debug();
1956 const u32 imag_pressure = (do_MHD_debug) ? pdl.
get_field_idx<Tvec>(
"mag_pressure") : -1;
1957 const u32 imag_tension = (do_MHD_debug) ? pdl.
get_field_idx<Tvec>(
"mag_tension") : -1;
1958 const u32 igas_pressure = (do_MHD_debug) ? pdl.
get_field_idx<Tvec>(
"gas_pressure") : -1;
1959 const u32 itensile_corr = (do_MHD_debug) ? pdl.
get_field_idx<Tvec>(
"tensile_corr") : -1;
1960 const u32 ipsi_propag = (do_MHD_debug) ? pdl.
get_field_idx<Tscal>(
"psi_propag") : -1;
1961 const u32 ipsi_diff = (do_MHD_debug) ? pdl.
get_field_idx<Tscal>(
"psi_diff") : -1;
1962 const u32 ipsi_cons = (do_MHD_debug) ? pdl.
get_field_idx<Tscal>(
"psi_cons") : -1;
1963 const u32 iu_mhd = (do_MHD_debug) ? pdl.
get_field_idx<Tscal>(
"u_mhd") : -1;
1965 const u32 iepsilon = (has_epsilon_field) ? pdl.
get_field_idx<Tscal>(
"epsilon") : 0;
1966 const u32 ideltav = (has_deltav_field) ? pdl.
get_field_idx<Tvec>(
"deltav") : 0;
1969 auto &ghost_layout_ptr = storage.ghost_layout;
1976 const u32 iaxyz_interf
1977 = (solver_config.has_axyz_in_ghost()) ? ghost_layout.
get_field_idx<Tvec>(
"axyz") : 0;
1979 const u32 isoundspeed_interf
1980 = (has_soundspeed_field) ? ghost_layout.
get_field_idx<Tscal>(
"soundspeed") : 0;
1982 const u32 iB_interf = (has_B_field) ? ghost_layout.
get_field_idx<Tvec>(
"B/rho") : 0;
1983 const u32 ipsi_interf = (has_psi_field) ? ghost_layout.
get_field_idx<Tscal>(
"psi/ch") : 0;
1984 const u32 icurlB_interf = (has_curlB_field) ? ghost_layout.
get_field_idx<Tvec>(
"curlB") : 0;
1986 const u32 iepsilon_interf
1987 = (has_epsilon_field) ? ghost_layout.
get_field_idx<Tscal>(
"epsilon") : 0;
1988 const u32 ideltav_interf = (has_deltav_field) ? ghost_layout.
get_field_idx<Tvec>(
"deltav") : 0;
1989 const u32 is_j_interf = (has_s_j_field) ? ghost_layout.
get_field_idx<Tscal>(
"s_j") : 0;
1996 auto pdat_interf = ghost_handle.template build_interface_native<PatchDataLayer>(
1997 storage.ghost_patch_cache.get(),
1999 PatchDataLayer pdat(ghost_layout_ptr);
2006 ghost_handle.template modify_interface_native<PatchDataLayer>(
2007 storage.ghost_patch_cache.get(),
2011 InterfaceBuildInfos binfo,
2015 PatchDataLayer &sender_patch = scheduler().patch_data.get_pdat(sender);
2016 PatchDataField<Tscal> &sender_omega = omega.get(sender);
2018 sender_patch.get_field<Tscal>(ihpart).append_subset_to(
2019 buf_idx, cnt, pdat.get_field<Tscal>(ihpart_interf));
2020 sender_patch.get_field<Tscal>(iuint).append_subset_to(
2021 buf_idx, cnt, pdat.get_field<Tscal>(iuint_interf));
2023 if (solver_config.has_axyz_in_ghost()) {
2024 sender_patch.get_field<Tvec>(iaxyz).append_subset_to(
2025 buf_idx, cnt, pdat.get_field<Tvec>(iaxyz_interf));
2028 sender_patch.get_field<Tvec>(ivxyz).append_subset_to(
2029 buf_idx, cnt, pdat.get_field<Tvec>(ivxyz_interf));
2031 sender_omega.append_subset_to(buf_idx, cnt, pdat.get_field<Tscal>(iomega_interf));
2033 if (has_soundspeed_field) {
2034 sender_patch.get_field<Tscal>(isoundspeed)
2035 .append_subset_to(buf_idx, cnt, pdat.get_field<Tscal>(isoundspeed_interf));
2039 sender_patch.get_field<Tvec>(iB_on_rho).append_subset_to(
2040 buf_idx, cnt, pdat.get_field<Tvec>(iB_interf));
2043 if (has_psi_field) {
2044 sender_patch.get_field<Tscal>(ipsi_on_ch)
2045 .append_subset_to(buf_idx, cnt, pdat.get_field<Tscal>(ipsi_interf));
2048 if (has_curlB_field) {
2049 sender_patch.get_field<Tvec>(icurlB).append_subset_to(
2050 buf_idx, cnt, pdat.get_field<Tvec>(icurlB_interf));
2053 if (has_epsilon_field) {
2054 sender_patch.get_field<Tscal>(iepsilon).append_subset_to(
2055 buf_idx, cnt, pdat.get_field<Tscal>(iepsilon_interf));
2058 if (has_deltav_field) {
2059 sender_patch.get_field<Tvec>(ideltav).append_subset_to(
2060 buf_idx, cnt, pdat.get_field<Tvec>(ideltav_interf));
2063 if (has_s_j_field) {
2064 sender_patch.get_field<Tscal>(is_j).append_subset_to(
2065 buf_idx, cnt, pdat.get_field<Tscal>(is_j_interf));
2069 ghost_handle.template modify_interface_native<PatchDataLayer>(
2070 storage.ghost_patch_cache.get(),
2074 InterfaceBuildInfos binfo,
2078 if (sycl::length(binfo.offset_speed) > 0) {
2079 pdat.get_field<Tvec>(ivxyz_interf).apply_offset(binfo.offset_speed);
2085 std::move(pdat_interf),
2086 storage.exchange_gz_node,
2087 solver_config.show_ghost_zone_graph);
2089 std::map<u64, u64> sz_interf_map;
2094 storage.merged_patchdata_ghost.set(
2095 ghost_handle.template merge_native<PatchDataLayer, PatchDataLayer>(
2096 std::move(interf_pdat),
2098 PatchDataLayer pdat_new(ghost_layout_ptr);
2100 u32 or_elem = pdat.get_obj_cnt();
2101 pdat_new.reserve(or_elem + sz_interf_map[p.id_patch]);
2102 u32 total_elements = or_elem;
2104 PatchDataField<Tscal> &cur_omega = omega.get(p.id_patch);
2106 pdat_new.get_field<Tscal>(ihpart_interf).insert(pdat.get_field<Tscal>(ihpart));
2107 pdat_new.get_field<Tscal>(iuint_interf).insert(pdat.get_field<Tscal>(iuint));
2108 pdat_new.get_field<Tvec>(ivxyz_interf).insert(pdat.get_field<Tvec>(ivxyz));
2110 if (solver_config.has_axyz_in_ghost()) {
2111 pdat_new.get_field<Tvec>(iaxyz_interf).insert(pdat.get_field<Tvec>(iaxyz));
2114 pdat_new.get_field<Tscal>(iomega_interf).insert(cur_omega);
2116 if (has_soundspeed_field) {
2117 pdat_new.get_field<Tscal>(isoundspeed_interf)
2118 .insert(pdat.get_field<Tscal>(isoundspeed));
2122 pdat_new.get_field<Tvec>(iB_interf).insert(pdat.get_field<Tvec>(iB_on_rho));
2125 if (has_psi_field) {
2126 pdat_new.get_field<Tscal>(ipsi_interf)
2127 .insert(pdat.get_field<Tscal>(ipsi_on_ch));
2130 if (has_curlB_field) {
2131 pdat_new.get_field<Tvec>(icurlB_interf).insert(pdat.get_field<Tvec>(icurlB));
2134 if (has_epsilon_field) {
2135 pdat_new.get_field<Tscal>(iepsilon_interf)
2136 .insert(pdat.get_field<Tscal>(iepsilon));
2139 if (has_deltav_field) {
2140 pdat_new.get_field<Tvec>(ideltav_interf).insert(pdat.get_field<Tvec>(ideltav));
2143 if (has_s_j_field) {
2144 pdat_new.get_field<Tscal>(is_j_interf).insert(pdat.get_field<Tscal>(is_j));
2147 pdat_new.check_field_obj_cnt_match();
2152 pdat.insert_elements(pdat_interf);
2155 timer_interf.stop();
2156 storage.timings_details.interface += timer_interf.elapsed_sec();
2474 for (
auto &callbacks : timestep_callbacks) {
2475 if (callbacks.step_begin_callback) {
2480 Tscal t_current = get_time();
2481 Tscal dt = get_dt_sph();
2487 sham::format(
"---------------- t = {}, dt = {} ----------------", t_current, dt));
2495 .update_load_balancing();
2496 scheduler().scheduler_step(
true,
true);
2498 .update_load_balancing();
2500 scheduler().scheduler_step(
false,
false);
2506 using namespace shamrock::patch;
2508 bool has_B_field = solver_config.has_field_B_on_rho();
2509 bool has_psi_field = solver_config.has_field_psi_on_ch();
2510 bool has_epsilon_field = solver_config.dust_config.has_epsilon_field();
2511 bool has_deltav_field = solver_config.dust_config.has_deltav_field();
2512 bool has_s_j_field = solver_config.dust_config.has_s_j_field();
2522 const u32 iB_on_rho = (has_B_field) ? pdl.
get_field_idx<Tvec>(
"B/rho") : 0;
2523 const u32 idB_on_rho = (has_B_field) ? pdl.
get_field_idx<Tvec>(
"dB/rho") : 0;
2524 const u32 ipsi_on_ch = (has_psi_field) ? pdl.
get_field_idx<Tscal>(
"psi/ch") : 0;
2525 const u32 idpsi_on_ch = (has_psi_field) ? pdl.
get_field_idx<Tscal>(
"dpsi/ch") : 0;
2526 const u32 iepsilon = (has_epsilon_field) ? pdl.
get_field_idx<Tscal>(
"epsilon") : 0;
2527 const u32 idtepsilon = (has_epsilon_field) ? pdl.
get_field_idx<Tscal>(
"dtepsilon") : 0;
2529 const u32 ids_j_dt = (has_s_j_field) ? pdl.
get_field_idx<Tscal>(
"ds_j_dt") : 0;
2530 const u32 ideltav = (has_deltav_field) ? pdl.
get_field_idx<Tvec>(
"deltav") : 0;
2531 const u32 idtdeltav = (has_deltav_field) ? pdl.
get_field_idx<Tvec>(
"dtdeltav") : 0;
2538 using namespace shamrock::solvergraph;
2552 gen_serial_patch_tree();
2556 u64 Npart_all = scheduler().get_total_obj_count();
2558 if (solver_config.enable_particle_reordering
2559 && solve_logs.step_count % solver_config.particle_reordering_step_freq == 0) {
2560 logger::info_ln(
"SPH",
"Reordering particles at step ", solve_logs.step_count);
2568 using namespace shamrock::solvergraph;
2579 if (solver_config.self_grav_config.is_sg_on()) {
2580 using namespace shamrock::solvergraph;
2586 auto &merged_xyzh = storage.merged_xyzh.get();
2596 u32 iB_on_rho_interf = (has_B_field) ? ghost_layout.
get_field_idx<Tvec>(
"B/rho") : 0;
2597 u32 ipsi_on_rho_interf = (has_psi_field) ? ghost_layout.
get_field_idx<Tscal>(
"psi/ch") : 0;
2604 u32 corrector_iter_cnt = 0;
2605 bool need_rerun_corrector =
false;
2611 if (corrector_iter_cnt == 50) {
2613 "the corrector has made over 50 loops, either their is a bug, either you are using "
2614 "a dt that is too large");
2620 if (solver_config.has_field_alphaAV()) {
2622 std::shared_ptr<shamrock::solvergraph::PatchDataLayerRefs> patchdatas
2623 = std::make_shared<shamrock::solvergraph::PatchDataLayerRefs>(
2624 "patchdata_layer_ref",
"patchdata_layer_ref");
2626 auto node_set_edge = scheduler().get_node_set_edge_patchdata_layer_refs();
2627 node_set_edge->set_edges(patchdatas);
2628 node_set_edge->evaluate();
2631 scheduler().get_layout_ptr_old(),
"alpha_AV");
2632 node_copy.set_edges(patchdatas, storage.alpha_av_updated);
2636 if (solver_config.has_field_dtdivv()) {
2638 if (solver_config.combined_dtdiv_divcurlv_compute) {
2639 if (solver_config.has_field_dtdivv()) {
2641 .update_dtdivv(
true);
2645 if (solver_config.has_field_divv()) {
2650 if (solver_config.has_field_curlv()) {
2655 if (solver_config.has_field_dtdivv()) {
2657 .update_dtdivv(
false);
2662 if (solver_config.has_field_divv()) {
2667 if (solver_config.has_field_curlv()) {
2684 if (solver_config.has_field_alphaAV()) {
2689 using InterfaceBuildInfos =
2693 time_interf.
start();
2695 auto field_interf = ghost_handle.template build_interface_native<PatchDataField<Tscal>>(
2696 storage.ghost_patch_cache.get(),
2699 InterfaceBuildInfos binfo,
2704 return sender_field.make_new_from_subset(buf_idx, cnt);
2708 = ghost_handle.communicate_pdatfield(
2709 std::move(field_interf), 1, storage.exchange_gz_alpha);
2713 std::move(interf_pdat),
2716 = comp_field_send.
get_field(p.id_patch);
2717 return receiver_field.duplicate();
2720 mpdat.insert(pdat_interf);
2724 storage.timings_details.interface += time_interf.
elapsed_sec();
2726 storage.alpha_av_ghost.set(std::move(merged_field));
2732 constexpr bool debug_interfaces =
false;
2733 if constexpr (debug_interfaces) {
2735 if (solver_config.do_debug_dump) {
2738 = storage.merged_patchdata_ghost.get();
2745 merged_xyzh.get(cur_p.
id_patch).field_pos.get_buf());
2746 sycl::buffer<Tvec> &buf_vxyz = mpdat.get_field_buf_ref<Tvec>(ivxyz_interf);
2747 sycl::buffer<Tscal> &buf_hpart = mpdat.get_field_buf_ref<Tscal>(ihpart_interf);
2751 SHAM_ASSERT(merged_patch.total_elements == total_elements);
2755 solver_config.gpart_mass,
2761 make_interface_debug_phantom_dump(info).gen_file().write_to_file(
2762 solver_config.debug_dump_filename);
2769 shamlog_debug_ln(
"sph::BasicGas",
"compute force");
2779 bool compute_GW = solver_config.compute_gw;
2782 using namespace shamrock::solvergraph;
2785 auto central_pos = IDataEdge<Tvec>::make_shared(
"x_0",
"\\mathbf{x}_0");
2786 central_pos->data = Tvec{0, 0, 0};
2788 auto central_vel = IDataEdge<Tvec>::make_shared(
"v_0",
"\\mathbf{v}_0");
2789 central_vel->data = Tvec{0, 0, 0};
2791 auto central_acc = IDataEdge<Tvec>::make_shared(
"a_0",
"\\mathbf{a}_0");
2792 central_acc->data = Tvec{0, 0, 0};
2794 auto gw_prefactor = IDataEdge<Tscal>::make_shared(
"gw_prefactor",
"gw_prefactor");
2795 gw_prefactor->data = Tscal(1);
2797 auto theta_gw = IDataEdge<Tscal>::make_shared(
"theta_gw",
"\\theta_{\\rm gw}");
2798 theta_gw->data = Tscal(0);
2800 auto phi_gw = IDataEdge<Tscal>::make_shared(
"phi_gw",
"\\phi_{\\rm gw}");
2801 phi_gw->data = Tscal(0);
2806 auto spans_masses = std::make_shared<FieldRefs<Tscal>>(
"m",
"m");
2807 map_field_refs_ext(scheduler(), gw_mass_field, *spans_masses);
2810 auto spans_accel_ext
2811 = std::make_shared<FieldRefs<Tvec>>(
"axyz_ext",
"\\mathbf{a}_{\\rm ext}");
2812 map_field_refs(scheduler(), iaxyz_ext, *spans_accel_ext);
2814 auto ddq = IDataEdge<typename GW::Tddq>::make_shared(
"ddq",
"\\ddot{Q}");
2815 auto ddq_xy = IDataEdge<typename GW::Tddqxy>::make_shared(
"ddq_xy",
"\\ddot{Q}_{xy}");
2816 auto hx = IDataEdge<typename GW::Th>::make_shared(
"hx",
"h_x");
2817 auto hp = IDataEdge<typename GW::Th>::make_shared(
"hp",
"h_+");
2819 GW node_computeGW{};
2820 node_computeGW.set_edges(
2832 storage.part_counts,
2838 node_computeGW.evaluate();
2847 bool has_luminosity = solver_config.compute_luminosity;
2849 if (has_luminosity) {
2853 .set_refs(storage.merged_xyzh.get()
2856 return std::ref(mpdat.get_field<Tscal>(
2861 .set_refs(storage.merged_xyzh.get()
2864 return std::ref(mpdat.get_field<Tscal>(1));
2867 auto uint_with_ghost = shamrock::solvergraph::FieldRefs<Tscal>::make_shared(
"",
"");
2869 set_uint_with_ghost_refs(
2872 = storage.merged_patchdata_ghost.
get();
2877 scheduler().for_each_patchdata_nonempty(
2881 auto &field = mpdat.get_field<Tscal>(iuint_interf);
2882 field_uint_with_ghost_refs.
add_obj(p.id_patch, std::ref(field));
2885 field_uint_with_ghost_edge.set_refs(field_uint_with_ghost_refs);
2888 set_uint_with_ghost_refs.
set_edges(uint_with_ghost);
2890 auto omega_with_ghost = shamrock::solvergraph::FieldRefs<Tscal>::make_shared(
"",
"");
2893 &field_omega_with_ghost_edge) {
2895 = storage.merged_patchdata_ghost.
get();
2900 scheduler().for_each_patchdata_nonempty(
2904 auto &field = mpdat.get_field<Tscal>(iomega_interf);
2905 field_omega_with_ghost_refs.
add_obj(p.id_patch, std::ref(field));
2908 field_omega_with_ghost_edge.set_refs(field_omega_with_ghost_refs);
2911 set_omega_with_ghost_refs.
set_edges(omega_with_ghost);
2913 auto luminosity = shamrock::solvergraph::FieldRefs<Tscal>::make_shared(
"",
"");
2916 set_luminosity_refs(
2919 = storage.merged_patchdata_ghost.
get();
2924 scheduler().for_each_patchdata_nonempty(
2926 auto &field = pdat.get_field<Tscal>(iluminosity);
2927 field_luminosity_refs.
add_obj(p.id_patch, std::ref(field));
2929 field_luminosity_edge.set_refs(field_luminosity_refs);
2932 set_luminosity_refs.
set_edges(luminosity);
2934 set_uint_with_ghost_refs.evaluate();
2935 set_omega_with_ghost_refs.evaluate();
2936 set_luminosity_refs.evaluate();
2938 Tscal alpha_u = solver_config.artif_viscosity.get_alpha_u().value();
2941 solver_config.gpart_mass, alpha_u};
2943 compute_luminosity.set_edges(
2944 storage.part_counts,
2945 storage.part_counts_with_ghost,
2946 storage.neigh_cache,
2947 storage.positions_with_ghosts,
2948 storage.hpart_with_ghosts,
2966 shamlog_debug_ln(
"sph::BasicGas",
"leapfrog corrector");
2967 utility.fields_leapfrog_corrector<Tvec>(
2968 ivxyz, iaxyz, storage.old_axyz.get(), vepsilon_v_sq, dt / 2);
2969 utility.fields_leapfrog_corrector<Tscal>(
2970 iuint, iduint, storage.old_duint.get(), uepsilon_u_sq, dt / 2);
2972 if (solver_config.has_field_B_on_rho()) {
2975 utility.fields_leapfrog_corrector<Tvec>(
2976 iB_on_rho, idB_on_rho, storage.old_dB_on_rho.get(), BOR_epsilon_BOR_sq, dt / 2);
2978 if (solver_config.has_field_B_on_rho()) {
2981 utility.fields_leapfrog_corrector<Tscal>(
2982 ipsi_on_ch, idpsi_on_ch, storage.old_dpsi_on_ch.get(), POC_epsilon_POC_sq, dt / 2);
2985 if (solver_config.dust_config.has_epsilon_field()) {
2988 utility.fields_leapfrog_corrector<Tscal>(
2989 iepsilon, idtepsilon, storage.old_dtepsilon.get(), epsilon_epsilon_sq, dt / 2);
2992 if (solver_config.dust_config.has_deltav_field()) {
2995 utility.fields_leapfrog_corrector<Tvec>(
2996 ideltav, idtdeltav, storage.old_dtdeltav.get(), epsilon_deltav_sq, dt / 2);
2999 if (solver_config.dust_config.has_s_j_field()) {
3001 "s_j s_j^2", solver_config.dust_config.get_dust_nvar());
3002 bool ensure_positivity
3003 = solver_config.dust_config.get_monofluid_tva().ensure_s_j_positivity;
3004 if (ensure_positivity) {
3005 utility.fields_leapfrog_corrector_positive_only<Tscal>(
3006 is_j, ids_j_dt, storage.old_ds_j_dt.get(), s_j_s_j_sq, dt / 2);
3008 utility.fields_leapfrog_corrector<Tscal>(
3009 is_j, ids_j_dt, storage.old_ds_j_dt.get(), s_j_s_j_sq, dt / 2);
3012 auto &monofluid_tva_cfg = solver_config.dust_config.get_monofluid_tva();
3013 if (monofluid_tva_cfg.should_clamp_dust_density()) {
3015 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"hfactd",
"hfactd");
3016 hfactd_edge->data = Kernel::hfactd;
3018 auto clamp_frac_edge = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
3019 "clamp_frac",
"clamp_frac");
3020 clamp_frac_edge->data = monofluid_tva_cfg.get_clamp_dust_frac();
3023 solver_config.dust_config.get_dust_nvar());
3024 density_clamp.set_edges(
3027 storage.solver_graph
3032 storage.solver_graph
3034 storage.solver_graph
3040 storage.old_axyz.reset();
3041 storage.old_duint.reset();
3042 if (solver_config.has_field_B_on_rho()) {
3043 storage.old_dB_on_rho.reset();
3045 if (solver_config.has_field_B_on_rho()) {
3046 storage.old_dpsi_on_ch.reset();
3049 if (solver_config.dust_config.has_epsilon_field()) {
3050 storage.old_dtepsilon.reset();
3053 if (solver_config.dust_config.has_deltav_field()) {
3054 storage.old_dtdeltav.reset();
3057 if (solver_config.dust_config.has_s_j_field()) {
3058 storage.old_ds_j_dt.reset();
3061 Tscal rank_veps_v = sycl::sqrt(vepsilon_v_sq.compute_rank_max());
3066 Tscal sum_vsq = utility.compute_rank_dot_sum<Tvec>(ivxyz);
3068 Tscal vmean_sq = shamalgs::collective::allreduce_sum(sum_vsq) / Tscal(Npart_all);
3070 Tscal vmean = sycl::sqrt(vmean_sq);
3072 Tscal rank_eps_v = rank_veps_v / vmean;
3078 Tscal eps_v = shamalgs::collective::allreduce_max(rank_eps_v);
3080 shamlog_debug_ln(
"BasicGas",
"epsilon v :", eps_v);
3087 "the corrector tolerance are broken the step will "
3088 "be re rerunned\n eps_v = {}",
3091 need_rerun_corrector =
true;
3092 set_cfl_multipler(get_cfl_multipler() / 2);
3096 need_rerun_corrector =
false;
3099 if (!need_rerun_corrector) {
3101 storage.solver_graph.get_node_ref_base(
"sink corrector").evaluate();
3104 if (solver_config.has_field_alphaAV()) {
3112 = pdat.get_field<Tscal>(ialpha_AV).get_buf();
3114 = alpha_av_updated.get_field(cur_p.
id_patch).get_buf();
3116 auto &q = shamsys::instance::get_compute_scheduler().
get_queue();
3120 auto alpha_av_updated = buf_alpha_av_updated.
get_read_access(depends_list);
3122 auto e = q.submit(depends_list, [&](sycl::handler &cgh) {
3123 shambase::parallel_for(
3125 alpha_av[id_a] = alpha_av_updated[id_a];
3134 shamlog_debug_ln(
"BasicGas",
"computing next CFL");
3138 = storage.merged_xyzh.get().template map<u32>(
3140 return scheduler().patch_data.get_pdat(
id).get_obj_cnt();
3143 std::shared_ptr<shamrock::solvergraph::Field<Tscal>> vsig_max_dt
3144 = std::make_shared<shamrock::solvergraph::Field<Tscal>>(
3145 1,
"vsig_a",
"v_{\\rm sig}");
3148 std::shared_ptr<shamrock::solvergraph::Field<Tscal>> vclean_dt;
3149 if (has_psi_field) {
3150 vclean_dt = std::make_shared<shamrock::solvergraph::Field<Tscal>>(
3151 1,
"vclean_a",
"v_{\\rm clean}");
3156 = storage.merged_patchdata_ghost.
get();
3162 = merged_xyzh.get(cur_p.
id_patch).template get_field_buf_ref<Tvec>(0);
3165 = mpdat.get_field_buf_ref<Tscal>(ihpart_interf);
3184 auto &q = shamsys::instance::get_compute_scheduler().get_queue();
3189 auto hpart = buf_hpart.get_read_access(depends_list);
3190 auto u = buf_uint.get_read_access(depends_list);
3191 auto pressure = buf_pressure.get_read_access(depends_list);
3194 auto particle_looper_ptrs = pcache.get_read_access(depends_list);
3197 auto e = q.submit(depends_list, [&](sycl::handler &cgh) {
3198 const Tscal pmass = solver_config.gpart_mass;
3199 const Tscal alpha_u = 1.0;
3200 const Tscal alpha_AV = 1.0;
3201 const Tscal beta_AV = 2.0;
3205 constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
3207 shambase::parallel_for(
3209 using namespace shamrock::sph;
3211 Tvec sum_axyz = {0, 0, 0};
3213 Tscal h_a = hpart[id_a];
3215 Tvec xyz_a = xyz[id_a];
3216 Tvec vxyz_a = vxyz[id_a];
3218 Tscal rho_a = rho_h(pmass, h_a, Kernel::hfactd);
3219 Tscal rho_a_sq = rho_a * rho_a;
3220 Tscal rho_a_inv = 1. / rho_a;
3222 Tscal P_a = pressure[id_a];
3224 const Tscal u_a = u[id_a];
3226 Tscal cs_a = cs[id_a];
3230 particle_looper.for_each_object(id_a, [&](
u32 id_b) {
3232 Tvec dr = xyz_a - xyz[id_b];
3233 Tscal rab2 = sycl::dot(dr, dr);
3234 Tscal h_b = hpart[id_b];
3236 if (rab2 > h_a * h_a * Rker2 && rab2 > h_b * h_b * Rker2) {
3240 Tscal rab = sycl::sqrt(rab2);
3241 Tvec vxyz_b = vxyz[id_b];
3242 Tvec v_ab = vxyz_a - vxyz_b;
3243 const Tscal u_b = u[id_b];
3245 Tvec r_ab_unit = dr / rab;
3248 r_ab_unit = {0, 0, 0};
3251 Tscal rho_b = rho_h(pmass, h_b, Kernel::hfactd);
3252 Tscal P_b = pressure[id_b];
3253 Tscal cs_b = cs[id_b];
3254 Tscal v_ab_r_ab = sycl::dot(v_ab, r_ab_unit);
3255 Tscal abs_v_ab_r_ab = sycl::fabs(v_ab_r_ab);
3260 const Tscal alpha_a = alpha_AV;
3261 const Tscal alpha_b = alpha_AV;
3263 Tscal vsig_a = alpha_a * cs_a + beta_AV * abs_v_ab_r_ab;
3265 vsig_max = sycl::fmax(vsig_max, vsig_a);
3268 vsig[id_a] = vsig_max;
3272 if (has_psi_field) {
3274 Tscal
const mu_0 = solver_config.get_constant_mu_0();
3277 Tvec *B_on_rho = mpdat.get_field_buf_ref<Tvec>(iB_on_rho_interf)
3278 .get_write_access(depends_list);
3282 auto e = q.submit(depends_list, [&](sycl::handler &cgh) {
3283 const Tscal pmass = solver_config.gpart_mass;
3287 constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
3289 shambase::parallel_for(
3291 using namespace shamrock::sph;
3293 Tscal h_a = hpart[id_a];
3294 Tscal rho_a = rho_h(pmass, h_a, Kernel::hfactd);
3295 const Tscal u_a = u[id_a];
3296 Tscal cs_a = cs[id_a];
3297 Tvec B_a = B_on_rho[id_a] * rho_a;
3299 Tscal vclean_a = shamphys::MHD_physics<Tvec, Tscal>::v_shock(
3300 cs_a, B_a, rho_a, mu_0);
3302 vclean[id_a] = vclean_a;
3305 mpdat.get_field_buf_ref<Tvec>(iB_on_rho_interf).complete_event_state(e);
3311 buf_hpart.complete_event_state(e);
3312 buf_uint.complete_event_state(e);
3313 buf_pressure.complete_event_state(e);
3319 pcache.complete_event_state(resulting_events);
3323 std::shared_ptr<shamrock::solvergraph::Field<Tscal>> cfl_dt
3324 = std::make_shared<shamrock::solvergraph::Field<Tscal>>(
3325 1,
"cfl_dt",
"\\Delta t_{cfl}");
3328 std::shared_ptr<shamrock::solvergraph::FieldRefs<Tvec>> axyz_refs
3329 = std::make_shared<shamrock::solvergraph::FieldRefs<Tvec>>(
"axyz",
"\\mathbf{a}");
3330 std::shared_ptr<shamrock::solvergraph::FieldRefs<Tscal>> hpart_refs
3331 = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"hpart",
"h");
3333 map_field_refs(scheduler(), iaxyz, *axyz_refs);
3334 map_field_refs_ext(scheduler(), mpdats, ihpart_interf, *hpart_refs);
3336 auto &q = shamsys::instance::get_compute_scheduler().get_queue();
3338 auto reset_dt_part_field = [&]() {
3339 if (solver_config.should_save_dt_to_fields()) {
3343 = pdat.get_field_buf_ref<Tscal>(idt_part);
3344 buf_dt_part.
fill(shambase::get_infty<Tscal>());
3349 auto save_dt_min_to_dt_part = [&]() {
3350 if (solver_config.should_save_dt_to_fields()) {
3354 = pdat.get_field_buf_ref<Tscal>(idt_part);
3362 [](
u32 id_a,
const Tscal *dt, Tscal *dt_part) {
3363 dt_part[id_a] = sycl::min(dt_part[id_a], dt[id_a]);
3370 auto reset_cfl_dt = [&]() {
3372 cfl_dt->get_buf(cur_p.
id_patch).fill(shambase::get_infty<Tscal>());
3376 Tscal C_cour = solver_config.cfl_config.cfl_cour * get_cfl_multipler();
3377 Tscal C_force = solver_config.cfl_config.cfl_force * get_cfl_multipler();
3378 Tscal eta_phi = solver_config.cfl_config.eta_sink;
3380 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> C_cour_edge
3381 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"C_cour",
"C_{cour}");
3382 C_cour_edge->data = C_cour;
3383 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> C_force_edge
3384 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"C_force",
"C_{force}");
3385 C_force_edge->data = C_force;
3386 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> eta_phi_edge
3387 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"eta_phi",
"\\eta_{\\phi}");
3388 eta_phi_edge->data = eta_phi;
3390 std::shared_ptr<ComputeCFLCourant<Tscal>> compute_cfl_courant
3391 = std::make_shared<ComputeCFLCourant<Tscal>>();
3392 compute_cfl_courant->set_edges(
3393 storage.part_counts, C_cour_edge, hpart_refs, vsig_max_dt, cfl_dt);
3395 std::shared_ptr<ComputeCFLForce<Tvec>> compute_cfl_force
3396 = std::make_shared<ComputeCFLForce<Tvec>>();
3397 compute_cfl_force->set_edges(
3398 storage.part_counts, C_force_edge, hpart_refs, axyz_refs, cfl_dt);
3400 std::shared_ptr<ComputeCFLDivBCleaning<Tscal>> compute_cfl_divB_cleaning;
3401 if (has_psi_field) {
3402 compute_cfl_divB_cleaning = std::make_shared<ComputeCFLDivBCleaning<Tscal>>();
3403 compute_cfl_divB_cleaning->set_edges(
3404 storage.part_counts, C_cour_edge, hpart_refs, vclean_dt, cfl_dt);
3407 std::shared_ptr<ComputeCFLDust1Fluid<Tvec>> compute_cfl_dust1_fluid;
3408 std::shared_ptr<shamrock::solvergraph::FieldRefs<Tscal>> s_j_refs;
3409 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> hfactd_edge;
3411 if (solver_config.dust_config.has_s_j_field()) {
3412 u32 ndust = solver_config.dust_config.get_dust_nvar();
3414 compute_cfl_dust1_fluid = std::make_shared<ComputeCFLDust1Fluid<Tvec>>(ndust);
3417 = storage.solver_graph
3418 .template get_edge_ptr<shamrock::solvergraph::Field<Tscal>>(
"Ts_j");
3421 = storage.solver_graph
3422 .template get_edge_ptr<shamrock::solvergraph::IDataEdge<Tscal>>(
3425 s_j_refs = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"s_j",
"s_j");
3428 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"hfactd",
"hfactd");
3429 hfactd_edge->data = Kernel::hfactd;
3431 map_field_refs(scheduler(), is_j, *s_j_refs);
3433 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> C_1fluid_edge
3434 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
3435 "C_1fluid",
"C_{1fluid}");
3437 = solver_config.dust_config.get_monofluid_tva().C_1_fluid * get_cfl_multipler();
3439 compute_cfl_dust1_fluid->set_edges(
3440 storage.part_counts,
3451 std::shared_ptr<ComputeCFLDustDrift<Tvec>> compute_cfl_dust_drift;
3452 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> C_drift_edge;
3453 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> cfl_density_threshold_edge;
3454 std::shared_ptr<shamrock::solvergraph::FieldRefs<Tvec>> delta_v_refs;
3456 if (solver_config.dust_config.has_s_j_field()) {
3457 u32 ndust = solver_config.dust_config.get_dust_nvar();
3459 compute_cfl_dust_drift = std::make_shared<ComputeCFLDustDrift<Tvec>>(ndust);
3461 delta_v_refs = std::make_shared<shamrock::solvergraph::FieldRefs<Tvec>>(
3462 "delta_v",
"delta_v");
3464 map_field_refs(scheduler(), idelta_v, *delta_v_refs);
3466 auto &cfg_monofluid_tva = solver_config.dust_config.get_monofluid_tva();
3469 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"C_drift",
"C_{drift}");
3470 C_drift_edge->data = cfg_monofluid_tva.C_drift * get_cfl_multipler();
3472 cfl_density_threshold_edge = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
3473 "cfl_density_threshold",
"cfl_density_threshold");
3474 cfl_density_threshold_edge->data = cfg_monofluid_tva.cfl_density_threshold;
3477 = storage.solver_graph
3478 .template get_edge_ptr<shamrock::solvergraph::IDataEdge<Tscal>>(
3481 compute_cfl_dust_drift->set_edges(
3482 storage.part_counts,
3484 cfl_density_threshold_edge,
3493 bool show_cfl_detail = solver_config.show_cfl_detail;
3494 std::vector<std::pair<std::string, Tscal>> cfl_detail;
3496 auto save_cfl_detail = [&](
const char *key) {
3497 if (show_cfl_detail) {
3498 save_dt_min_to_dt_part();
3499 cfl_detail.push_back(
3500 {std::string(key), cfl_dt->get_native().compute_rank_min()});
3505 reset_dt_part_field();
3508 compute_cfl_courant->evaluate();
3509 save_cfl_detail(
"courant");
3511 compute_cfl_force->evaluate();
3512 save_cfl_detail(
"force");
3514 if (has_psi_field) {
3515 compute_cfl_divB_cleaning->evaluate();
3516 save_cfl_detail(
"divB_cleaning");
3519 if (solver_config.dust_config.has_s_j_field()) {
3520 compute_cfl_dust1_fluid->evaluate();
3521 save_cfl_detail(
"dust1_fluid");
3523 compute_cfl_dust_drift->evaluate();
3524 save_cfl_detail(
"dust_drift");
3527 if (!show_cfl_detail) {
3528 save_dt_min_to_dt_part();
3529 cfl_detail.push_back({
"all SPH", cfl_dt->get_native().compute_rank_min()});
3532 auto &sync = scheduler().synchronized_data;
3537 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> G_edge
3538 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"G",
"G");
3539 G_edge->data = solver_config.get_constant_G();
3541 std::shared_ptr<shamrock::solvergraph::IDataEdge<Tscal>> sink_sink_cfl
3542 = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
3543 "sink_sink_cfl",
"\\Delta t_{\\rm sink-sink}");
3550 compute_cfl_sink_sink.set_edges(
3554 sync.template get_edge_ptr<SinkVecEdge>(
"sink_pos"),
3555 sync.template get_edge_ptr<SinkScalEdge>(
"sink_mass"),
3556 sync.template get_edge_ptr<SinkVecEdge>(
"sink_acc_ext"),
3560 cfl_detail.push_back({
"sink_sink", sink_sink_cfl->data});
3563 Tscal rank_dt = shambase::get_infty<Tscal>();
3564 for (
auto &[key, value] : cfl_detail) {
3565 rank_dt = sham::min(rank_dt, value);
3568 if (show_cfl_detail) {
3569 for (
auto &[key, value] : cfl_detail) {
3570 value = shamalgs::collective::allreduce_min(value);
3575 table.add_double_rule();
3576 table.add_data({
"key",
"value"}, shambase::table::center);
3577 table.add_double_rule();
3578 for (
auto &[key, value] : cfl_detail) {
3580 {key, sham::format(
"{:.2e}", value)}, shambase::table::right);
3587 next_cfl = shamalgs::collective::allreduce_min(rank_dt);
3591 "sph::Model",
"cfl dt =", next_cfl,
"cfl multiplier :", get_cfl_multipler());
3597 if (solver_config.has_field_soundspeed()) {
3612 auto &q = shamsys::instance::get_compute_scheduler().get_queue();
3615 auto cs_in = buf_cs_in.get_read_access(depends_list);
3618 auto e = q.submit(depends_list, [&](sycl::handler &cgh) {
3619 const Tscal pmass = solver_config.gpart_mass;
3622 sycl::range<1>{pdat.
get_obj_cnt()}, [=](sycl::item<1> item) {
3623 cs[item] = cs_in[item];
3627 buf_cs_in.complete_event_state(e);
3634 corrector_iter_cnt++;
3636 if (solver_config.has_field_alphaAV()) {
3637 storage.alpha_av_ghost.reset();
3639 }
while (need_rerun_corrector);
3641 reset_merge_ghosts_fields();
3648 for (
auto it = timestep_callbacks.rbegin(); it != timestep_callbacks.rend(); ++it) {
3649 if (it->step_end_callback) {
3662 = (mem_perf_infos_end.
time_alloc_device - mem_perf_infos_start.time_alloc_device)
3663 + (mem_perf_infos_end.
time_free_device - mem_perf_infos_start.time_free_device);
3664 f64 t_host_alloc = (mem_perf_infos_end.
time_alloc_host - mem_perf_infos_start.time_alloc_host)
3665 + (mem_perf_infos_end.
time_free_host - mem_perf_infos_start.time_free_host);
3667 u64 rank_count = scheduler().get_rank_count();
3668 f64 rate =
f64(rank_count) / tstep.elapsed_sec();
3670 u64 npatch = scheduler().patch_list.local.size();
3674 std::string log_step = report_perf_timestep(
3678 tstep.elapsed_sec(),
3684 system_metrics_delta,
3685 shamsys::has_reporter());
3690 "sph::Model",
"estimated rate :", dt * (3600 / tstep.elapsed_sec()),
"(tsim/hr)");
3693 solve_logs.register_log(
3699 tstep.elapsed_sec(),
3701 system_metrics_delta});
3703 storage.timings_details.reset();
3705 reset_serial_patch_tree();
3706 reset_ghost_handler();
3712 storage.merged_xyzh.reset();
3714 clear_merged_pos_trees();
3715 clear_ghost_cache();
3716 reset_presteps_rint();
3717 reset_neighbors_cache();
3721 set_next_dt(next_cfl);
3722 set_time(t_current + dt);
3724 auto get_next_cfl_mult = [&]() {
3725 Tscal cfl_m = get_cfl_multipler();
3726 Tscal stiff = solver_config.cfl_config.cfl_multiplier_stiffness;
3728 return (cfl_m * stiff + 1.) / (stiff + 1.);
3731 set_cfl_multipler(get_next_cfl_mult());
3736 log.npart = rank_count;
3737 log.tcompute = tstep.elapsed_sec();