Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
InterpolateToFace.cpp
Go to the documentation of this file.
1// -------------------------------------------------------//
2//
3// SHAMROCK code for hydrodynamics
4// Copyright (c) 2021-2026 Timothée David--Cléris <tim.shamrock@proton.me>
5// SPDX-License-Identifier: CeCILL Free Software License Agreement v2.1
6// Shamrock is licensed under the CeCILL 2.1 License, see LICENSE for more information
7//
8// -------------------------------------------------------//
9
16
17#include "shambase/assert.hpp"
18#include "shammath/AABB.hpp"
19#include "shammodels/common/amr/AMRBlock.hpp"
20#include "shammodels/ramses/modules/InterpolateToFace.hpp"
21#include "shamrock/patch/PatchDataField.hpp"
22
23namespace {
24
25 template<class Tvec, class TgridVec, class AMRBlock>
26 class GetShift {
27
28 public:
29 using Tscal = shambase::VecComponent<Tvec>;
30 using Tgridscal = shambase::VecComponent<TgridVec>;
31
32 const Tvec *acc_aabb_block_lower;
33 const Tscal *acc_aabb_cell_size;
34
35 GetShift(const Tvec *aabb_block_lower, const Tscal *aabb_cell_size)
36 : acc_aabb_block_lower{aabb_block_lower}, acc_aabb_cell_size{aabb_cell_size} {}
37
38 shammath::AABB<Tvec> get_cell_aabb(u32 id) const {
39
40 const u32 cell_global_id = (u32) id;
41
42 const u32 block_id = cell_global_id / AMRBlock::block_size;
43 const u32 cell_loc_id = cell_global_id % AMRBlock::block_size;
44
45 // fetch current block info
46 const Tvec cblock_min = acc_aabb_block_lower[block_id];
47 const Tscal delta_cell = acc_aabb_cell_size[block_id];
48
49 std::array<u32, 3> lcoord_arr = AMRBlock::get_coord(cell_loc_id);
50 Tvec offset = Tvec{lcoord_arr[0], lcoord_arr[1], lcoord_arr[2]} * delta_cell;
51
52 Tvec aabb_min = cblock_min + offset;
53 Tvec aabb_max = aabb_min + delta_cell;
54
55 return {aabb_min, aabb_max};
56 }
57
58 std::pair<Tvec, Tvec> get_shifts(u32 id_a, u32 id_b) const {
59
60 shammath::AABB<Tvec> aabb_cell_a = get_cell_aabb(id_a);
61 shammath::AABB<Tvec> aabb_cell_b = get_cell_aabb(id_b);
62
63 shammath::AABB<Tvec> face_aabb = aabb_cell_a.get_intersect(aabb_cell_b);
64
65 Tvec face_center = face_aabb.get_center();
66
67 Tvec shift_a = face_center - aabb_cell_a.get_center();
68 Tvec shift_b = face_center - aabb_cell_b.get_center();
69
70 return {shift_a, shift_b};
71 }
72 };
73
74 template<class Tvec, class TgridVec, class AMRBlock>
75 class RhoInterpolate {
76 using Tscal = shambase::VecComponent<Tvec>;
77
78 public:
79 shamrock::PatchDataFieldSpanPointer<Tvec> aabb_block_lower;
80 shamrock::PatchDataFieldSpanPointer<Tscal> aabb_cell_size;
81 shamrock::PatchDataFieldSpanPointer<Tscal> rho_cell;
82 shamrock::PatchDataFieldSpanPointer<Tvec> grad_rho_cell;
83 // For time interpolation
84 Tscal dt_interp;
85 shamrock::PatchDataFieldSpanPointer<Tscal> dt_rho_cell;
86
87 class acc {
88 public:
89 GetShift<Tvec, TgridVec, AMRBlock> shift_get;
90
91 const Tscal *acc_rho_cell;
92 const Tvec *acc_grad_rho_cell;
93
94 // For time interpolation
95 const Tscal *acc_dt_rho_cell;
96
97 Tscal dt_interp;
98
99 acc(const Tvec *aabb_block_lower,
100 const Tscal *aabb_cell_size,
101 const Tscal *rho_cell,
102 const Tvec *grad_rho_cell,
103 // For time interpolation
104 Tscal dt_interp,
105 const Tscal *dt_rho_cell)
106 : shift_get(aabb_block_lower, aabb_cell_size), acc_rho_cell{rho_cell},
107 acc_grad_rho_cell{grad_rho_cell}, acc_dt_rho_cell{dt_rho_cell},
108 dt_interp(dt_interp) {}
109
110 std::array<Tscal, 2> get_link_field_val(u32 id_a, u32 id_b) const {
111
112 auto [shift_a, shift_b] = shift_get.get_shifts(id_a, id_b);
113
114 Tscal rho_a = acc_rho_cell[id_a];
115 Tvec grad_rho_a = acc_grad_rho_cell[id_a];
116 Tscal rho_b = acc_rho_cell[id_b];
117 Tvec grad_rho_b = acc_grad_rho_cell[id_b];
118
119 // Spatial interpolate
120 Tscal rho_face_a = rho_a + sycl::dot(grad_rho_a, shift_a);
121 Tscal rho_face_b = rho_b + sycl::dot(grad_rho_b, shift_b);
122
123 // Interpolate also to half a timestep
124 rho_face_a += acc_dt_rho_cell[id_a] * dt_interp;
125 rho_face_b += acc_dt_rho_cell[id_b] * dt_interp;
126
127 return {rho_face_a, rho_face_b};
128 }
129 };
130
131 inline acc get_read_access(sham::EventList &deps) {
132 return acc(
133 aabb_block_lower.get_read_access(deps),
134 aabb_cell_size.get_read_access(deps),
135 rho_cell.get_read_access(deps),
136 grad_rho_cell.get_read_access(deps),
137 // For time interpolation
138 dt_interp,
139 dt_rho_cell.get_read_access(deps));
140 }
141
142 inline void complete_event_state(sycl::event e) {
143 aabb_block_lower.complete_event_state(e);
144 aabb_cell_size.complete_event_state(e);
145 rho_cell.complete_event_state(e);
146 grad_rho_cell.complete_event_state(e);
147 dt_rho_cell.complete_event_state(e);
148 }
149 };
150
151 template<class Tvec, class TgridVec, class AMRBlock>
152 class VelInterpolate {
153 using Tscal = shambase::VecComponent<Tvec>;
154
155 public:
156 shamrock::PatchDataFieldSpanPointer<Tvec> aabb_block_lower;
157 shamrock::PatchDataFieldSpanPointer<Tscal> aabb_cell_size;
158
159 shamrock::PatchDataFieldSpanPointer<Tvec> vel_cell;
160 shamrock::PatchDataFieldSpanPointer<Tvec> dx_v_cell;
161 shamrock::PatchDataFieldSpanPointer<Tvec> dy_v_cell;
162 shamrock::PatchDataFieldSpanPointer<Tvec> dz_v_cell;
163 // For time interpolation
164 Tscal dt_interp;
165 shamrock::PatchDataFieldSpanPointer<Tvec> dt_vel_cell;
166
167 class acc {
168 public:
169 GetShift<Tvec, TgridVec, AMRBlock> shift_get;
170
171 const Tvec *acc_vel_cell;
172 const Tvec *acc_dx_v_cell;
173 const Tvec *acc_dy_v_cell;
174 const Tvec *acc_dz_v_cell;
175
176 // For time interpolation
177 const Tvec *acc_dt_vel_cell;
178
179 Tscal dt_interp;
180
181 acc(const Tvec *aabb_block_lower,
182 const Tscal *aabb_cell_size,
183 const Tvec *vel_cell,
184 const Tvec *dx_v_cell,
185 const Tvec *dy_v_cell,
186 const Tvec *dz_v_cell,
187 // For time interpolation
188 Tscal dt_interp,
189 const Tvec *dt_vel_cell)
190 : shift_get(aabb_block_lower, aabb_cell_size), acc_vel_cell{vel_cell},
191 acc_dx_v_cell{dx_v_cell}, acc_dy_v_cell{dy_v_cell}, acc_dz_v_cell{dz_v_cell},
192 acc_dt_vel_cell{dt_vel_cell}, dt_interp(dt_interp) {}
193
194 std::array<Tvec, 2> get_link_field_val(u32 id_a, u32 id_b) const {
195
196 auto [shift_a, shift_b] = shift_get.get_shifts(id_a, id_b);
197
198 Tvec v_a = acc_vel_cell[id_a];
199 Tvec dx_vel_a = acc_dx_v_cell[id_a];
200 Tvec dy_vel_a = acc_dy_v_cell[id_a];
201 Tvec dz_vel_a = acc_dz_v_cell[id_a];
202
203 Tvec v_b = acc_vel_cell[id_b];
204 Tvec dx_vel_b = acc_dx_v_cell[id_b];
205 Tvec dy_vel_b = acc_dy_v_cell[id_b];
206 Tvec dz_vel_b = acc_dz_v_cell[id_b];
207
208 Tvec dx_v_a_dot_shift
209 = shift_a.x() * dx_vel_a + shift_a.y() * dy_vel_a + shift_a.z() * dz_vel_a;
210 Tvec dx_v_b_dot_shift
211 = shift_b.x() * dx_vel_b + shift_b.y() * dy_vel_b + shift_b.z() * dz_vel_b;
212
213 Tvec dt_v_a = acc_dt_vel_cell[id_a];
214 Tvec dt_v_b = acc_dt_vel_cell[id_b];
215
216 Tvec vel_face_a = v_a + dx_v_a_dot_shift + dt_v_a * dt_interp;
217 Tvec vel_face_b = v_b + dx_v_b_dot_shift + dt_v_b * dt_interp;
218
219 return {vel_face_a, vel_face_b};
220 }
221 };
222
223 inline acc get_read_access(sham::EventList &deps) {
224 return acc(
225 aabb_block_lower.get_read_access(deps),
226 aabb_cell_size.get_read_access(deps),
227 vel_cell.get_read_access(deps),
228 dx_v_cell.get_read_access(deps),
229 dy_v_cell.get_read_access(deps),
230 dz_v_cell.get_read_access(deps),
231 // For time interpolation
232 dt_interp,
233 dt_vel_cell.get_read_access(deps));
234 }
235
236 inline void complete_event_state(sycl::event e) {
237 aabb_block_lower.complete_event_state(e);
238 aabb_cell_size.complete_event_state(e);
239 vel_cell.complete_event_state(e);
240 dx_v_cell.complete_event_state(e);
241 dy_v_cell.complete_event_state(e);
242 dz_v_cell.complete_event_state(e);
243 dt_vel_cell.complete_event_state(e);
244 }
245 };
246
247 template<class Tvec, class TgridVec, class AMRBlock>
248 class PressInterpolate {
249 using Tscal = shambase::VecComponent<Tvec>;
250
251 public:
252 shamrock::PatchDataFieldSpanPointer<Tvec> aabb_block_lower;
253 shamrock::PatchDataFieldSpanPointer<Tscal> aabb_cell_size;
254 shamrock::PatchDataFieldSpanPointer<Tscal> P_cell;
255 shamrock::PatchDataFieldSpanPointer<Tvec> grad_P_cell;
256 // For time interpolation
257 Tscal dt_interp;
258 shamrock::PatchDataFieldSpanPointer<Tscal> dt_P_cell;
259
260 class acc {
261 public:
262 GetShift<Tvec, TgridVec, AMRBlock> shift_get;
263
264 const Tscal *acc_P_cell;
265 const Tvec *acc_grad_P_cell;
266
267 // For time interpolation
268 const Tscal *acc_dt_P_cell;
269
270 Tscal dt_interp;
271
272 acc(const Tvec *aabb_block_lower,
273 const Tscal *aabb_cell_size,
274 const Tscal *P_cell,
275 const Tvec *grad_P_cell,
276 // For time interpolation
277 Tscal dt_interp,
278 const Tscal *dt_P_cell)
279 : shift_get(aabb_block_lower, aabb_cell_size), acc_P_cell{P_cell},
280 acc_grad_P_cell{grad_P_cell}, acc_dt_P_cell{dt_P_cell}, dt_interp(dt_interp) {}
281
282 std::array<Tscal, 2> get_link_field_val(u32 id_a, u32 id_b) const {
283
284 auto [shift_a, shift_b] = shift_get.get_shifts(id_a, id_b);
285
286 Tscal P_a = acc_P_cell[id_a];
287 Tvec grad_P_a = acc_grad_P_cell[id_a];
288 Tscal P_b = acc_P_cell[id_b];
289 Tvec grad_P_b = acc_grad_P_cell[id_b];
290
291 Tscal dtP_cell_a = acc_dt_P_cell[id_a];
292 Tscal dtP_cell_b = acc_dt_P_cell[id_b];
293
294 Tscal P_face_a = P_a + sycl::dot(grad_P_a, shift_a) + dtP_cell_a * dt_interp;
295 Tscal P_face_b = P_b + sycl::dot(grad_P_b, shift_b) + dtP_cell_b * dt_interp;
296
297 SHAM_ASSERT(P_face_a >= 0.0);
298 SHAM_ASSERT(P_face_b >= 0.0);
299
300 return {P_face_a, P_face_b};
301 }
302 };
303
304 inline acc get_read_access(sham::EventList &deps) {
305 return acc(
306 aabb_block_lower.get_read_access(deps),
307 aabb_cell_size.get_read_access(deps),
308 P_cell.get_read_access(deps),
309 grad_P_cell.get_read_access(deps),
310 dt_interp,
311 dt_P_cell.get_read_access(deps));
312 }
313
314 inline void complete_event_state(sycl::event e) {
315 aabb_block_lower.complete_event_state(e);
316 aabb_cell_size.complete_event_state(e);
317 P_cell.complete_event_state(e);
318 grad_P_cell.complete_event_state(e);
319 dt_P_cell.complete_event_state(e);
320 }
321 };
322
323 template<class Tvec, class TgridVec, class AMRBlock>
324 class RhoDustInterpolate {
325 using Tscal = shambase::VecComponent<Tvec>;
326
327 public:
328 u32 nvar;
329 shamrock::PatchDataFieldSpanPointer<Tvec> aabb_block_lower;
330 shamrock::PatchDataFieldSpanPointer<Tscal> aabb_cell_size;
331 shamrock::PatchDataFieldSpanPointer<Tscal> rho_dust_cell;
332 shamrock::PatchDataFieldSpanPointer<Tvec> grad_rho_dust_cell;
333 // For time interpolation
334 Tscal dt_interp;
335 shamrock::PatchDataFieldSpanPointer<Tscal> dt_rho_dust_cell;
336
337 class acc {
338 public:
339 GetShift<Tvec, TgridVec, AMRBlock> shift_get;
340 u32 nvar;
341
342 const Tscal *acc_rho_dust_cell;
343 const Tvec *acc_grad_rho_dust_cell;
344
345 // For time interpolation
346 const Tscal *acc_dt_rho_dust_cell;
347
348 Tscal dt_interp;
349
350 acc(u32 nvar,
351 const Tvec *aabb_block_lower,
352 const Tscal *aabb_cell_size,
353 const Tscal *rho_dust_cell,
354 const Tvec *grad_rho_dust_cell,
355 // For time interpolation
356 Tscal dt_interp,
357 const Tscal *dt_rho_dust_cell)
358 : shift_get(aabb_block_lower, aabb_cell_size), nvar(nvar),
359 acc_rho_dust_cell{rho_dust_cell}, acc_grad_rho_dust_cell{grad_rho_dust_cell},
360 acc_dt_rho_dust_cell{dt_rho_dust_cell}, dt_interp(dt_interp) {}
361
362 std::array<Tscal, 2> get_link_field_val(u32 id_a, u32 id_b) const {
363 const u32 icell_a = id_a / nvar;
364 const u32 icell_b = id_b / nvar;
365
366 auto [shift_a, shift_b] = shift_get.get_shifts(icell_a, icell_b);
367
368 Tscal rho_dust_a = acc_rho_dust_cell[id_a];
369 Tvec grad_rho_dust_a = acc_grad_rho_dust_cell[id_a];
370 Tscal rho_dust_b = acc_rho_dust_cell[id_b];
371 Tvec grad_rho_dust_b = acc_grad_rho_dust_cell[id_b];
372
373 Tscal rho_dust_face_a = rho_dust_a + sycl::dot(grad_rho_dust_a, shift_a);
374 Tscal rho_dust_face_b = rho_dust_b + sycl::dot(grad_rho_dust_b, shift_b);
375
376 rho_dust_face_a += acc_dt_rho_dust_cell[id_a] * dt_interp;
377 rho_dust_face_b += acc_dt_rho_dust_cell[id_b] * dt_interp;
378
379 return {rho_dust_face_a, rho_dust_face_b};
380 }
381 };
382
383 inline acc get_read_access(sham::EventList &deps) {
384 return acc(
385 nvar,
386 aabb_block_lower.get_read_access(deps),
387 aabb_cell_size.get_read_access(deps),
388 rho_dust_cell.get_read_access(deps),
389 grad_rho_dust_cell.get_read_access(deps),
390 // For time interpolation
391 dt_interp,
392 dt_rho_dust_cell.get_read_access(deps));
393 }
394
395 inline void complete_event_state(sycl::event e) {
396 aabb_block_lower.complete_event_state(e);
397 aabb_cell_size.complete_event_state(e);
398 rho_dust_cell.complete_event_state(e);
399 grad_rho_dust_cell.complete_event_state(e);
400 dt_rho_dust_cell.complete_event_state(e);
401 }
402 };
403
404 template<class Tvec, class TgridVec, class AMRBlock>
405 class VelDustInterpolate {
406 using Tscal = shambase::VecComponent<Tvec>;
407
408 public:
409 u32 nvar;
410 shamrock::PatchDataFieldSpanPointer<Tvec> aabb_block_lower;
411 shamrock::PatchDataFieldSpanPointer<Tscal> aabb_cell_size;
412 shamrock::PatchDataFieldSpanPointer<Tvec> vel_dust_cell;
413 shamrock::PatchDataFieldSpanPointer<Tvec> dx_v_dust_cell;
414 shamrock::PatchDataFieldSpanPointer<Tvec> dy_v_dust_cell;
415 shamrock::PatchDataFieldSpanPointer<Tvec> dz_v_dust_cell;
416 // For time interpolation
417 Tscal dt_interp;
418 shamrock::PatchDataFieldSpanPointer<Tvec> dt_vel_dust_cell;
419
420 class acc {
421 public:
422 GetShift<Tvec, TgridVec, AMRBlock> shift_get;
423 u32 nvar;
424
425 const Tvec *acc_vel_dust_cell;
426 const Tvec *acc_dx_v_dust_cell;
427 const Tvec *acc_dy_v_dust_cell;
428 const Tvec *acc_dz_v_dust_cell;
429
430 // For time interpolation
431 const Tvec *acc_dt_vel_dust_cell;
432
433 Tscal dt_interp;
434
435 acc(u32 nvar,
436 const Tvec *aabb_block_lower,
437 const Tscal *aabb_cell_size,
438 const Tvec *vel_dust_cell,
439 const Tvec *dx_v_dust_cell,
440 const Tvec *dy_v_dust_cell,
441 const Tvec *dz_v_dust_cell,
442 // For time interpolation
443 Tscal dt_interp,
444 const Tvec *dt_vel_dust_cell)
445 : shift_get(aabb_block_lower, aabb_cell_size), nvar(nvar),
446 acc_vel_dust_cell{vel_dust_cell}, acc_dx_v_dust_cell{dx_v_dust_cell},
447 acc_dy_v_dust_cell{dy_v_dust_cell}, acc_dz_v_dust_cell{dz_v_dust_cell},
448 acc_dt_vel_dust_cell{dt_vel_dust_cell}, dt_interp(dt_interp) {}
449
450 std::array<Tvec, 2> get_link_field_val(u32 id_a, u32 id_b) const {
451 const u32 icell_a = id_a / nvar;
452 const u32 icell_b = id_b / nvar;
453
454 auto [shift_a, shift_b] = shift_get.get_shifts(icell_a, icell_b);
455
456 Tvec v_dust_a = acc_vel_dust_cell[id_a];
457 Tvec dx_vel_dust_a = acc_dx_v_dust_cell[id_a];
458 Tvec dy_vel_dust_a = acc_dy_v_dust_cell[id_a];
459 Tvec dz_vel_dust_a = acc_dz_v_dust_cell[id_a];
460
461 Tvec v_dust_b = acc_vel_dust_cell[id_b];
462 Tvec dx_vel_dust_b = acc_dx_v_dust_cell[id_b];
463 Tvec dy_vel_dust_b = acc_dy_v_dust_cell[id_b];
464 Tvec dz_vel_dust_b = acc_dz_v_dust_cell[id_b];
465
466 Tvec dx_v_dust_a_dot_shift = shift_a.x() * dx_vel_dust_a
467 + shift_a.y() * dy_vel_dust_a
468 + shift_a.z() * dz_vel_dust_a;
469 Tvec dx_v_dust_b_dot_shift = shift_b.x() * dx_vel_dust_b
470 + shift_b.y() * dy_vel_dust_b
471 + shift_b.z() * dz_vel_dust_b;
472
473 Tvec dt_v_dust_a = acc_dt_vel_dust_cell[id_a];
474 Tvec dt_v_dust_b = acc_dt_vel_dust_cell[id_b];
475
476 Tvec vel_dust_face_a = v_dust_a + dx_v_dust_a_dot_shift + dt_v_dust_a * dt_interp;
477 Tvec vel_dust_face_b = v_dust_b + dx_v_dust_b_dot_shift + dt_v_dust_b * dt_interp;
478
479 return {vel_dust_face_a, vel_dust_face_b};
480 }
481 };
482
483 inline acc get_read_access(sham::EventList &deps) {
484 return acc(
485 nvar,
486 aabb_block_lower.get_read_access(deps),
487 aabb_cell_size.get_read_access(deps),
488 vel_dust_cell.get_read_access(deps),
489 dx_v_dust_cell.get_read_access(deps),
490 dy_v_dust_cell.get_read_access(deps),
491 dz_v_dust_cell.get_read_access(deps),
492 // For time interpolation
493 dt_interp,
494 dt_vel_dust_cell.get_read_access(deps));
495 }
496
497 inline void complete_event_state(sycl::event e) {
498 aabb_block_lower.complete_event_state(e);
499 aabb_cell_size.complete_event_state(e);
500 vel_dust_cell.complete_event_state(e);
501 dx_v_dust_cell.complete_event_state(e);
502 dy_v_dust_cell.complete_event_state(e);
503 dz_v_dust_cell.complete_event_state(e);
504 dt_vel_dust_cell.complete_event_state(e);
505 }
506 };
507
508} // namespace
509
510template<class Tvec, class TgridVec>
511void shammodels::basegodunov::modules::InterpolateToFaceRho<Tvec, TgridVec>::
512 _impl_evaluate_internal() {
513 StackEntry stack_loc{};
514
515 using OrientedAMRGraph = OrientedAMRGraph<Tvec, TgridVec>;
516
517 static constexpr u32 NsideBlockPow = 1;
518 using AMRBlock = amr::AMRBlock<Tvec, TgridVec, NsideBlockPow>;
519
520 SHAM_ASSERT(AMRBlock::block_size == block_size);
521
522 auto edges = get_edges();
523
524 auto dt_interp = edges.dt_interp.value;
525
526 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_face_xp = edges.rho_face_xp;
527 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_face_xm = edges.rho_face_xm;
528 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_face_yp = edges.rho_face_yp;
529 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_face_ym = edges.rho_face_ym;
530 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_face_zp = edges.rho_face_zp;
531 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_face_zm = edges.rho_face_zm;
532
533 rho_face_xp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xp));
534 rho_face_xm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xm));
535 rho_face_yp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::yp));
536 rho_face_ym.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::ym));
537 rho_face_zp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zp));
538 rho_face_zm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zm));
539
540 auto spans_block_cell_sizes = edges.spans_block_cell_sizes.get_spans();
541 auto spans_cell0block_aabb_lower = edges.spans_cell0block_aabb_lower.get_spans();
542 auto spans_rhos = edges.spans_rhos.get_spans();
543 auto spans_grad_rho = edges.spans_grad_rho.get_spans();
544 auto spans_dt_rho = edges.spans_dt_rho.get_spans();
545
546 using Interp = RhoInterpolate<Tvec, TgridVec, AMRBlock>;
547 auto interpolators
548 = spans_block_cell_sizes.template map<Interp>([&](u64 id, auto &csize) -> Interp {
549 return {
550 spans_cell0block_aabb_lower.get(id),
551 spans_block_cell_sizes.get(id),
552 spans_rhos.get(id),
553 spans_grad_rho.get(id),
554 dt_interp,
555 spans_dt_rho.get(id)};
556 });
557
558 auto graphs_xp = edges.cell_neigh_graph.get_refs_dir(Direction::xp);
559 auto graphs_xm = edges.cell_neigh_graph.get_refs_dir(Direction::xm);
560 auto graphs_yp = edges.cell_neigh_graph.get_refs_dir(Direction::yp);
561 auto graphs_ym = edges.cell_neigh_graph.get_refs_dir(Direction::ym);
562 auto graphs_zp = edges.cell_neigh_graph.get_refs_dir(Direction::zp);
563 auto graphs_zm = edges.cell_neigh_graph.get_refs_dir(Direction::zm);
564
565 shambase::DistributedData<u32> counts_xp
566 = graphs_xp.template map<u32>([&](u64 id, auto &graph) {
567 return graph.get().obj_cnt;
568 });
569 shambase::DistributedData<u32> counts_xm
570 = graphs_xm.template map<u32>([&](u64 id, auto &graph) {
571 return graph.get().obj_cnt;
572 });
573 shambase::DistributedData<u32> counts_yp
574 = graphs_yp.template map<u32>([&](u64 id, auto &graph) {
575 return graph.get().obj_cnt;
576 });
577 shambase::DistributedData<u32> counts_ym
578 = graphs_ym.template map<u32>([&](u64 id, auto &graph) {
579 return graph.get().obj_cnt;
580 });
581 shambase::DistributedData<u32> counts_zp
582 = graphs_zp.template map<u32>([&](u64 id, auto &graph) {
583 return graph.get().obj_cnt;
584 });
585 shambase::DistributedData<u32> counts_zm
586 = graphs_zm.template map<u32>([&](u64 id, auto &graph) {
587 return graph.get().obj_cnt;
588 });
589
590 sham::distributed_data_kernel_call(
591 shamsys::instance::get_compute_scheduler_ptr(),
592 sham::DDMultiRef{graphs_xp, interpolators},
593 sham::DDMultiRef{rho_face_xp.link_fields},
594 counts_xp,
595 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
596 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
597 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
598 });
599 });
600 sham::distributed_data_kernel_call(
601 shamsys::instance::get_compute_scheduler_ptr(),
602 sham::DDMultiRef{graphs_xm, interpolators},
603 sham::DDMultiRef{rho_face_xm.link_fields},
604 counts_xm,
605 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
606 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
607 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
608 });
609 });
610 sham::distributed_data_kernel_call(
611 shamsys::instance::get_compute_scheduler_ptr(),
612 sham::DDMultiRef{graphs_yp, interpolators},
613 sham::DDMultiRef{rho_face_yp.link_fields},
614 counts_yp,
615 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
616 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
617 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
618 });
619 });
620 sham::distributed_data_kernel_call(
621 shamsys::instance::get_compute_scheduler_ptr(),
622 sham::DDMultiRef{graphs_ym, interpolators},
623 sham::DDMultiRef{rho_face_ym.link_fields},
624 counts_ym,
625 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
626 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
627 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
628 });
629 });
630 sham::distributed_data_kernel_call(
631 shamsys::instance::get_compute_scheduler_ptr(),
632 sham::DDMultiRef{graphs_zp, interpolators},
633 sham::DDMultiRef{rho_face_zp.link_fields},
634 counts_zp,
635 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
636 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
637 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
638 });
639 });
640 sham::distributed_data_kernel_call(
641 shamsys::instance::get_compute_scheduler_ptr(),
642 sham::DDMultiRef{graphs_zm, interpolators},
643 sham::DDMultiRef{rho_face_zm.link_fields},
644 counts_zm,
645 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
646 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
647 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
648 });
649 });
650}
651
652template<class Tvec, class TgridVec>
657
658template class shammodels::basegodunov::modules::InterpolateToFaceRho<f64_3, i64_3>;
659
660template<class Tvec, class TgridVec>
661void shammodels::basegodunov::modules::InterpolateToFaceVel<Tvec, TgridVec>::
662 _impl_evaluate_internal() {
663 StackEntry stack_loc{};
664
665 using OrientedAMRGraph = OrientedAMRGraph<Tvec, TgridVec>;
666
667 static constexpr u32 NsideBlockPow = 1;
668 using AMRBlock = amr::AMRBlock<Tvec, TgridVec, NsideBlockPow>;
669
670 SHAM_ASSERT(AMRBlock::block_size == block_size);
671
672 auto edges = get_edges();
673
674 auto dt_interp = edges.dt_interp.value;
675
676 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_face_xp = edges.vel_face_xp;
677 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_face_xm = edges.vel_face_xm;
678 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_face_yp = edges.vel_face_yp;
679 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_face_ym = edges.vel_face_ym;
680 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_face_zp = edges.vel_face_zp;
681 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_face_zm = edges.vel_face_zm;
682
683 vel_face_xp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xp));
684 vel_face_xm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xm));
685 vel_face_yp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::yp));
686 vel_face_ym.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::ym));
687 vel_face_zp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zp));
688 vel_face_zm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zm));
689
690 auto spans_block_cell_sizes = edges.spans_block_cell_sizes.get_spans();
691 auto spans_cell0block_aabb_lower = edges.spans_cell0block_aabb_lower.get_spans();
692 auto spans_vel = edges.spans_vel.get_spans();
693 auto spans_dx_vel = edges.spans_dx_vel.get_spans();
694 auto spans_dy_vel = edges.spans_dy_vel.get_spans();
695 auto spans_dz_vel = edges.spans_dz_vel.get_spans();
696 auto spans_dt_vel = edges.spans_dt_vel.get_spans();
697
698 using Interp = VelInterpolate<Tvec, TgridVec, AMRBlock>;
699 auto interpolators
700 = spans_block_cell_sizes.template map<Interp>([&](u64 id, auto &csize) -> Interp {
701 return {
702 spans_cell0block_aabb_lower.get(id),
703 spans_block_cell_sizes.get(id),
704 spans_vel.get(id),
705 spans_dx_vel.get(id),
706 spans_dy_vel.get(id),
707 spans_dz_vel.get(id),
708 dt_interp,
709 spans_dt_vel.get(id)};
710 });
711
712 auto graphs_xp = edges.cell_neigh_graph.get_refs_dir(Direction::xp);
713 auto graphs_xm = edges.cell_neigh_graph.get_refs_dir(Direction::xm);
714 auto graphs_yp = edges.cell_neigh_graph.get_refs_dir(Direction::yp);
715 auto graphs_ym = edges.cell_neigh_graph.get_refs_dir(Direction::ym);
716 auto graphs_zp = edges.cell_neigh_graph.get_refs_dir(Direction::zp);
717 auto graphs_zm = edges.cell_neigh_graph.get_refs_dir(Direction::zm);
718
719 shambase::DistributedData<u32> counts_xp
720 = graphs_xp.template map<u32>([&](u64 id, auto &graph) {
721 return graph.get().obj_cnt;
722 });
723 shambase::DistributedData<u32> counts_xm
724 = graphs_xm.template map<u32>([&](u64 id, auto &graph) {
725 return graph.get().obj_cnt;
726 });
727 shambase::DistributedData<u32> counts_yp
728 = graphs_yp.template map<u32>([&](u64 id, auto &graph) {
729 return graph.get().obj_cnt;
730 });
731 shambase::DistributedData<u32> counts_ym
732 = graphs_ym.template map<u32>([&](u64 id, auto &graph) {
733 return graph.get().obj_cnt;
734 });
735 shambase::DistributedData<u32> counts_zp
736 = graphs_zp.template map<u32>([&](u64 id, auto &graph) {
737 return graph.get().obj_cnt;
738 });
739 shambase::DistributedData<u32> counts_zm
740 = graphs_zm.template map<u32>([&](u64 id, auto &graph) {
741 return graph.get().obj_cnt;
742 });
743
744 sham::distributed_data_kernel_call(
745 shamsys::instance::get_compute_scheduler_ptr(),
746 sham::DDMultiRef{graphs_xp, interpolators},
747 sham::DDMultiRef{vel_face_xp.link_fields},
748 counts_xp,
749 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
750 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
751 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
752 });
753 });
754 sham::distributed_data_kernel_call(
755 shamsys::instance::get_compute_scheduler_ptr(),
756 sham::DDMultiRef{graphs_xm, interpolators},
757 sham::DDMultiRef{vel_face_xm.link_fields},
758 counts_xm,
759 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
760 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
761 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
762 });
763 });
764 sham::distributed_data_kernel_call(
765 shamsys::instance::get_compute_scheduler_ptr(),
766 sham::DDMultiRef{graphs_yp, interpolators},
767 sham::DDMultiRef{vel_face_yp.link_fields},
768 counts_yp,
769 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
770 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
771 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
772 });
773 });
774 sham::distributed_data_kernel_call(
775 shamsys::instance::get_compute_scheduler_ptr(),
776 sham::DDMultiRef{graphs_ym, interpolators},
777 sham::DDMultiRef{vel_face_ym.link_fields},
778 counts_ym,
779 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
780 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
781 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
782 });
783 });
784 sham::distributed_data_kernel_call(
785 shamsys::instance::get_compute_scheduler_ptr(),
786 sham::DDMultiRef{graphs_zp, interpolators},
787 sham::DDMultiRef{vel_face_zp.link_fields},
788 counts_zp,
789 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
790 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
791 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
792 });
793 });
794 sham::distributed_data_kernel_call(
795 shamsys::instance::get_compute_scheduler_ptr(),
796 sham::DDMultiRef{graphs_zm, interpolators},
797 sham::DDMultiRef{vel_face_zm.link_fields},
798 counts_zm,
799 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
800 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
801 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
802 });
803 });
804}
805
806template<class Tvec, class TgridVec>
811
812template class shammodels::basegodunov::modules::InterpolateToFaceVel<f64_3, i64_3>;
813
814template<class Tvec, class TgridVec>
815void shammodels::basegodunov::modules::InterpolateToFacePress<Tvec, TgridVec>::
816 _impl_evaluate_internal() {
817 StackEntry stack_loc{};
818
819 using OrientedAMRGraph = OrientedAMRGraph<Tvec, TgridVec>;
820
821 static constexpr u32 NsideBlockPow = 1;
822 using AMRBlock = amr::AMRBlock<Tvec, TgridVec, NsideBlockPow>;
823
824 SHAM_ASSERT(AMRBlock::block_size == block_size);
825
826 auto edges = get_edges();
827
828 auto dt_interp = edges.dt_interp.value;
829
830 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &press_face_xp = edges.press_face_xp;
831 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &press_face_xm = edges.press_face_xm;
832 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &press_face_yp = edges.press_face_yp;
833 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &press_face_ym = edges.press_face_ym;
834 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &press_face_zp = edges.press_face_zp;
835 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &press_face_zm = edges.press_face_zm;
836
837 press_face_xp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xp));
838 press_face_xm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xm));
839 press_face_yp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::yp));
840 press_face_ym.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::ym));
841 press_face_zp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zp));
842 press_face_zm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zm));
843
844 auto spans_block_cell_sizes = edges.spans_block_cell_sizes.get_spans();
845 auto spans_cell0block_aabb_lower = edges.spans_cell0block_aabb_lower.get_spans();
846 auto spans_press = edges.spans_press.get_spans();
847 auto spans_grad_P = edges.spans_grad_P.get_spans();
848 auto spans_dt_press = edges.spans_dt_press.get_spans();
849
850 using Interp = PressInterpolate<Tvec, TgridVec, AMRBlock>;
851 auto interpolators
852 = spans_block_cell_sizes.template map<Interp>([&](u64 id, auto &csize) -> Interp {
853 return {
854 spans_cell0block_aabb_lower.get(id),
855 spans_block_cell_sizes.get(id),
856 spans_press.get(id),
857 spans_grad_P.get(id),
858 dt_interp,
859 spans_dt_press.get(id)};
860 });
861
862 auto graphs_xp = edges.cell_neigh_graph.get_refs_dir(Direction::xp);
863 auto graphs_xm = edges.cell_neigh_graph.get_refs_dir(Direction::xm);
864 auto graphs_yp = edges.cell_neigh_graph.get_refs_dir(Direction::yp);
865 auto graphs_ym = edges.cell_neigh_graph.get_refs_dir(Direction::ym);
866 auto graphs_zp = edges.cell_neigh_graph.get_refs_dir(Direction::zp);
867 auto graphs_zm = edges.cell_neigh_graph.get_refs_dir(Direction::zm);
868
869 shambase::DistributedData<u32> counts_xp
870 = graphs_xp.template map<u32>([&](u64 id, auto &graph) {
871 return graph.get().obj_cnt;
872 });
873 shambase::DistributedData<u32> counts_xm
874 = graphs_xm.template map<u32>([&](u64 id, auto &graph) {
875 return graph.get().obj_cnt;
876 });
877 shambase::DistributedData<u32> counts_yp
878 = graphs_yp.template map<u32>([&](u64 id, auto &graph) {
879 return graph.get().obj_cnt;
880 });
881 shambase::DistributedData<u32> counts_ym
882 = graphs_ym.template map<u32>([&](u64 id, auto &graph) {
883 return graph.get().obj_cnt;
884 });
885 shambase::DistributedData<u32> counts_zp
886 = graphs_zp.template map<u32>([&](u64 id, auto &graph) {
887 return graph.get().obj_cnt;
888 });
889 shambase::DistributedData<u32> counts_zm
890 = graphs_zm.template map<u32>([&](u64 id, auto &graph) {
891 return graph.get().obj_cnt;
892 });
893
894 sham::distributed_data_kernel_call(
895 shamsys::instance::get_compute_scheduler_ptr(),
896 sham::DDMultiRef{graphs_xp, interpolators},
897 sham::DDMultiRef{press_face_xp.link_fields},
898 counts_xp,
899 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
900 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
901 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
902 });
903 });
904 sham::distributed_data_kernel_call(
905 shamsys::instance::get_compute_scheduler_ptr(),
906 sham::DDMultiRef{graphs_xm, interpolators},
907 sham::DDMultiRef{press_face_xm.link_fields},
908 counts_xm,
909 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
910 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
911 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
912 });
913 });
914 sham::distributed_data_kernel_call(
915 shamsys::instance::get_compute_scheduler_ptr(),
916 sham::DDMultiRef{graphs_yp, interpolators},
917 sham::DDMultiRef{press_face_yp.link_fields},
918 counts_yp,
919 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
920 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
921 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
922 });
923 });
924 sham::distributed_data_kernel_call(
925 shamsys::instance::get_compute_scheduler_ptr(),
926 sham::DDMultiRef{graphs_ym, interpolators},
927 sham::DDMultiRef{press_face_ym.link_fields},
928 counts_ym,
929 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
930 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
931 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
932 });
933 });
934 sham::distributed_data_kernel_call(
935 shamsys::instance::get_compute_scheduler_ptr(),
936 sham::DDMultiRef{graphs_zp, interpolators},
937 sham::DDMultiRef{press_face_zp.link_fields},
938 counts_zp,
939 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
940 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
941 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
942 });
943 });
944 sham::distributed_data_kernel_call(
945 shamsys::instance::get_compute_scheduler_ptr(),
946 sham::DDMultiRef{graphs_zm, interpolators},
947 sham::DDMultiRef{press_face_zm.link_fields},
948 counts_zm,
949 [](u32 id_a, auto link_iter, auto compute, auto acc_link_field) {
950 link_iter.for_each_object_link_id(id_a, [&](u32 id_b, u32 link_id) {
951 acc_link_field[link_id] = compute.get_link_field_val(id_a, id_b);
952 });
953 });
954}
955
956template<class Tvec, class TgridVec>
961
962template class shammodels::basegodunov::modules::InterpolateToFacePress<f64_3, i64_3>;
963
964template<class Tvec, class TgridVec>
965void shammodels::basegodunov::modules::InterpolateToFaceRhoDust<Tvec, TgridVec>::
966 _impl_evaluate_internal() {
967 __shamrock_stack_entry();
968
969 using OrientedAMRGraph = OrientedAMRGraph<Tvec, TgridVec>;
970
971 static constexpr u32 NsideBlockPow = 1;
972 using AMRBlock = amr::AMRBlock<Tvec, TgridVec, NsideBlockPow>;
973
974 SHAM_ASSERT(AMRBlock::block_size == block_size);
975
976 auto edges = get_edges();
977
978 auto dt_interp = edges.dt_interp.value;
979 auto ndust = this->ndust;
980
981 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_dust_face_xp
982 = edges.rho_dust_face_xp;
983 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_dust_face_xm
984 = edges.rho_dust_face_xm;
985 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_dust_face_yp
986 = edges.rho_dust_face_yp;
987 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_dust_face_ym
988 = edges.rho_dust_face_ym;
989 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_dust_face_zp
990 = edges.rho_dust_face_zp;
991 solvergraph::NeighGraphLinkFieldEdge<std::array<Tscal, 2>> &rho_dust_face_zm
992 = edges.rho_dust_face_zm;
993
994 rho_dust_face_xp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xp));
995 rho_dust_face_xm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xm));
996 rho_dust_face_yp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::yp));
997 rho_dust_face_ym.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::ym));
998 rho_dust_face_zp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zp));
999 rho_dust_face_zm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zm));
1000
1001 auto spans_block_cell_sizes = edges.spans_block_cell_sizes.get_spans();
1002 auto spans_cell0block_aabb_lower = edges.spans_cell0block_aabb_lower.get_spans();
1003 auto spans_rhos_dust = edges.spans_rhos_dust.get_spans();
1004 auto spans_grad_rho_dust = edges.spans_grad_rho_dust.get_spans();
1005 auto spans_dt_rho_dust = edges.spans_dt_rho_dust.get_spans();
1006
1007 using Interp = RhoDustInterpolate<Tvec, TgridVec, AMRBlock>;
1008 auto interpolators
1009 = spans_block_cell_sizes.template map<Interp>([&](u64 id, auto &csize) -> Interp {
1010 return {
1011 ndust,
1012 spans_cell0block_aabb_lower.get(id),
1013 spans_block_cell_sizes.get(id),
1014 spans_rhos_dust.get(id),
1015 spans_grad_rho_dust.get(id),
1016 dt_interp,
1017 spans_dt_rho_dust.get(id)};
1018 });
1019
1020 auto graphs_xp = edges.cell_neigh_graph.get_refs_dir(Direction::xp);
1021 auto graphs_xm = edges.cell_neigh_graph.get_refs_dir(Direction::xm);
1022 auto graphs_yp = edges.cell_neigh_graph.get_refs_dir(Direction::yp);
1023 auto graphs_ym = edges.cell_neigh_graph.get_refs_dir(Direction::ym);
1024 auto graphs_zp = edges.cell_neigh_graph.get_refs_dir(Direction::zp);
1025 auto graphs_zm = edges.cell_neigh_graph.get_refs_dir(Direction::zm);
1026
1027 shambase::DistributedData<u32> counts_xp
1028 = graphs_xp.template map<u32>([&](u64 id, auto &graph) {
1029 return graph.get().obj_cnt * ndust;
1030 });
1031 shambase::DistributedData<u32> counts_xm
1032 = graphs_xm.template map<u32>([&](u64 id, auto &graph) {
1033 return graph.get().obj_cnt * ndust;
1034 });
1035 shambase::DistributedData<u32> counts_yp
1036 = graphs_yp.template map<u32>([&](u64 id, auto &graph) {
1037 return graph.get().obj_cnt * ndust;
1038 });
1039 shambase::DistributedData<u32> counts_ym
1040 = graphs_ym.template map<u32>([&](u64 id, auto &graph) {
1041 return graph.get().obj_cnt * ndust;
1042 });
1043 shambase::DistributedData<u32> counts_zp
1044 = graphs_zp.template map<u32>([&](u64 id, auto &graph) {
1045 return graph.get().obj_cnt * ndust;
1046 });
1047 shambase::DistributedData<u32> counts_zm
1048 = graphs_zm.template map<u32>([&](u64 id, auto &graph) {
1049 return graph.get().obj_cnt * ndust;
1050 });
1051
1052 sham::distributed_data_kernel_call(
1053 shamsys::instance::get_compute_scheduler_ptr(),
1054 sham::DDMultiRef{graphs_xp, interpolators},
1055 sham::DDMultiRef{rho_dust_face_xp.link_fields},
1056 counts_xp,
1057 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1058 const u32 id_cell_a = idvar_a / ndust;
1059 const u32 nvar_loc = idvar_a % ndust;
1060 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1061 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1062 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1063 });
1064 });
1065 sham::distributed_data_kernel_call(
1066 shamsys::instance::get_compute_scheduler_ptr(),
1067 sham::DDMultiRef{graphs_xm, interpolators},
1068 sham::DDMultiRef{rho_dust_face_xm.link_fields},
1069 counts_xm,
1070 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1071 const u32 id_cell_a = idvar_a / ndust;
1072 const u32 nvar_loc = idvar_a % ndust;
1073 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1074 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1075 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1076 });
1077 });
1078 sham::distributed_data_kernel_call(
1079 shamsys::instance::get_compute_scheduler_ptr(),
1080 sham::DDMultiRef{graphs_yp, interpolators},
1081 sham::DDMultiRef{rho_dust_face_yp.link_fields},
1082 counts_yp,
1083 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1084 const u32 id_cell_a = idvar_a / ndust;
1085 const u32 nvar_loc = idvar_a % ndust;
1086 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1087 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1088 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1089 });
1090 });
1091 sham::distributed_data_kernel_call(
1092 shamsys::instance::get_compute_scheduler_ptr(),
1093 sham::DDMultiRef{graphs_ym, interpolators},
1094 sham::DDMultiRef{rho_dust_face_ym.link_fields},
1095 counts_ym,
1096 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1097 const u32 id_cell_a = idvar_a / ndust;
1098 const u32 nvar_loc = idvar_a % ndust;
1099 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1100 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1101 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1102 });
1103 });
1104 sham::distributed_data_kernel_call(
1105 shamsys::instance::get_compute_scheduler_ptr(),
1106 sham::DDMultiRef{graphs_zp, interpolators},
1107 sham::DDMultiRef{rho_dust_face_zp.link_fields},
1108 counts_zp,
1109 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1110 const u32 id_cell_a = idvar_a / ndust;
1111 const u32 nvar_loc = idvar_a % ndust;
1112 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1113 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1114 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1115 });
1116 });
1117 sham::distributed_data_kernel_call(
1118 shamsys::instance::get_compute_scheduler_ptr(),
1119 sham::DDMultiRef{graphs_zm, interpolators},
1120 sham::DDMultiRef{rho_dust_face_zm.link_fields},
1121 counts_zm,
1122 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1123 const u32 id_cell_a = idvar_a / ndust;
1124 const u32 nvar_loc = idvar_a % ndust;
1125 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1126 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1127 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1128 });
1129 });
1130}
1131
1132template<class Tvec, class TgridVec>
1137
1138template class shammodels::basegodunov::modules::InterpolateToFaceRhoDust<f64_3, i64_3>;
1139
1140template<class Tvec, class TgridVec>
1141void shammodels::basegodunov::modules::InterpolateToFaceVelDust<Tvec, TgridVec>::
1142 _impl_evaluate_internal() {
1143 StackEntry stack_loc{};
1144
1145 using OrientedAMRGraph = OrientedAMRGraph<Tvec, TgridVec>;
1146
1147 static constexpr u32 NsideBlockPow = 1;
1148 using AMRBlock = amr::AMRBlock<Tvec, TgridVec, NsideBlockPow>;
1149
1150 SHAM_ASSERT(AMRBlock::block_size == block_size);
1151
1152 auto edges = get_edges();
1153
1154 auto dt_interp = edges.dt_interp.value;
1155 auto ndust = this->ndust;
1156
1157 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_dust_face_xp
1158 = edges.vel_dust_face_xp;
1159 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_dust_face_xm
1160 = edges.vel_dust_face_xm;
1161 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_dust_face_yp
1162 = edges.vel_dust_face_yp;
1163 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_dust_face_ym
1164 = edges.vel_dust_face_ym;
1165 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_dust_face_zp
1166 = edges.vel_dust_face_zp;
1167 solvergraph::NeighGraphLinkFieldEdge<std::array<Tvec, 2>> &vel_dust_face_zm
1168 = edges.vel_dust_face_zm;
1169
1170 vel_dust_face_xp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xp));
1171 vel_dust_face_xm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::xm));
1172 vel_dust_face_yp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::yp));
1173 vel_dust_face_ym.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::ym));
1174 vel_dust_face_zp.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zp));
1175 vel_dust_face_zm.resize_according_to(edges.cell_neigh_graph.get_refs_dir(Direction::zm));
1176
1177 auto spans_block_cell_sizes = edges.spans_block_cell_sizes.get_spans();
1178 auto spans_cell0block_aabb_lower = edges.spans_cell0block_aabb_lower.get_spans();
1179 auto spans_vel_dust = edges.spans_vel_dust.get_spans();
1180 auto spans_dx_vel_dust = edges.spans_dx_vel_dust.get_spans();
1181 auto spans_dy_vel_dust = edges.spans_dy_vel_dust.get_spans();
1182 auto spans_dz_vel_dust = edges.spans_dz_vel_dust.get_spans();
1183 auto spans_dt_vel_dust = edges.spans_dt_vel_dust.get_spans();
1184
1185 using Interp = VelDustInterpolate<Tvec, TgridVec, AMRBlock>;
1186 auto interpolators
1187 = spans_block_cell_sizes.template map<Interp>([&](u64 id, auto &csize) -> Interp {
1188 return {
1189 ndust,
1190 spans_cell0block_aabb_lower.get(id),
1191 spans_block_cell_sizes.get(id),
1192 spans_vel_dust.get(id),
1193 spans_dx_vel_dust.get(id),
1194 spans_dy_vel_dust.get(id),
1195 spans_dz_vel_dust.get(id),
1196 dt_interp,
1197 spans_dt_vel_dust.get(id)};
1198 });
1199
1200 auto graphs_xp = edges.cell_neigh_graph.get_refs_dir(Direction::xp);
1201 auto graphs_xm = edges.cell_neigh_graph.get_refs_dir(Direction::xm);
1202 auto graphs_yp = edges.cell_neigh_graph.get_refs_dir(Direction::yp);
1203 auto graphs_ym = edges.cell_neigh_graph.get_refs_dir(Direction::ym);
1204 auto graphs_zp = edges.cell_neigh_graph.get_refs_dir(Direction::zp);
1205 auto graphs_zm = edges.cell_neigh_graph.get_refs_dir(Direction::zm);
1206
1207 shambase::DistributedData<u32> counts_xp
1208 = graphs_xp.template map<u32>([&](u64 id, auto &graph) {
1209 return graph.get().obj_cnt * ndust;
1210 });
1211 shambase::DistributedData<u32> counts_xm
1212 = graphs_xm.template map<u32>([&](u64 id, auto &graph) {
1213 return graph.get().obj_cnt * ndust;
1214 });
1215 shambase::DistributedData<u32> counts_yp
1216 = graphs_yp.template map<u32>([&](u64 id, auto &graph) {
1217 return graph.get().obj_cnt * ndust;
1218 });
1219 shambase::DistributedData<u32> counts_ym
1220 = graphs_ym.template map<u32>([&](u64 id, auto &graph) {
1221 return graph.get().obj_cnt * ndust;
1222 });
1223 shambase::DistributedData<u32> counts_zp
1224 = graphs_zp.template map<u32>([&](u64 id, auto &graph) {
1225 return graph.get().obj_cnt * ndust;
1226 });
1227 shambase::DistributedData<u32> counts_zm
1228 = graphs_zm.template map<u32>([&](u64 id, auto &graph) {
1229 return graph.get().obj_cnt * ndust;
1230 });
1231
1232 sham::distributed_data_kernel_call(
1233 shamsys::instance::get_compute_scheduler_ptr(),
1234 sham::DDMultiRef{graphs_xp, interpolators},
1235 sham::DDMultiRef{vel_dust_face_xp.link_fields},
1236 counts_xp,
1237 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1238 const u32 id_cell_a = idvar_a / ndust;
1239 const u32 nvar_loc = idvar_a % ndust;
1240 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1241 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1242 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1243 });
1244 });
1245 sham::distributed_data_kernel_call(
1246 shamsys::instance::get_compute_scheduler_ptr(),
1247 sham::DDMultiRef{graphs_xm, interpolators},
1248 sham::DDMultiRef{vel_dust_face_xm.link_fields},
1249 counts_xm,
1250 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1251 const u32 id_cell_a = idvar_a / ndust;
1252 const u32 nvar_loc = idvar_a % ndust;
1253 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1254 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1255 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1256 });
1257 });
1258 sham::distributed_data_kernel_call(
1259 shamsys::instance::get_compute_scheduler_ptr(),
1260 sham::DDMultiRef{graphs_yp, interpolators},
1261 sham::DDMultiRef{vel_dust_face_yp.link_fields},
1262 counts_yp,
1263 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1264 const u32 id_cell_a = idvar_a / ndust;
1265 const u32 nvar_loc = idvar_a % ndust;
1266 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1267 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1268 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1269 });
1270 });
1271 sham::distributed_data_kernel_call(
1272 shamsys::instance::get_compute_scheduler_ptr(),
1273 sham::DDMultiRef{graphs_ym, interpolators},
1274 sham::DDMultiRef{vel_dust_face_ym.link_fields},
1275 counts_ym,
1276 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1277 const u32 id_cell_a = idvar_a / ndust;
1278 const u32 nvar_loc = idvar_a % ndust;
1279 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1280 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1281 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1282 });
1283 });
1284 sham::distributed_data_kernel_call(
1285 shamsys::instance::get_compute_scheduler_ptr(),
1286 sham::DDMultiRef{graphs_zp, interpolators},
1287 sham::DDMultiRef{vel_dust_face_zp.link_fields},
1288 counts_zp,
1289 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1290 const u32 id_cell_a = idvar_a / ndust;
1291 const u32 nvar_loc = idvar_a % ndust;
1292 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1293 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1294 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1295 });
1296 });
1297 sham::distributed_data_kernel_call(
1298 shamsys::instance::get_compute_scheduler_ptr(),
1299 sham::DDMultiRef{graphs_zm, interpolators},
1300 sham::DDMultiRef{vel_dust_face_zm.link_fields},
1301 counts_zm,
1302 [ndust](u32 idvar_a, auto link_iter, auto compute, auto acc_link_field) {
1303 const u32 id_cell_a = idvar_a / ndust;
1304 const u32 nvar_loc = idvar_a % ndust;
1305 link_iter.for_each_object_link_id(id_cell_a, [&](u32 id_cell_b, u32 link_id) {
1306 acc_link_field[link_id * ndust + nvar_loc] = compute.get_link_field_val(
1307 id_cell_a * ndust + nvar_loc, id_cell_b * ndust + nvar_loc);
1308 });
1309 });
1310}
1311
1312template<class Tvec, class TgridVec>
1317
1318template class shammodels::basegodunov::modules::InterpolateToFaceVelDust<f64_3, i64_3>;
AMRBlock.hpp
utility to manipulate AMR blocks
u32
std::uint32_t u32
32 bit unsigned integer
Definition aliases_int.hpp:27
u64
std::uint64_t u64
64 bit unsigned integer
Definition aliases_int.hpp:26
assert.hpp
Shamrock assertion utility.
SHAM_ASSERT
#define SHAM_ASSERT(x)
Shorthand for SHAM_ASSERT_NAMED without a message.
Definition assert.hpp:67
shambase::DistributedData
Represents a collection of objects distributed across patches identified by a u64 id.
Definition DistributedData.hpp:43
shammodels::basegodunov::modules::InterpolateToFacePress::_impl_get_tex
virtual std::string _impl_get_tex() const
get the tex of the node
Definition InterpolateToFace.cpp:958
shammodels::basegodunov::modules::InterpolateToFaceRhoDust::_impl_get_tex
virtual std::string _impl_get_tex() const
get the tex of the node
Definition InterpolateToFace.cpp:1134
shammodels::basegodunov::modules::InterpolateToFaceRho::_impl_get_tex
virtual std::string _impl_get_tex() const
get the tex of the node
Definition InterpolateToFace.cpp:653
shammodels::basegodunov::modules::InterpolateToFaceVelDust::_impl_get_tex
virtual std::string _impl_get_tex() const
get the tex of the node
Definition InterpolateToFace.cpp:1314
shammodels::basegodunov::modules::InterpolateToFaceVel::_impl_get_tex
virtual std::string _impl_get_tex() const
get the tex of the node
Definition InterpolateToFace.cpp:807
sham::distributed_data_kernel_call
void distributed_data_kernel_call(sham::DeviceScheduler_ptr dev_sched, RefIn in, RefOut in_out, const shambase::DistributedData< index_t > &thread_counts, Functor &&func)
A variant of sham::kernel_call for distributed data.
Definition kernel_call_distrib.hpp:79
__shamrock_stack_entry
#define __shamrock_stack_entry()
Macro to create a stack entry.
Definition stacktrace.hpp:257
StackEntry
shambase::details::BasicStackEntry StackEntry
Alias for shambase::details::BasicStackEntry.
Definition stacktrace.hpp:242
sham::DDMultiRef
A variant of sham::MultiRef for distributed data.
Definition kernel_call_distrib.hpp:33
shammath::AABB::get_center
T get_center() const noexcept
Returns the center of the AABB.
Definition AABB.hpp:174
shammath::AABB::get_intersect
AABB get_intersect(AABB other) const noexcept
Compute the intersection of two AABB.
Definition AABB.hpp:234
shammodels::amr::AMRBlock
utility class to handle AMR blocks
Definition AMRBlock.hpp:35
Generated on for Shamrock by doxygen 1.15.0