29template<
class Tvec,
class Tgr
idVec>
30template<
class UserAcc,
class... T>
37 using namespace shamrock::patch;
43 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
49 u32 obj_cnt = pdat.get_obj_cnt();
57 UserAcc uacc(depends_list, id_patch, cur_p, pdat,
args...);
59 auto refine_acc = refine_flags.get_write_access(depends_list);
60 auto derefine_acc = derefine_flags.get_write_access(depends_list);
63 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
64 cgh.parallel_for(sycl::range<1>(obj_cnt), [=](sycl::item<1> gid) {
65 bool flag_refine =
false;
66 bool flag_derefine =
false;
67 uacc.refine_criterion(gid.get_linear_id(), uacc, flag_refine, flag_derefine);
70 if (flag_refine && flag_derefine) {
71 flag_derefine =
false;
74 refine_acc[gid] = (flag_refine) ? 1 : 0;
75 derefine_acc[gid] = (flag_derefine) ? 1 : 0;
82 refine_flags.complete_event_state(resulting_events);
83 derefine_flags.complete_event_state(resulting_events);
85 uacc.finalize(resulting_events, id_patch, cur_p, pdat,
args...);
93 auto acc_max = buf_cell_max.get_read_access(depends_list);
94 auto acc_merge_flag = derefine_flags.get_write_access(depends_list);
97 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
98 cgh.parallel_for(sycl::range<1>(obj_cnt), [=](sycl::item<1> gid) {
99 u32 id = gid.get_linear_id();
101 std::array<BlockCoord, split_count> blocks;
102 bool do_merge =
true;
106 if (
id + split_count <= obj_cnt) {
107 bool all_want_to_merge =
true;
109 for (
u32 lid = 0; lid < split_count; lid++) {
110 blocks[lid] = BlockCoord{acc_min[gid + lid], acc_max[gid + lid]};
111 all_want_to_merge = all_want_to_merge && acc_merge_flag[gid + lid];
114 do_merge = all_want_to_merge && BlockCoord::are_mergeable(blocks);
120 acc_merge_flag[gid] = do_merge;
125 buf_cell_max.complete_event_state(e);
126 derefine_flags.complete_event_state(e);
136 "AMRGrid",
"patch ", id_patch,
"refine block count = ", buf_refine.get_size());
138 tot_refine += buf_refine.get_size();
141 dd_refine_list.add_obj(id_patch, std::move(buf_refine));
151 "AMRGrid",
"patch ", id_patch,
"merge block count = ", buf_derefine.get_size());
153 tot_derefine += buf_derefine.get_size();
156 dd_derefine_list.add_obj(id_patch, std::move(buf_derefine));
159 logger::info_ln(
"AMRGrid",
"on this process", tot_refine,
"blocks were refined");
161 "AMRGrid",
"on this process", tot_derefine * split_count,
"blocks were derefined");
163template<
class Tvec,
class Tgr
idVec>
164template<
class UserAcc>
168 using namespace shamrock::patch;
170 u64 sum_block_count = 0;
172 bool new_cell_were_added =
false;
177 u32 old_obj_cnt = pdat.get_obj_cnt();
184 pdat.expand(refine_flags.
get_size() * (split_count - 1));
191 auto block_bound_high = buf_cell_max.get_write_access(depends_list);
192 UserAcc uacc(depends_list, pdat);
196 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
197 u32 start_index_push = old_obj_cnt;
199 constexpr u32 new_splits = split_count - 1;
201 cgh.parallel_for(sycl::range<1>(refine_flags.
get_size()), [=](sycl::item<1> gid) {
202 u32 tid = gid.get_linear_id();
204 u32 idx_to_refine = index_to_ref[tid];
207 BlockCoord cur_block{
208 block_bound_low[idx_to_refine], block_bound_high[idx_to_refine]};
210 std::array<BlockCoord, split_count> block_coords
211 = BlockCoord::get_split(cur_block.bmin, cur_block.bmax);
214 std::array<u32, split_count> blocks_ids;
215 blocks_ids[0] = idx_to_refine;
220 for (
u32 pid = 0; pid < new_splits; pid++) {
221 blocks_ids[pid + 1] = start_index_push + tid * new_splits + pid;
227 for (
u32 pid = 0; pid < split_count; pid++) {
228 block_bound_low[blocks_ids[pid]] = block_coords[pid].bmin;
229 block_bound_high[blocks_ids[pid]] = block_coords[pid].bmax;
233 uacc.apply_refine(idx_to_refine, cur_block, blocks_ids, block_coords, uacc);
240 buf_cell_max.complete_event_state(resulting_events);
242 uacc.finalize(resulting_events, pdat);
247 sum_block_count += pdat.get_obj_cnt();
248 new_cell_were_added = new_cell_were_added || refine_flags.
get_size() > 0;
253 return new_cell_were_added;
256template<
class Tvec,
class Tgr
idVec>
257template<
class UserAcc>
262 using namespace shamrock::patch;
264 bool cell_were_removed =
false;
267 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
270 u32 old_obj_cnt = pdat.get_obj_cnt();
274 if (derefine_flags.
get_size() > 0) {
278 keep_block_flag.fill(1);
285 auto block_bound_high = buf_cell_max.get_write_access(depends_list);
286 UserAcc uacc(depends_list, pdat);
288 auto flag_keep = keep_block_flag.get_write_access(depends_list);
291 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
292 cgh.parallel_for(sycl::range<1>(derefine_flags.
get_size()), [=](sycl::item<1> gid) {
293 u32 tid = gid.get_linear_id();
295 u32 idx_to_derefine = index_to_deref[gid];
298 std::array<u32, split_count> old_indexes;
300 for (u32 pid = 0; pid < split_count; pid++) {
301 old_indexes[pid] = idx_to_derefine + pid;
305 std::array<BlockCoord, split_count> block_coords;
307 for (
u32 pid = 0; pid < split_count; pid++) {
308 block_coords[pid] = BlockCoord{
309 block_bound_low[old_indexes[pid]], block_bound_high[old_indexes[pid]]};
313 BlockCoord merged_block_coord = BlockCoord::get_merge(block_coords);
316 block_bound_low[idx_to_derefine] = merged_block_coord.bmin;
317 block_bound_high[idx_to_derefine] = merged_block_coord.bmax;
321 for (
u32 pid = 1; pid < split_count; pid++) {
322 flag_keep[idx_to_derefine + pid] = 0;
327 old_indexes, block_coords, idx_to_derefine, merged_block_coord, uacc);
334 buf_cell_max.complete_event_state(resulting_events);
336 uacc.finalize(resulting_events, pdat);
338 keep_block_flag.complete_event_state(resulting_events);
349 "derefine block count ",
352 buf_keep.get_size());
354 if (buf_keep.get_size() == 0) {
355 throw std::runtime_error(
"buf keep must contain something at this point");
361 cell_were_removed = cell_were_removed || derefine_flags.
get_size() > 0;
365 return cell_were_removed;
368template<
class Tvec,
class Tgr
idVec>
369template<
class UserAccCrit,
class UserAccSplit,
class UserAccMerge>
374 AMRSortBlocks block_sorter(context, solver_config, storage);
375 block_sorter.reorder_amr_blocks();
381 gen_refine_block_changes<UserAccCrit>(dd_refine_list, dd_derefine_list);
385 internal_refine_grid<UserAccSplit>(std::move(dd_refine_list));
392 internal_derefine_grid<UserAccMerge>(std::move(dd_derefine_list));
395template<
class Tvec,
class Tgr
idVec>
399 class RefineCritBlock {
401 const TgridVec *block_low_bound;
402 const TgridVec *block_high_bound;
403 const Tscal *block_density_field;
405 Tscal one_over_Nside = 1. / AMRBlock::Nside;
411 sham::EventList &depends_list,
413 shamrock::patch::Patch p,
414 shamrock::patch::PatchDataLayer &pdat,
417 : dxfact(dxfact), wanted_mass(wanted_mass) {
419 block_low_bound = pdat.get_field<TgridVec>(0).get_buf().get_read_access(depends_list);
420 block_high_bound = pdat.get_field<TgridVec>(1).get_buf().get_read_access(depends_list);
421 block_density_field = pdat.get_field<Tscal>(pdat.pdl().get_field_idx<Tscal>(
"rho"))
423 .get_read_access(depends_list);
427 sham::EventList &resulting_events,
429 shamrock::patch::Patch p,
430 shamrock::patch::PatchDataLayer &pdat,
434 sham::DeviceBuffer<i64_3> &buf_cell_low_bound = pdat.get_field<i64_3>(0).get_buf();
435 sham::DeviceBuffer<i64_3> &buf_cell_high_bound = pdat.get_field<i64_3>(1).get_buf();
439 pdat.get_field<Tscal>(pdat.pdl().get_field_idx<Tscal>(
"rho"))
441 .complete_event_state(resulting_events);
444 void refine_criterion(
445 u32 block_id, RefineCritBlock acc,
bool &should_refine,
bool &should_derefine)
const {
447 TgridVec low_bound = acc.block_low_bound[block_id];
448 TgridVec high_bound = acc.block_high_bound[block_id];
450 Tvec lower_flt = low_bound.template convert<Tscal>() * dxfact;
451 Tvec upper_flt = high_bound.template convert<Tscal>() * dxfact;
453 Tvec block_cell_size = (upper_flt - lower_flt) * one_over_Nside;
456 for (
u32 i = 0; i < AMRBlock::block_size; i++) {
457 sum_mass += acc.block_density_field[i + block_id * AMRBlock::block_size];
459 sum_mass *= block_cell_size.x() * block_cell_size.y() * block_cell_size.z();
461 if (sum_mass > wanted_mass * 8) {
462 should_refine =
true;
463 should_derefine =
false;
464 }
else if (sum_mass < wanted_mass) {
465 should_refine =
false;
466 should_derefine =
true;
468 should_refine =
false;
469 should_derefine =
false;
472 should_refine = should_refine && (high_bound.x() - low_bound.x() > AMRBlock::Nside);
473 should_refine = should_refine && (high_bound.y() - low_bound.y() > AMRBlock::Nside);
474 should_refine = should_refine && (high_bound.z() - low_bound.z() > AMRBlock::Nside);
478 class RefineCellAccessor {
484 RefineCellAccessor(sham::EventList &depends_list, shamrock::patch::PatchDataLayer &pdat) {
486 rho = pdat.get_field<
f64>(2).get_buf().get_write_access(depends_list);
487 rho_vel = pdat.get_field<f64_3>(3).get_buf().get_write_access(depends_list);
488 rhoE = pdat.get_field<
f64>(4).get_buf().get_write_access(depends_list);
491 void finalize(sham::EventList &resulting_events, shamrock::patch::PatchDataLayer &pdat) {
492 pdat.get_field<
f64>(2).get_buf().complete_event_state(resulting_events);
493 pdat.get_field<f64_3>(3).get_buf().complete_event_state(resulting_events);
494 pdat.get_field<
f64>(4).get_buf().complete_event_state(resulting_events);
499 BlockCoord cur_coords,
500 std::array<u32, 8> new_blocks,
501 std::array<BlockCoord, 8> new_block_coords,
502 RefineCellAccessor acc)
const {
504 auto get_coord_ref = [](
u32 i) -> std::array<u32, dim> {
505 constexpr u32 NsideBlockPow = 1;
506 constexpr u32 Nside = 1U << NsideBlockPow;
508 if constexpr (dim == 3) {
509 const u32 tmp = i >> NsideBlockPow;
510 return {i % Nside, (tmp) % Nside, (tmp) >> NsideBlockPow};
514 auto get_index_block = [](std::array<u32, dim> coord) ->
u32 {
515 constexpr u32 NsideBlockPow = 1;
516 constexpr u32 Nside = 1U << NsideBlockPow;
518 if constexpr (dim == 3) {
519 return coord[0] + Nside * coord[1] + Nside * Nside * coord[2];
523 auto get_gid_write = [&](std::array<u32, dim> &glid) ->
u32 {
524 std::array<u32, dim> bid
525 = {glid[0] >> AMRBlock::NsideBlockPow,
526 glid[1] >> AMRBlock::NsideBlockPow,
527 glid[2] >> AMRBlock::NsideBlockPow};
530 return new_blocks[get_index_block(bid)] * AMRBlock::block_size
531 + AMRBlock::get_index(
532 {glid[0] % AMRBlock::Nside,
533 glid[1] % AMRBlock::Nside,
534 glid[2] % AMRBlock::Nside});
537 std::array<f64, AMRBlock::block_size> old_rho_block;
538 std::array<f64_3, AMRBlock::block_size> old_rho_vel_block;
539 std::array<f64, AMRBlock::block_size> old_rhoE_block;
542 for (
u32 loc_id = 0; loc_id < AMRBlock::block_size; loc_id++) {
544 auto [lx, ly, lz] = get_coord_ref(loc_id);
545 u32 old_cell_idx = cur_idx * AMRBlock::block_size + loc_id;
546 old_rho_block[loc_id] = acc.rho[old_cell_idx];
547 old_rho_vel_block[loc_id] = acc.rho_vel[old_cell_idx];
548 old_rhoE_block[loc_id] = acc.rhoE[old_cell_idx];
551 for (
u32 loc_id = 0; loc_id < AMRBlock::block_size; loc_id++) {
553 auto [lx, ly, lz] = get_coord_ref(loc_id);
554 u32 old_cell_idx = cur_idx * AMRBlock::block_size + loc_id;
556 Tscal rho_block = old_rho_block[loc_id];
557 Tvec rho_vel_block = old_rho_vel_block[loc_id];
558 Tscal rhoE_block = old_rhoE_block[loc_id];
559 for (
u32 subdiv_lid = 0; subdiv_lid < 8; subdiv_lid++) {
561 auto [sx, sy, sz] = get_coord_ref(subdiv_lid);
563 std::array<u32, 3> glid = {lx * 2 + sx, ly * 2 + sy, lz * 2 + sz};
565 u32 new_cell_idx = get_gid_write(glid);
580 acc.rho[new_cell_idx] = rho_block;
581 acc.rho_vel[new_cell_idx] = rho_vel_block;
582 acc.rhoE[new_cell_idx] = rhoE_block;
588 std::array<u32, 8> old_blocks,
589 std::array<BlockCoord, 8> old_coords,
591 BlockCoord new_coord,
593 RefineCellAccessor acc)
const {
595 std::array<f64, AMRBlock::block_size> rho_block;
596 std::array<f64_3, AMRBlock::block_size> rho_vel_block;
597 std::array<f64, AMRBlock::block_size> rhoE_block;
599 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
600 rho_block[cell_id] = {};
601 rho_vel_block[cell_id] = {};
602 rhoE_block[cell_id] = {};
605 for (
u32 pid = 0; pid < 8; pid++) {
606 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
607 rho_block[cell_id] += acc.rho[old_blocks[pid] * AMRBlock::block_size + cell_id];
608 rho_vel_block[cell_id]
609 += acc.rho_vel[old_blocks[pid] * AMRBlock::block_size + cell_id];
611 += acc.rhoE[old_blocks[pid] * AMRBlock::block_size + cell_id];
615 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
616 rho_block[cell_id] /= 8;
617 rho_vel_block[cell_id] /= 8;
618 rhoE_block[cell_id] /= 8;
621 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
622 u32 newcell_idx = new_cell * AMRBlock::block_size + cell_id;
623 acc.rho[newcell_idx] = rho_block[cell_id];
624 acc.rho_vel[newcell_idx] = rho_vel_block[cell_id];
625 acc.rhoE[newcell_idx] = rhoE_block[cell_id];
630 using AMRmode_None =
typename AMRMode<Tvec, TgridVec>::None;
631 using AMRmode_DensityBased =
typename AMRMode<Tvec, TgridVec>::DensityBased;
633 bool has_cell_order_changed =
false;
635 if (AMRmode_None *cfg = std::get_if<AMRmode_None>(&solver_config.amr_mode.config)) {
638 AMRmode_DensityBased *cfg
639 = std::get_if<AMRmode_DensityBased>(&solver_config.amr_mode.config)) {
640 Tscal dxfact(solver_config.grid_coord_to_pos_fact);
646 gen_refine_block_changes_old<RefineCritBlock>(
647 refine_list, derefine_list, dxfact, cfg->crit_mass);
651 bool change_refine = internal_refine_grid_old<RefineCellAccessor>(std::move(refine_list));
659 = internal_derefine_grid_old<RefineCellAccessor>(std::move(derefine_list));
661 has_cell_order_changed = has_cell_order_changed || (change_refine || change_derefine);
664 if (has_cell_order_changed) {
666 AMRSortBlocks block_sorter(context, solver_config, storage);
667 block_sorter.reorder_amr_blocks();
671template<
class Tvec,
class Tgr
idVec>
672template<
class UserAcc,
class... T>
680 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
686 u32 obj_cnt = pdat.get_obj_cnt();
694 UserAcc uacc(depends_list, storage, id_patch, cur_p, pdat,
args...);
700 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
701 cgh.parallel_for(sycl::range<1>(obj_cnt), [=](sycl::item<1> gid) {
702 bool flag_refine =
false;
703 bool flag_derefine =
false;
704 uacc.refine_criterion_new(
705 gid.get_linear_id(), uacc, flag_refine, flag_derefine);
708 if (flag_refine && flag_derefine) {
709 flag_derefine =
false;
712 refine_acc[gid] = (flag_refine) ? 1 : 0;
713 derefine_acc[gid] = (flag_derefine) ? 1 : 0;
723 uacc.finalize_new(resulting_events, storage, id_patch, cur_p, pdat,
args...);
726 dd_refine_flags.add_obj(id_patch, std::move(refine_flags));
727 dd_derefine_flags.add_obj(id_patch, std::move(derefine_flags));
737template<
class Tvec,
class Tgr
idVec>
746 u32 obj_cnt = pdat.get_obj_cnt();
750 .get_refs_dir(Direction_::xp)
753 .get_refs_dir(Direction_::xm)
756 .get_refs_dir(Direction_::yp)
759 .get_refs_dir(Direction_::ym)
762 .get_refs_dir(Direction_::zp)
765 .get_refs_dir(Direction_::zm)
771 std::shared_ptr<sham::DeviceScheduler> dev_sched
772 = shamsys::instance::get_compute_scheduler_ptr();
776 for (
u32 pass = 0; pass < 100; pass++) {
778 changed_buf.set_val_at_idx(0, 0);
781 AMRGraphLinkiterator block_graph_xp
782 = block_graph_neighs_xp.get_read_access(depend_list);
783 AMRGraphLinkiterator block_graph_xm
784 = block_graph_neighs_xm.get_read_access(depend_list);
785 AMRGraphLinkiterator block_graph_yp
786 = block_graph_neighs_yp.get_read_access(depend_list);
787 AMRGraphLinkiterator block_graph_ym
788 = block_graph_neighs_ym.get_read_access(depend_list);
789 AMRGraphLinkiterator block_graph_zp
790 = block_graph_neighs_zp.get_read_access(depend_list);
791 AMRGraphLinkiterator block_graph_zm
792 = block_graph_neighs_zm.get_read_access(depend_list);
794 auto acc_amr_levels = buf_amr_block_levels.
get_read_access(depend_list);
795 auto acc_changed = changed_buf.get_write_access(depend_list);
799 auto e = q.
submit(depend_list, [&](sycl::handler &cgh) {
800 cgh.parallel_for(sycl::range<1>(obj_cnt), [=](sycl::item<1> gid) {
801 u32 block_id = gid.get_linear_id();
803 u32 cur_ref_flag = acc_ref_flags[block_id];
808 auto cur_block_level = acc_amr_levels[block_id];
810 auto check_2To1_ref = [&](
u32 nid) {
815 u32 neigh_ref_flag = acc_ref_flags[nid];
816 auto neigh_block_level = acc_amr_levels[nid];
818 auto cur_future = cur_block_level + (cur_ref_flag ? 1 : 0);
820 auto neigh_future = neigh_block_level + (neigh_ref_flag ? 1 : 0);
822 if (cur_ref_flag && (cur_future > neigh_future + 1)) {
824 if (!neigh_ref_flag) {
827 sycl::memory_order::relaxed,
828 sycl::memory_scope::system>
829 atomic_neigh_flag(acc_ref_flags[nid]);
830 atomic_neigh_flag.exchange(1);
834 sycl::memory_order::relaxed,
835 sycl::memory_scope::system>
836 atomic_changed(acc_changed[0]);
837 atomic_changed.exchange(1);
842 block_graph_xp.for_each_object_link(block_id, check_2To1_ref);
843 block_graph_xm.for_each_object_link(block_id, check_2To1_ref);
844 block_graph_yp.for_each_object_link(block_id, check_2To1_ref);
845 block_graph_ym.for_each_object_link(block_id, check_2To1_ref);
846 block_graph_zp.for_each_object_link(block_id, check_2To1_ref);
847 block_graph_zm.for_each_object_link(block_id, check_2To1_ref);
850 block_graph_neighs_xp.complete_event_state(e);
851 block_graph_neighs_xm.complete_event_state(e);
852 block_graph_neighs_yp.complete_event_state(e);
853 block_graph_neighs_ym.complete_event_state(e);
854 block_graph_neighs_zp.complete_event_state(e);
855 block_graph_neighs_zm.complete_event_state(e);
858 changed_buf.complete_event_state(e);
864 if (changed_buf.get_val_at_idx(0) == 0) {
865 logger::raw_ln(
"Refinement 2:1 balance converged in ", pass + 1,
" sweeps");
879 "AMRGrid",
"patch ", id_patch, dev_buf_ref.get_size(),
"marked for refinement + 2:1");
891template<
class Tvec,
class Tgr
idVec>
897 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
906 u32 obj_cnt = pdat.get_obj_cnt();
910 .get_refs_dir(Direction_::xp)
913 .get_refs_dir(Direction_::xm)
916 .get_refs_dir(Direction_::yp)
919 .get_refs_dir(Direction_::ym)
922 .get_refs_dir(Direction_::zp)
925 .get_refs_dir(Direction_::zm)
938 " Count block's flag for derefinement [No geometry validity check and no 2:1 check] \t "
940 dev_buf_deref_0.get_size(),
950 auto acc_max = buf_cell_max.get_read_access(depends_list);
951 auto acc_amr_levels = buf_amr_block_levels.
get_read_access(depends_list);
954 auto acc_refine_flag = patch_refine_flag.
get_read_access(depends_list);
956 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
957 cgh.parallel_for(sycl::range<1>(obj_cnt), [=](sycl::item<1> gid) {
958 u32 id = gid.get_linear_id();
960 std::array<BlockCoord, split_count> blocks;
961 bool do_merge =
true;
962 bool all_same_level =
true;
966 if (
id + split_count <= obj_cnt) {
967 bool all_want_to_merge =
true;
969 auto get_coord = [](
u32 i) -> std::array<u32, dim> {
970 constexpr u32 NsideBlockPow = 1;
971 constexpr u32 Nside = 1U << NsideBlockPow;
972 constexpr u32 side_size = Nside;
975 if constexpr (dim == 3) {
976 const u32 tmp = i >> NsideBlockPow;
979 return {(tmp) >> NsideBlockPow, (tmp) % Nside, i % Nside};
984 = [=](BlockCoord target_block) -> std::array<BlockCoord, split_count> {
985 std::array<BlockCoord, split_count> ret;
986 auto bmin = target_block.bmin;
987 auto bmax = target_block.bmax;
988 auto split = bmin + (bmax - bmin) / 2;
989 std::array<TgridVec, 3> szs = {bmin, split, bmax};
990 for (
u32 i = 0; i < split_count; i++) {
991 auto [lx, ly, lz] = get_coord(i);
993 ret[i].bmin = TgridVec{szs[lx].x(), szs[ly].y(), szs[lz].z()};
995 = TgridVec{szs[lx + 1].x(), szs[ly + 1].y(), szs[lz + 1].z()};
1001 for (
u32 b_lid = 0; b_lid < split_count; b_lid++) {
1002 blocks[b_lid] = BlockCoord{acc_min[
id + b_lid], acc_max[
id + b_lid]};
1003 all_want_to_merge = all_want_to_merge && acc_merge_flag[
id + b_lid];
1005 = all_same_level && (acc_amr_levels[id] == acc_amr_levels[
id + b_lid]);
1008 BlockCoord merged = BlockCoord::get_merge(blocks);
1009 std::array<BlockCoord, split_count> splitted = get_split(merged);
1010 for (
u32 lid = 0; lid < split_count; lid++) {
1011 do_merge = do_merge && sham::equals(blocks[lid].bmin, splitted[lid].bmin)
1012 && sham::equals(blocks[lid].bmax, splitted[lid].bmax);
1015 do_merge = do_merge && all_want_to_merge && all_same_level;
1016 if (acc_refine_flag[
id] && do_merge) {
1023 acc_merge_flag[id] = do_merge;
1027 buf_cell_max.complete_event_state(e);
1037 " Count block's flag for derefinement [After geometry validity check and before 2:1 "
1040 buf_derefine_1.get_size(),
1049 std::shared_ptr<sham::DeviceScheduler> dev_sched
1050 = shamsys::instance::get_compute_scheduler_ptr();
1056 patch_derefine_flag.
copy_range(0, obj_cnt, patch_derefine_flag_old);
1061 for (
int it = 0; it < 100; it++) {
1062 changed_buf.set_val_at_idx(0, 0);
1066 AMRGraphLinkiterator block_graph_xp
1067 = block_graph_neighs_xp.get_read_access(depend_list);
1069 AMRGraphLinkiterator block_graph_xm
1070 = block_graph_neighs_xm.get_read_access(depend_list);
1071 AMRGraphLinkiterator block_graph_yp
1072 = block_graph_neighs_yp.get_read_access(depend_list);
1073 AMRGraphLinkiterator block_graph_ym
1074 = block_graph_neighs_ym.get_read_access(depend_list);
1075 AMRGraphLinkiterator block_graph_zp
1076 = block_graph_neighs_zp.get_read_access(depend_list);
1077 AMRGraphLinkiterator block_graph_zm
1078 = block_graph_neighs_zm.get_read_access(depend_list);
1080 auto acc_amr_levels = buf_amr_block_levels.
get_read_access(depend_list);
1082 auto acc_changed = changed_buf.get_write_access(depend_list);
1084 auto acc_deref_old = patch_derefine_flag_old.get_read_access(depend_list);
1085 auto acc_deref_new = patch_derefine_flag_new.get_write_access(depend_list);
1087 auto e_2to1 = q.
submit(depend_list, [&](sycl::handler &cgh) {
1088 cgh.parallel_for(sycl::range<1>(obj_cnt), [=](sycl::item<1> gid) {
1089 auto lid = gid.get_linear_id();
1091 auto old_flag = acc_deref_old[lid];
1092 auto new_flag = old_flag;
1094 auto check_2To1_der = [&](
u32 nid) {
1099 auto neigh_future = acc_amr_levels[nid] + (acc_ref_flag[nid] ? 1 : 0)
1100 - (acc_deref_old[nid] ? 1 : 0);
1102 auto my_future = acc_amr_levels[lid] - 1;
1104 if (neigh_future > my_future + 1) {
1112 for (
u32 i = 0; i < AMRBlock::block_size; i++) {
1113 block_graph_xp.for_each_object_link((lid + i), check_2To1_der);
1114 block_graph_xm.for_each_object_link((lid + i), check_2To1_der);
1115 block_graph_yp.for_each_object_link((lid + i), check_2To1_der);
1116 block_graph_ym.for_each_object_link((lid + i), check_2To1_der);
1117 block_graph_zp.for_each_object_link((lid + i), check_2To1_der);
1118 block_graph_zm.for_each_object_link((lid + i), check_2To1_der);
1121 if (old_flag != new_flag) {
1124 sycl::memory_order::relaxed,
1125 sycl::memory_scope::system>
1126 atomic_changed(acc_changed[0]);
1127 atomic_changed.exchange(1);
1131 acc_deref_new[lid] = new_flag;
1134 block_graph_neighs_xp.complete_event_state(e_2to1);
1135 block_graph_neighs_xm.complete_event_state(e_2to1);
1136 block_graph_neighs_yp.complete_event_state(e_2to1);
1137 block_graph_neighs_ym.complete_event_state(e_2to1);
1138 block_graph_neighs_zp.complete_event_state(e_2to1);
1139 block_graph_neighs_zm.complete_event_state(e_2to1);
1142 changed_buf.complete_event_state(e_2to1);
1144 patch_derefine_flag_old.complete_event_state(e_2to1);
1145 patch_derefine_flag_new.complete_event_state(e_2to1);
1148 std::swap(patch_derefine_flag_old, patch_derefine_flag_new);
1150 if (changed_buf.get_val_at_idx(0) == 0) {
1153 "Derefinement 2:1 balance converge in \t ", it + 1,
"\t sweeps \n\n");
1159 patch_derefine_flag_old.copy_range(0, obj_cnt, patch_derefine_flag);
1169 " Count block's flag for derefinement [After geometry validity check and after 2:1 "
1171 buf_derefine.get_size(),
1175 "AMRGrid",
"patch ", id_patch, buf_derefine.get_size(),
"marked for derefinement ");
1179template<
class Tvec,
class Tgr
idVec>
1180template<
class UserAcc>
1184 u64 sum_block_count = 0;
1186 bool new_cell_were_added =
false;
1188 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
1191 u32 old_obj_cnt = pdat.get_obj_cnt();
1193 dev_sched, dd_refine_flags.get(id_patch), old_obj_cnt);
1195 if (stream_compaction_result.get_size() > 0) {
1199 "Will refine \t", stream_compaction_result.get_size(),
" \t blocks \n\n");
1201 pdat.expand(
static_cast<u32>(stream_compaction_result.get_size()) * (split_count - 1));
1207 auto block_bound_high = buf_cell_max.get_write_access(depends_list);
1208 UserAcc uacc(depends_list, storage, id_patch, pdat);
1209 auto index_to_ref = stream_compaction_result.get_read_access(depends_list);
1212 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
1213 u32 start_index_push = old_obj_cnt;
1215 constexpr u32 new_splits = split_count - 1;
1218 sycl::range<1>(stream_compaction_result.get_size()), [=](sycl::item<1> gid) {
1219 u32 tid = gid.get_linear_id();
1221 u32 idx_to_refine = index_to_ref[gid];
1224 BlockCoord cur_block{
1225 block_bound_low[idx_to_refine], block_bound_high[idx_to_refine]};
1227 std::array<BlockCoord, split_count> block_coords
1228 = BlockCoord::get_split(cur_block.bmin, cur_block.bmax);
1231 std::array<u32, split_count> blocks_ids;
1232 blocks_ids[0] = idx_to_refine;
1237 for (
u32 pid = 0; pid < new_splits; pid++) {
1238 blocks_ids[pid + 1] = start_index_push + tid * new_splits + pid;
1244 for (
u32 pid = 0; pid < split_count; pid++) {
1245 block_bound_low[blocks_ids[pid]] = block_coords[pid].bmin;
1246 block_bound_high[blocks_ids[pid]] = block_coords[pid].bmax;
1250 uacc.apply_refine_new(
1251 idx_to_refine, cur_block, blocks_ids, block_coords, uacc);
1258 buf_cell_max.complete_event_state(resulting_events);
1259 stream_compaction_result.complete_event_state(e);
1260 uacc.finalize_new(resulting_events, storage, id_patch, pdat);
1263 shamlog_debug_ln(
"AMRGrid",
"patch ", id_patch,
"new block count = ", pdat.get_obj_cnt());
1264 sum_block_count += pdat.get_obj_cnt();
1265 new_cell_were_added = new_cell_were_added || (stream_compaction_result.get_size() > 0);
1270 return new_cell_were_added;
1273template<
class Tvec,
class Tgr
idVec>
1274template<
class UserAcc>
1279 using namespace shamrock::patch;
1281 bool cell_were_removed =
false;
1284 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
1287 u32 old_obj_cnt = pdat.get_obj_cnt();
1288 u32 old_obj_cnt_before_refinement = dd_derefine_flags.get(id_patch).get_size();
1290 dev_sched, dd_derefine_flags.get(id_patch), old_obj_cnt_before_refinement);
1291 if (stream_compact_results.get_size() > 0) {
1293 sycl::buffer<u32> keep_block_flag
1302 auto block_bound_high = buf_cell_max.get_write_access(depends_list);
1303 UserAcc uacc(depends_list, storage, id_patch, pdat);
1304 auto index_to_deref = stream_compact_results.get_read_access(depends_list);
1307 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
1308 sycl::accessor flag_keep{keep_block_flag, cgh, sycl::read_write};
1310 sycl::range<1>(stream_compact_results.get_size()), [=](sycl::item<1> gid) {
1311 u32 tid = gid.get_linear_id();
1313 u32 idx_to_derefine = index_to_deref[gid];
1316 std::array<u32, split_count> old_indexes;
1318 for (u32 pid = 0; pid < split_count; pid++) {
1319 old_indexes[pid] = idx_to_derefine + pid;
1323 std::array<BlockCoord, split_count> block_coords;
1325 for (
u32 pid = 0; pid < split_count; pid++) {
1326 block_coords[pid] = BlockCoord{
1327 block_bound_low[old_indexes[pid]],
1328 block_bound_high[old_indexes[pid]]};
1332 BlockCoord merged_block_coord = BlockCoord::get_merge(block_coords);
1335 block_bound_low[idx_to_derefine] = merged_block_coord.bmin;
1336 block_bound_high[idx_to_derefine] = merged_block_coord.bmax;
1340 for (
u32 pid = 1; pid < split_count; pid++) {
1341 flag_keep[idx_to_derefine + pid] = 0;
1346 uacc.apply_derefine_new(
1347 old_indexes, block_coords, idx_to_derefine, merged_block_coord, uacc);
1354 buf_cell_max.complete_event_state(resulting_events);
1355 uacc.finalize_new(resulting_events, storage, id_patch, pdat);
1357 stream_compact_results.complete_event_state(resulting_events);
1367 "derefine block count = ",
1369 "new block count = ",
1373 throw std::runtime_error(
"opt buf must contain something at this point");
1380 cell_were_removed = cell_were_removed || stream_compact_results.get_size() > 0;
1384 return cell_were_removed;
1387template<
class Tvec,
class Tgr
idVec>
1388template<
class UserAccCrit,
class UserAccSplit,
class UserAccMerge>
1393 AMRSortBlocks block_sorter(context, solver_config, storage);
1394 block_sorter.reorder_amr_blocks();
1400 gen_refine_block_changes_new<UserAccCrit>(dd_refine_list, dd_derefine_list);
1404 internal_refine_grid_new<UserAccSplit>(std::move(dd_refine_list));
1411 internal_derefine_grid_new<UserAccMerge>(std::move(dd_derefine_list));
1414template<
class Tvec,
class Tgr
idVec>
1418 class RefineCritBlock {
1420 const TgridVec *block_low_bound;
1421 const TgridVec *block_high_bound;
1422 const Tscal *block_density_field;
1424 Tscal one_over_Nside = 1. / AMRBlock::Nside;
1437 : dxfact(dxfact), wanted_mass(wanted_mass) {
1439 block_low_bound = pdat.get_field<TgridVec>(0).get_buf().get_read_access(depends_list);
1440 block_high_bound = pdat.get_field<TgridVec>(1).get_buf().get_read_access(depends_list);
1441 block_density_field = pdat.get_field<Tscal>(pdat.pdl().get_field_idx<Tscal>(
"rho"))
1443 .get_read_access(depends_list);
1453 Tscal wanted_mass) {
1455 pdat.get_field<TgridVec>(0).get_buf().complete_event_state(resulting_events);
1456 pdat.get_field<TgridVec>(1).get_buf().complete_event_state(resulting_events);
1457 pdat.get_field<Tscal>(pdat.pdl().get_field_idx<Tscal>(
"rho"))
1459 .complete_event_state(resulting_events);
1462 void refine_criterion_new(
1463 u32 block_id, RefineCritBlock acc,
bool &should_refine,
bool &should_derefine)
const {
1465 TgridVec low_bound = acc.block_low_bound[block_id];
1466 TgridVec high_bound = acc.block_high_bound[block_id];
1468 Tvec lower_flt = low_bound.template convert<Tscal>() * dxfact;
1469 Tvec upper_flt = high_bound.template convert<Tscal>() * dxfact;
1471 Tvec block_cell_size = (upper_flt - lower_flt) * one_over_Nside;
1474 for (
u32 i = 0; i < AMRBlock::block_size; i++) {
1475 sum_mass += acc.block_density_field[i + block_id * AMRBlock::block_size];
1477 sum_mass *= block_cell_size.x() * block_cell_size.y() * block_cell_size.z();
1479 if (sum_mass > wanted_mass * 8) {
1480 should_refine =
true;
1481 should_derefine =
false;
1482 }
else if (sum_mass < wanted_mass) {
1483 should_refine =
false;
1484 should_derefine =
true;
1486 should_refine =
false;
1487 should_derefine =
false;
1490 should_refine = should_refine && (high_bound.x() - low_bound.x() > AMRBlock::Nside);
1491 should_refine = should_refine && (high_bound.y() - low_bound.y() > AMRBlock::Nside);
1492 should_refine = should_refine && (high_bound.z() - low_bound.z() > AMRBlock::Nside);
1496 class RefineCellAccessor {
1505 AMRGraphLinkiterator cell_graph_xp;
1506 AMRGraphLinkiterator cell_graph_xm;
1507 AMRGraphLinkiterator cell_graph_yp;
1508 AMRGraphLinkiterator cell_graph_ym;
1509 AMRGraphLinkiterator cell_graph_zp;
1510 AMRGraphLinkiterator cell_graph_zm;
1519 .get_refs_dir(Direction::xp)
1522 .get_read_access(depends_list)),
1525 .get_refs_dir(Direction::xm)
1528 .get_read_access(depends_list)),
1531 .get_refs_dir(Direction::yp)
1534 .get_read_access(depends_list)),
1537 .get_refs_dir(Direction::ym)
1540 .get_read_access(depends_list)),
1543 .get_refs_dir(Direction::zp)
1546 .get_read_access(depends_list)),
1549 .get_refs_dir(Direction::zm)
1552 .get_read_access(depends_list))
1556 rho = pdat.get_field<
f64>(2).get_buf().get_write_access(depends_list);
1557 rho_vel = pdat.get_field<f64_3>(3).get_buf().get_write_access(depends_list);
1558 rhoE = pdat.get_field<
f64>(4).get_buf().get_write_access(depends_list);
1569 pdat.get_field<
f64>(2).get_buf().complete_event_state(resulting_events);
1570 pdat.get_field<f64_3>(3).get_buf().complete_event_state(resulting_events);
1571 pdat.get_field<
f64>(4).get_buf().complete_event_state(resulting_events);
1574 .get_refs_dir(Direction_::xp)
1577 .complete_event_state(resulting_events);
1579 .get_refs_dir(Direction_::xm)
1582 .complete_event_state(resulting_events);
1584 .get_refs_dir(Direction_::yp)
1587 .complete_event_state(resulting_events);
1589 .get_refs_dir(Direction_::ym)
1592 .complete_event_state(resulting_events);
1594 .get_refs_dir(Direction_::zp)
1597 .complete_event_state(resulting_events);
1599 .get_refs_dir(Direction_::zm)
1602 .complete_event_state(resulting_events);
1609 void apply_refine_new(
1611 BlockCoord cur_coords,
1612 std::array<u32, 8> new_blocks,
1613 std::array<BlockCoord, 8> new_block_coords,
1614 RefineCellAccessor acc)
const {
1616 auto get_coord_ref = [](
u32 i) -> std::array<u32, dim> {
1617 constexpr u32 NsideBlockPow = 1;
1618 constexpr u32 Nside = 1U << NsideBlockPow;
1620 if constexpr (dim == 3) {
1621 const u32 tmp = i >> NsideBlockPow;
1622 return {i % Nside, (tmp) % Nside, (tmp) >> NsideBlockPow};
1626 auto get_index_block = [](std::array<u32, dim> coord) ->
u32 {
1627 constexpr u32 NsideBlockPow = 1;
1628 constexpr u32 Nside = 1U << NsideBlockPow;
1630 if constexpr (dim == 3) {
1631 return coord[0] + Nside * coord[1] + Nside * Nside * coord[2];
1635 auto get_gid_write = [&](std::array<u32, dim> &glid) ->
u32 {
1638 std::array<u32, dim> bid
1639 = {glid[0] >> AMRBlock::NsideBlockPow,
1640 glid[1] >> AMRBlock::NsideBlockPow,
1641 glid[2] >> AMRBlock::NsideBlockPow};
1644 auto new_glob_id = new_blocks[get_index_block(bid)] * AMRBlock::block_size;
1649 + AMRBlock::get_index(
1650 {glid[0] % AMRBlock::Nside,
1651 glid[1] % AMRBlock::Nside,
1652 glid[2] % AMRBlock::Nside});
1655 std::array<f64, AMRBlock::block_size> old_rho_block;
1656 std::array<f64_3, AMRBlock::block_size> old_rho_vel_block;
1657 std::array<f64, AMRBlock::block_size> old_rhoE_block;
1660 for (
u32 loc_id = 0; loc_id < AMRBlock::block_size; loc_id++) {
1662 auto [lx, ly, lz] = get_coord_ref(loc_id);
1663 u32 old_cell_idx = cur_idx * AMRBlock::block_size + loc_id;
1664 old_rho_block[loc_id] = acc.rho[old_cell_idx];
1665 old_rho_vel_block[loc_id] = acc.rho_vel[old_cell_idx];
1666 old_rhoE_block[loc_id] = acc.rhoE[old_cell_idx];
1669 for (
u32 loc_id = 0; loc_id < AMRBlock::block_size; loc_id++) {
1671 auto [lx, ly, lz] = get_coord_ref(loc_id);
1672 u32 old_cell_idx = cur_idx * AMRBlock::block_size + loc_id;
1706 Tscal rho_block = old_rho_block[loc_id];
1707 Tvec rho_vel_block = old_rho_vel_block[loc_id];
1708 Tscal rhoE_block = old_rhoE_block[loc_id];
1709 for (
u32 subdiv_lid = 0; subdiv_lid < 8; subdiv_lid++) {
1711 auto [sx, sy, sz] = get_coord_ref(subdiv_lid);
1713 std::array<u32, 3> glid = {lx * 2 + sx, ly * 2 + sy, lz * 2 + sz};
1715 u32 new_cell_idx = get_gid_write(glid);
1726 acc.rho[new_cell_idx] = rho_block;
1727 acc.rho_vel[new_cell_idx] = rho_vel_block;
1728 acc.rhoE[new_cell_idx] = rhoE_block;
1733 void apply_derefine_new(
1734 std::array<u32, 8> old_blocks,
1735 std::array<BlockCoord, 8> old_coords,
1737 BlockCoord new_coord,
1739 RefineCellAccessor acc)
const {
1741 std::array<f64, AMRBlock::block_size> rho_block;
1742 std::array<f64_3, AMRBlock::block_size> rho_vel_block;
1743 std::array<f64, AMRBlock::block_size> rhoE_block;
1745 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
1746 rho_block[cell_id] = {};
1747 rho_vel_block[cell_id] = {};
1748 rhoE_block[cell_id] = {};
1752 for (
u32 pid = 0; pid < 8; pid++) {
1753 auto rho_pid = rho_block[pid];
1754 auto rho_vel_pid = rho_vel_block[pid];
1755 auto rhoe_pid = rhoE_block[pid];
1757 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
1758 rho_pid += acc.rho[old_blocks[pid] * AMRBlock::block_size + cell_id];
1759 rho_vel_pid += acc.rho_vel[old_blocks[pid] * AMRBlock::block_size + cell_id];
1760 rhoe_pid += acc.rhoE[old_blocks[pid] * AMRBlock::block_size + cell_id];
1762 rho_block[pid] = rho_pid * (1. / 8.);
1763 rho_vel_block[pid] = rho_vel_pid * (1. / 8.);
1764 rhoE_block[pid] = rhoe_pid * (1. / 8.);
1767 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
1768 u32 newcell_idx = new_cell * AMRBlock::block_size + cell_id;
1769 acc.rho[newcell_idx] = rho_block[cell_id];
1770 acc.rho_vel[newcell_idx] = rho_vel_block[cell_id];
1771 acc.rhoE[newcell_idx] = rhoE_block[cell_id];
1779 class RefineCritPseudoGradientAccessor {
1781 Tscal one_over_Nside = 1. / AMRBlock::Nside;
1782 const TgridVec *block_low_bound;
1783 const TgridVec *block_high_bound;
1784 const Tscal *block_rho;
1785 const f64 *block_pressure;
1786 const f64_3 *block_velocity;
1788 const Tscal *rho_cons;
1791 u32 nblock_per_patch;
1793 AMRGraphLinkiterator cell_graph_xp;
1794 AMRGraphLinkiterator cell_graph_xm;
1795 AMRGraphLinkiterator cell_graph_yp;
1796 AMRGraphLinkiterator cell_graph_ym;
1797 AMRGraphLinkiterator cell_graph_zp;
1798 AMRGraphLinkiterator cell_graph_zm;
1800 RefineCritPseudoGradientAccessor(
1808 : error_min(err_min), error_max(err_max),
1811 .get_refs_dir(Direction_::xp)
1814 .get_read_access(depends_list)),
1817 .get_refs_dir(Direction_::xm)
1820 .get_read_access(depends_list)),
1823 .get_refs_dir(Direction_::yp)
1826 .get_read_access(depends_list)),
1829 .get_refs_dir(Direction_::ym)
1832 .get_read_access(depends_list)),
1835 .get_refs_dir(Direction_::zp)
1838 .get_read_access(depends_list)),
1841 .get_refs_dir(Direction_::zm)
1844 .get_read_access(depends_list))
1847 block_low_bound = pdat.get_field<TgridVec>(0).get_buf().get_read_access(depends_list);
1848 block_high_bound = pdat.get_field<TgridVec>(1).get_buf().get_read_access(depends_list);
1849 rho_cons = pdat.get_field<Tscal>(pdat.pdl().get_field_idx<Tscal>(
"rho"))
1851 .get_read_access(depends_list);
1853 nblock_per_patch = pdat.get_obj_cnt();
1857 .get_read_access(depends_list);
1860 .get_read_access(depends_list);
1863 .get_read_access(depends_list);
1875 pdat.get_field<i64_3>(0).get_buf().complete_event_state(resulting_events);
1876 pdat.get_field<i64_3>(1).get_buf().complete_event_state(resulting_events);
1877 pdat.get_field<Tscal>(pdat.pdl().get_field_idx<Tscal>(
"rho"))
1879 .complete_event_state(resulting_events);
1882 .get_refs_dir(Direction_::xp)
1885 .complete_event_state(resulting_events);
1888 .get_refs_dir(Direction_::xm)
1891 .complete_event_state(resulting_events);
1894 .get_refs_dir(Direction_::yp)
1897 .complete_event_state(resulting_events);
1899 .get_refs_dir(Direction_::ym)
1902 .complete_event_state(resulting_events);
1904 .get_refs_dir(Direction_::zp)
1907 .complete_event_state(resulting_events);
1909 .get_refs_dir(Direction_::zm)
1912 .complete_event_state(resulting_events);
1916 .complete_event_state(resulting_events);
1920 .complete_event_state(resulting_events);
1924 .complete_event_state(resulting_events);
1927 void refine_criterion_new(
1929 RefineCritPseudoGradientAccessor acc,
1930 bool &should_refine,
1931 bool &should_derefine)
const {
1932 TgridVec low_bound = acc.block_low_bound[block_id];
1933 TgridVec high_bound = acc.block_high_bound[block_id];
1952 Tscal block_rho_slope = shambase::VectorProperties<Tscal>::get_zero();
1972 Tscal block_press_grad = shambase::VectorProperties<Tscal>::get_zero();
1973 for (
u32 i = 0; i < AMRBlock::block_size; i++) {
1974 block_press_grad = sham::details::g_sycl_max(
1976 get_pseudo_grad<Tscal, Tvec>(
1977 i + block_id * AMRBlock::block_size,
1985 return block_pressure[id];
1989 Tscal error = sham::details::g_sycl_max(
1990 block_press_grad, sham::details::g_sycl_max(block_rho_slope, 0.0));
1992 should_refine =
false;
1993 should_derefine =
false;
1994 if (error > error_max) {
1995 should_refine =
true;
1996 }
else if (error < (error_min * error_max)) {
1997 should_derefine =
true;
2000 should_refine = should_refine && (high_bound.x() - low_bound.x() > AMRBlock::Nside);
2001 should_refine = should_refine && (high_bound.y() - low_bound.y() > AMRBlock::Nside);
2002 should_refine = should_refine && (high_bound.z() - low_bound.z() > AMRBlock::Nside);
2009 class RefineCritShearAccessor {
2011 Tscal one_over_Nside = 1. / AMRBlock::Nside;
2012 const TgridVec *block_low_bound;
2013 const TgridVec *block_high_bound;
2014 const Tscal *block_rho;
2015 const f64 *block_pressure;
2016 const f64_3 *block_velocity;
2022 AMRGraphLinkiterator cell_graph_xp;
2023 AMRGraphLinkiterator cell_graph_xm;
2024 AMRGraphLinkiterator cell_graph_yp;
2025 AMRGraphLinkiterator cell_graph_ym;
2026 AMRGraphLinkiterator cell_graph_zp;
2027 AMRGraphLinkiterator cell_graph_zm;
2029 RefineCritShearAccessor(
2038 : threshold(threshold), gamma(gamma), dxfact(dxfact),
2041 .get_refs_dir(Direction_::xp)
2044 .get_read_access(depends_list)),
2047 .get_refs_dir(Direction_::xm)
2050 .get_read_access(depends_list)),
2053 .get_refs_dir(Direction_::yp)
2056 .get_read_access(depends_list)),
2059 .get_refs_dir(Direction_::ym)
2062 .get_read_access(depends_list)),
2065 .get_refs_dir(Direction_::zp)
2068 .get_read_access(depends_list)),
2071 .get_refs_dir(Direction_::zm)
2074 .get_read_access(depends_list))
2077 block_low_bound = pdat.get_field<TgridVec>(0).get_buf().get_read_access(depends_list);
2078 block_high_bound = pdat.get_field<TgridVec>(1).get_buf().get_read_access(depends_list);
2082 .get_read_access(depends_list);
2085 .get_read_access(depends_list);
2088 .get_read_access(depends_list);
2101 pdat.get_field<i64_3>(0).get_buf().complete_event_state(resulting_events);
2102 pdat.get_field<i64_3>(1).get_buf().complete_event_state(resulting_events);
2105 .get_refs_dir(Direction_::xp)
2108 .complete_event_state(resulting_events);
2111 .get_refs_dir(Direction_::xm)
2114 .complete_event_state(resulting_events);
2117 .get_refs_dir(Direction_::yp)
2120 .complete_event_state(resulting_events);
2122 .get_refs_dir(Direction_::ym)
2125 .complete_event_state(resulting_events);
2127 .get_refs_dir(Direction_::zp)
2130 .complete_event_state(resulting_events);
2132 .get_refs_dir(Direction_::zm)
2135 .complete_event_state(resulting_events);
2139 .complete_event_state(resulting_events);
2143 .complete_event_state(resulting_events);
2147 .complete_event_state(resulting_events);
2150 void refine_criterion_new(
2152 RefineCritShearAccessor acc,
2153 bool &should_refine,
2154 bool &should_derefine)
const {
2155 TgridVec low_bound = acc.block_low_bound[block_id];
2156 TgridVec high_bound = acc.block_high_bound[block_id];
2158 Tvec lower_flt = low_bound.template convert<Tscal>() * dxfact;
2159 Tvec upper_flt = high_bound.template convert<Tscal>() * dxfact;
2161 Tvec block_cell_size = (upper_flt - lower_flt) * one_over_Nside;
2163 Tscal block_normalized_shear = shambase::VectorProperties<Tscal>::get_zero();
2164 for (
u32 i = 0; i < AMRBlock::block_size; i++) {
2165 auto cell_id = i + block_id * AMRBlock::block_size;
2166 auto cs = sycl::sqrt(gamma * acc.block_pressure[cell_id] / acc.block_rho[cell_id]);
2167 block_normalized_shear = sham::details::g_sycl_max(
2168 block_normalized_shear,
2169 normalized_shear<Tvec>(
2181 return acc.block_velocity[id];
2184 should_refine =
false;
2185 should_derefine =
false;
2186 if (block_normalized_shear > threshold * threshold) {
2187 should_refine =
true;
2188 }
else if (block_normalized_shear < 0.25 * threshold * threshold) {
2189 should_derefine =
true;
2192 should_refine = should_refine && (high_bound.x() - low_bound.x() > AMRBlock::Nside);
2193 should_refine = should_refine && (high_bound.y() - low_bound.y() > AMRBlock::Nside);
2194 should_refine = should_refine && (high_bound.z() - low_bound.z() > AMRBlock::Nside);
2201 class RefineCellAccessorAutogravity {
2213 AMRGraphLinkiterator cell_graph_xp;
2214 AMRGraphLinkiterator cell_graph_xm;
2215 AMRGraphLinkiterator cell_graph_yp;
2216 AMRGraphLinkiterator cell_graph_ym;
2217 AMRGraphLinkiterator cell_graph_zp;
2218 AMRGraphLinkiterator cell_graph_zm;
2220 RefineCellAccessorAutogravity(
2227 .get_refs_dir(Direction::xp)
2230 .get_read_access(depends_list)),
2233 .get_refs_dir(Direction::xm)
2236 .get_read_access(depends_list)),
2239 .get_refs_dir(Direction::yp)
2242 .get_read_access(depends_list)),
2245 .get_refs_dir(Direction::ym)
2248 .get_read_access(depends_list)),
2251 .get_refs_dir(Direction::zp)
2254 .get_read_access(depends_list)),
2257 .get_refs_dir(Direction::zm)
2260 .get_read_access(depends_list))
2264 rho = pdat.get_field<
f64>(2).get_buf().get_write_access(depends_list);
2265 rho_vel = pdat.get_field<f64_3>(3).get_buf().get_write_access(depends_list);
2266 rhoE = pdat.get_field<
f64>(4).get_buf().get_write_access(depends_list);
2267 phi_old = pdat.get_field<
f64>(pdat.pdl().get_field_idx<Tscal>(
"phi_old"))
2269 .get_write_access(depends_list);
2270 phi_new = pdat.get_field<
f64>(pdat.pdl().get_field_idx<Tscal>(
"phi"))
2272 .get_write_access(depends_list);
2280 pdat.get_field<
f64>(2).get_buf().complete_event_state(resulting_events);
2281 pdat.get_field<f64_3>(3).get_buf().complete_event_state(resulting_events);
2282 pdat.get_field<
f64>(4).get_buf().complete_event_state(resulting_events);
2283 pdat.get_field<
f64>(pdat.pdl().get_field_idx<Tscal>(
"phi_old"))
2285 .complete_event_state(resulting_events);
2286 pdat.get_field<
f64>(pdat.pdl().get_field_idx<Tscal>(
"phi"))
2288 .complete_event_state(resulting_events);
2291 .get_refs_dir(Direction_::xp)
2294 .complete_event_state(resulting_events);
2296 .get_refs_dir(Direction_::xm)
2299 .complete_event_state(resulting_events);
2301 .get_refs_dir(Direction_::yp)
2304 .complete_event_state(resulting_events);
2306 .get_refs_dir(Direction_::ym)
2309 .complete_event_state(resulting_events);
2311 .get_refs_dir(Direction_::zp)
2314 .complete_event_state(resulting_events);
2316 .get_refs_dir(Direction_::zm)
2319 .complete_event_state(resulting_events);
2322 void apply_refine_new(
2324 BlockCoord cur_coords,
2325 std::array<u32, 8> new_blocks,
2326 std::array<BlockCoord, 8> new_block_coords,
2327 RefineCellAccessorAutogravity acc)
const {
2329 auto get_coord_ref = [](
u32 i) -> std::array<u32, dim> {
2330 constexpr u32 NsideBlockPow = 1;
2331 constexpr u32 Nside = 1U << NsideBlockPow;
2333 if constexpr (dim == 3) {
2334 const u32 tmp = i >> NsideBlockPow;
2335 return {i % Nside, (tmp) % Nside, (tmp) >> NsideBlockPow};
2339 auto get_index_block = [](std::array<u32, dim> coord) ->
u32 {
2340 constexpr u32 NsideBlockPow = 1;
2341 constexpr u32 Nside = 1U << NsideBlockPow;
2343 if constexpr (dim == 3) {
2344 return coord[0] + Nside * coord[1] + Nside * Nside * coord[2];
2348 auto get_gid_write = [&](std::array<u32, dim> &glid) ->
u32 {
2351 std::array<u32, dim> bid
2352 = {glid[0] >> AMRBlock::NsideBlockPow,
2353 glid[1] >> AMRBlock::NsideBlockPow,
2354 glid[2] >> AMRBlock::NsideBlockPow};
2357 auto new_glob_id = new_blocks[get_index_block(bid)] * AMRBlock::block_size;
2362 + AMRBlock::get_index(
2363 {glid[0] % AMRBlock::Nside,
2364 glid[1] % AMRBlock::Nside,
2365 glid[2] % AMRBlock::Nside});
2368 std::array<f64, AMRBlock::block_size> old_rho_block;
2369 std::array<f64_3, AMRBlock::block_size> old_rho_vel_block;
2370 std::array<f64, AMRBlock::block_size> old_rhoE_block;
2371 std::array<f64, AMRBlock::block_size> old_phi_old_block;
2372 std::array<f64, AMRBlock::block_size> old_phi_new_block;
2375 for (
u32 loc_id = 0; loc_id < AMRBlock::block_size; loc_id++) {
2377 auto [lx, ly, lz] = get_coord_ref(loc_id);
2378 u32 old_cell_idx = cur_idx * AMRBlock::block_size + loc_id;
2379 old_rho_block[loc_id] = acc.rho[old_cell_idx];
2380 old_rho_vel_block[loc_id] = acc.rho_vel[old_cell_idx];
2381 old_rhoE_block[loc_id] = acc.rhoE[old_cell_idx];
2382 old_phi_old_block[loc_id] = acc.phi_old[old_cell_idx];
2383 old_phi_new_block[loc_id] = acc.phi_new[old_cell_idx];
2386 for (
u32 loc_id = 0; loc_id < AMRBlock::block_size; loc_id++) {
2388 auto [lx, ly, lz] = get_coord_ref(loc_id);
2389 u32 old_cell_idx = cur_idx * AMRBlock::block_size + loc_id;
2423 Tscal rho_block = old_rho_block[loc_id];
2424 Tvec rho_vel_block = old_rho_vel_block[loc_id];
2425 Tscal rhoE_block = old_rhoE_block[loc_id];
2426 Tscal phi_old_block = old_phi_old_block[loc_id];
2427 Tscal phi_new_block = old_phi_new_block[loc_id];
2429 for (
u32 subdiv_lid = 0; subdiv_lid < 8; subdiv_lid++) {
2431 auto [sx, sy, sz] = get_coord_ref(subdiv_lid);
2433 std::array<u32, 3> glid = {lx * 2 + sx, ly * 2 + sy, lz * 2 + sz};
2435 u32 new_cell_idx = get_gid_write(glid);
2446 acc.rho[new_cell_idx] = rho_block;
2447 acc.rho_vel[new_cell_idx] = rho_vel_block;
2448 acc.rhoE[new_cell_idx] = rhoE_block;
2449 acc.phi_old[new_cell_idx] = phi_old_block;
2450 acc.phi_new[new_cell_idx] = phi_new_block;
2455 void apply_derefine_new(
2456 std::array<u32, 8> old_blocks,
2457 std::array<BlockCoord, 8> old_coords,
2459 BlockCoord new_coord,
2461 RefineCellAccessorAutogravity acc)
const {
2463 std::array<f64, AMRBlock::block_size> rho_block;
2464 std::array<f64_3, AMRBlock::block_size> rho_vel_block;
2465 std::array<f64, AMRBlock::block_size> rhoE_block;
2466 std::array<f64, AMRBlock::block_size> phi_old_block;
2467 std::array<f64, AMRBlock::block_size> phi_new_block;
2469 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
2470 rho_block[cell_id] = {};
2471 rho_vel_block[cell_id] = {};
2472 rhoE_block[cell_id] = {};
2473 phi_old_block[cell_id] = {};
2474 phi_new_block[cell_id] = {};
2478 for (
u32 pid = 0; pid < 8; pid++) {
2479 auto rho_pid = rho_block[pid];
2480 auto rho_vel_pid = rho_vel_block[pid];
2481 auto rhoe_pid = rhoE_block[pid];
2482 auto phi_old_pid = phi_old_block[pid];
2483 auto phi_new_pid = phi_new_block[pid];
2485 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
2486 rho_pid += acc.rho[old_blocks[pid] * AMRBlock::block_size + cell_id];
2487 rho_vel_pid += acc.rho_vel[old_blocks[pid] * AMRBlock::block_size + cell_id];
2488 rhoe_pid += acc.rhoE[old_blocks[pid] * AMRBlock::block_size + cell_id];
2489 phi_old_pid += acc.phi_old[old_blocks[pid] * AMRBlock::block_size + cell_id];
2490 phi_new_pid += acc.phi_new[old_blocks[pid] * AMRBlock::block_size + cell_id];
2492 rho_block[pid] = rho_pid * (1. / 8.);
2493 rho_vel_block[pid] = rho_vel_pid * (1. / 8.);
2494 rhoE_block[pid] = rhoe_pid * (1. / 8.);
2499 for (
u32 cell_id = 0; cell_id < AMRBlock::block_size; cell_id++) {
2500 u32 newcell_idx = new_cell * AMRBlock::block_size + cell_id;
2501 acc.rho[newcell_idx] = rho_block[cell_id];
2502 acc.rho_vel[newcell_idx] = rho_vel_block[cell_id];
2503 acc.rhoE[newcell_idx] = rhoE_block[cell_id];
2517 bool has_cell_order_changed =
false;
2523 if (AMRmode_None *cfg = std::get_if<AMRmode_None>(&solver_config.amr_mode.config)) {
2526 if (AMRmode_DensityBased *cfg
2527 = std::get_if<AMRmode_DensityBased>(&solver_config.amr_mode.config)) {
2529 Tscal dxfact(solver_config.grid_coord_to_pos_fact);
2530 gen_refine_block_changes_new<RefineCritBlock>(
2531 refine_list, derefine_list, dxfact, cfg->crit_mass);
2535 AMRmode_PseudoGradientBased *cfg
2536 = std::get_if<AMRmode_PseudoGradientBased>(&solver_config.amr_mode.config)) {
2538 gen_refine_block_changes_new<RefineCritPseudoGradientAccessor>(
2539 refine_list, derefine_list, cfg->error_min, cfg->error_max);
2543 AMRmode_ShearBased *cfg
2544 = std::get_if<AMRmode_ShearBased>(&solver_config.amr_mode.config)) {
2545 Tscal dxfact(solver_config.grid_coord_to_pos_fact);
2546 Tscal gamma(solver_config.eos_gamma);
2548 gen_refine_block_changes_new<RefineCritShearAccessor>(
2549 refine_list, derefine_list, cfg->threshold, gamma, dxfact);
2553 enforce_two_to_one_refinement_new(std::move(refine_list));
2555 enforce_two_to_one_derefinement_new(std::move(derefine_list), std::move(refine_list));
2558 bool change_refine = internal_refine_grid_new<RefineCellAccessor>(std::move(refine_list));
2565 bool change_derefine
2566 = internal_derefine_grid_new<RefineCellAccessor>(std::move(derefine_list));
2568 has_cell_order_changed = has_cell_order_changed || (change_refine || change_derefine);
2570 if (has_cell_order_changed) {
2572 AMRSortBlocks block_sorter(context, solver_config, storage);
2573 block_sorter.reorder_amr_blocks();
double f64
Alias for double.
std::uint32_t u32
32 bit unsigned integer
std::uint64_t u64
64 bit unsigned integer
A buffer allocated in USM (Unified Shared Memory).
void complete_event_state(sycl::event e) const
Complete the event state of the buffer.
T * get_write_access(sham::EventList &depends_list, SourceLocation src_loc=SourceLocation{})
Get a read-write pointer to the buffer's data.
void copy_range(size_t begin, size_t end, sham::DeviceBuffer< T, dest_target > &dest) const
Copy a range of elements from the buffer to another buffer.
size_t get_size() const
Gets the number of elements in the buffer.
const T * get_read_access(sham::EventList &depends_list, SourceLocation src_loc=SourceLocation{}) const
Get a read-only pointer to the buffer's data.
A SYCL queue associated with a device and a context.
sycl::queue q
The SYCL queue associated with this context.
sycl::event submit(Fct &&fct)
Submits a kernel to the SYCL queue.
Class to manage a list of SYCL events.
void add_event(sycl::event e)
Add an event to the list of events.
Represents a collection of objects distributed across patches identified by a u64 id.
void update_refinement_new()
PatchDataLayer container class, the layout is described in patchdata_layout.
void index_remap_resize(sycl::buffer< u32 > &index_map, u32 len)
this function remaps the patchdatafield like so val[id] = val[index_map[id]] This function can be use...
main include file for the shamalgs algorithms
Slope mode enum + json serialization/deserialization.
alias namespace to simplify the use of log functions
sycl::buffer< typename std::invoke_result_t< Fct, u32 > > gen_buffer_device(sycl::queue &q, u32 len, Fct &&func)
generate a buffer from a lambda expression based on the indexes
std::tuple< std::optional< sycl::buffer< u32 > >, u32 > stream_compact(sycl::queue &q, sycl::buffer< u32 > &buf_flags, u32 len)
Stream compaction algorithm.
constexpr T pow_constexpr(T a) noexcept
Calculates the power of a number at compile time.
T & get_check_ref(const std::unique_ptr< T > &ptr, SourceLocation loc=SourceLocation())
Takes a std::unique_ptr and returns a reference to the object it holds. It throws a std::runtime_erro...
std::vector< std::string_view > args
Executable argument list (mapped from argv).
void raw_ln(Types... var2)
Prints a log message with multiple arguments followed by a newline.
void info_ln(std::string module_name, Types... var2)
Prints a log message with multiple arguments followed by a newline.
Patch object that contain generic patch information.
u64 id_patch
unique key that identify the patch