Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
ComputeEos.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
18
19#include "shambase/DistributedData.hpp"
20#include "shambase/exception.hpp"
21#include "shambackends/kernel_call.hpp"
22#include "shammath/sphkernels.hpp"
23#include "shammodels/sph/math/density.hpp"
24#include "shammodels/sph/modules/ComputeEos.hpp"
25#include "shamphys/eos.hpp"
26#include "shamrock/patch/PatchDataLayer.hpp"
27#include "shamrock/scheduler/SchedulerUtility.hpp"
28#include "shamsys/legacy/log.hpp"
29
30template<class Tscal>
31struct RhoGetterBase {
32 sham::DeviceBuffer<Tscal> &buf_h;
33 Tscal pmass;
34 Tscal hfact;
35
36 struct accessed {
37 const Tscal *h;
38 Tscal pmass;
39 Tscal hfact;
40
41 Tscal operator()(u32 i) const {
42 using namespace shamrock::sph;
43 return rho_h(pmass, h[i], hfact);
44 }
45 };
46
47 accessed get_read_access(sham::EventList &depends_list) {
48 auto h = buf_h.get_read_access(depends_list);
49 return accessed{h, pmass, hfact};
50 }
51
52 void complete_event_state(sycl::event e) { buf_h.complete_event_state(e); }
53};
54
55template<class Tscal>
56struct RhoGetterMonofluid {
57 sham::DeviceBuffer<Tscal> &buf_h;
58 sham::DeviceBuffer<Tscal> &buf_epsilon;
59 u32 nvar_dust;
60 Tscal pmass;
61 Tscal hfact;
62
63 struct accessed {
64 const Tscal *h;
65 const Tscal *buf_epsilon;
66 u32 nvar_dust;
67 Tscal pmass;
68 Tscal hfact;
69
70 Tscal operator()(u32 i) const {
71
72 Tscal epsilon_sum = 0;
73 for (u32 j = 0; j < nvar_dust; j++) {
74 epsilon_sum += buf_epsilon[i * nvar_dust + j];
75 }
76
77 using namespace shamrock::sph;
78 return (1 - epsilon_sum) * rho_h(pmass, h[i], hfact);
79 }
80 };
81
82 accessed get_read_access(sham::EventList &depends_list) {
83 auto h = buf_h.get_read_access(depends_list);
84 auto epsilon = buf_epsilon.get_read_access(depends_list);
85
86 return accessed{h, epsilon, nvar_dust, pmass, hfact};
87 }
88
89 void complete_event_state(sycl::event e) { buf_h.complete_event_state(e); }
90};
91
92template<class Tscal>
93struct RhoGetterSJ {
94 sham::DeviceBuffer<Tscal> &buf_h;
95 sham::DeviceBuffer<Tscal> &buf_s_j;
96 u32 nvar_dust;
97 Tscal pmass;
98 Tscal hfact;
99
100 struct accessed {
101 const Tscal *h;
102 const Tscal *buf_s_j;
103 u32 nvar_dust;
104 Tscal pmass;
105 Tscal hfact;
106
107 Tscal operator()(u32 i) const {
108 using namespace shamrock::sph;
109 Tscal rho = rho_h(pmass, h[i], hfact);
110
111 Tscal epsilon_sum = 0;
112 for (u32 j = 0; j < nvar_dust; j++) {
113 Tscal s_j = buf_s_j[i * nvar_dust + j];
114 epsilon_sum += s_j * s_j / rho;
115 }
116
117 return rho * (1 - epsilon_sum);
118 }
119 };
120
121 accessed get_read_access(sham::EventList &depends_list) {
122 auto h = buf_h.get_read_access(depends_list);
123 auto s_j = buf_s_j.get_read_access(depends_list);
124
125 return accessed{h, s_j, nvar_dust, pmass, hfact};
126 }
127
128 void complete_event_state(sycl::event e) {
129 buf_h.complete_event_state(e);
130 buf_s_j.complete_event_state(e);
131 }
132};
133
134template<class Tvec, template<class> class SPHKernel>
135template<class RhoGetGen>
136void shammodels::sph::modules::ComputeEos<Tvec, SPHKernel>::compute_eos_internal(
137 RhoGetGen &&rho_getter_gen) {
138
139 shamrock::patch::PatchDataLayerLayout &ghost_layout
140 = shambase::get_check_ref(storage.ghost_layout.get());
141 u32 iuint_interf = ghost_layout.get_field_idx<Tscal>("uint");
142
143 using namespace shamrock;
144 using namespace shamrock::patch;
145
146 using SolverConfigEOS = typename Config::EOSConfig;
147 using SolverEOS_Isothermal = typename SolverConfigEOS::Isothermal;
148 using SolverEOS_Adiabatic = typename SolverConfigEOS::Adiabatic;
149 using SolverEOS_Polytropic = typename SolverConfigEOS::Polytropic;
150 using SolverEOS_LocallyIsothermal = typename SolverConfigEOS::LocallyIsothermal;
151 using SolverEOS_LocallyIsothermalLP07 = typename SolverConfigEOS::LocallyIsothermalLP07;
152 using SolverEOS_LocallyIsothermalFA2014 = typename SolverConfigEOS::LocallyIsothermalFA2014;
153 using SolverEOS_LocallyIsothermalFA2014Extended =
154 typename SolverConfigEOS::LocallyIsothermalFA2014Extended;
155 using SolverEOS_Fermi = typename SolverConfigEOS::Fermi;
156
157 sham::DeviceQueue &q = shamsys::instance::get_compute_scheduler().get_queue();
158
159 if (SolverEOS_Isothermal *eos_config
160 = std::get_if<SolverEOS_Isothermal>(&solver_config.eos_config.config)) {
161
162 using EOS = shamphys::EOS_Isothermal<Tscal>;
163
164 storage.merged_patchdata_ghost.get().for_each([&](u64 id, PatchDataLayer &mpdat) {
165 sham::DeviceBuffer<Tscal> &buf_P
166 = shambase::get_check_ref(storage.pressure).get_field(id).get_buf();
167 sham::DeviceBuffer<Tscal> &buf_cs
168 = shambase::get_check_ref(storage.soundspeed).get_field(id).get_buf();
169 auto rho_getter = rho_getter_gen(mpdat);
170
171 u32 total_elements
172 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes.get(id);
173
174 sham::kernel_call(
175 q,
176 sham::MultiRef{rho_getter},
177 sham::MultiRef{buf_P, buf_cs},
178 total_elements,
179 [cs_cfg
180 = eos_config->cs](u32 i, auto rho, Tscal *__restrict P, Tscal *__restrict cs) {
181 using namespace shamrock::sph;
182 Tscal rho_a = rho(i);
183 Tscal P_a = EOS::pressure(cs_cfg, rho_a);
184 P[i] = P_a;
185 cs[i] = cs_cfg;
186 });
187 });
188 } else if (
189 SolverEOS_Adiabatic *eos_config
190 = std::get_if<SolverEOS_Adiabatic>(&solver_config.eos_config.config)) {
191
192 using EOS = shamphys::EOS_Adiabatic<Tscal>;
193
194 storage.merged_patchdata_ghost.get().for_each([&](u64 id, PatchDataLayer &mpdat) {
195 sham::DeviceBuffer<Tscal> &buf_P
196 = shambase::get_check_ref(storage.pressure).get_field(id).get_buf();
197 sham::DeviceBuffer<Tscal> &buf_cs
198 = shambase::get_check_ref(storage.soundspeed).get_field(id).get_buf();
199 sham::DeviceBuffer<Tscal> &buf_uint = mpdat.get_field_buf_ref<Tscal>(iuint_interf);
200 auto rho_getter = rho_getter_gen(mpdat);
201
202 u32 total_elements
203 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes.get(id);
204
205 sham::kernel_call(
206 q,
207 sham::MultiRef{rho_getter, buf_uint},
208 sham::MultiRef{buf_P, buf_cs},
209 total_elements,
210 [gamma = eos_config->gamma](
211 u32 i,
212 auto rho,
213 const Tscal *__restrict U,
214 Tscal *__restrict P,
215 Tscal *__restrict cs) {
216 using namespace shamrock::sph;
217 Tscal rho_a = rho(i);
218 Tscal P_a = EOS::pressure(gamma, rho_a, U[i]);
219 Tscal cs_a = EOS::cs_from_p(gamma, rho_a, P_a);
220 P[i] = P_a;
221 cs[i] = cs_a;
222 });
223 });
224
225 } else if (
226 SolverEOS_Polytropic *eos_config
227 = std::get_if<SolverEOS_Polytropic>(&solver_config.eos_config.config)) {
228
229 using EOS = shamphys::EOS_Polytropic<Tscal>;
230
231 storage.merged_patchdata_ghost.get().for_each([&](u64 id, PatchDataLayer &mpdat) {
232 sham::DeviceBuffer<Tscal> &buf_P
233 = shambase::get_check_ref(storage.pressure).get_field(id).get_buf();
234 sham::DeviceBuffer<Tscal> &buf_cs
235 = shambase::get_check_ref(storage.soundspeed).get_field(id).get_buf();
236 auto rho_getter = rho_getter_gen(mpdat);
237
238 u32 total_elements
239 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes.get(id);
240
241 sham::kernel_call(
242 q,
243 sham::MultiRef{rho_getter},
244 sham::MultiRef{buf_P, buf_cs},
245 total_elements,
246 [K = eos_config->K, gamma = eos_config->gamma](
247 u32 i, auto rho, Tscal *__restrict P, Tscal *__restrict cs) {
248 using namespace shamrock::sph;
249 Tscal rho_a = rho(i);
250 Tscal P_a = EOS::pressure(gamma, K, rho_a);
251 Tscal cs_a = EOS::soundspeed(gamma, K, rho_a);
252 P[i] = P_a;
253 cs[i] = cs_a;
254 });
255 });
256
257 } else if (
258 SolverEOS_LocallyIsothermal *eos_config
259 = std::get_if<SolverEOS_LocallyIsothermal>(&solver_config.eos_config.config)) {
260
261 using EOS = shamphys::EOS_LocallyIsothermal<Tscal>;
262
263 u32 isoundspeed_interf = ghost_layout.get_field_idx<Tscal>("soundspeed");
264
265 storage.merged_patchdata_ghost.get().for_each([&](u64 id, PatchDataLayer &mpdat) {
266 sham::DeviceBuffer<Tscal> &buf_P
267 = shambase::get_check_ref(storage.pressure).get_field(id).get_buf();
268 sham::DeviceBuffer<Tscal> &buf_cs
269 = shambase::get_check_ref(storage.soundspeed).get_field(id).get_buf();
270 sham::DeviceBuffer<Tscal> &buf_uint = mpdat.get_field_buf_ref<Tscal>(iuint_interf);
271 auto rho_getter = rho_getter_gen(mpdat);
272 sham::DeviceBuffer<Tscal> &buf_cs0 = mpdat.get_field_buf_ref<Tscal>(isoundspeed_interf);
273
274 u32 total_elements
275 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes.get(id);
276
277 sham::kernel_call(
278 q,
279 sham::MultiRef{rho_getter, buf_uint, buf_cs0},
280 sham::MultiRef{buf_P, buf_cs},
281 total_elements,
282 [](u32 i,
283 auto rho,
284 const Tscal *__restrict U,
285 const Tscal *__restrict cs0,
286 Tscal *__restrict P,
287 Tscal *__restrict cs) {
288 using namespace shamrock::sph;
289
290 Tscal cs_out = cs0[i];
291 Tscal rho_a = rho(i);
292
293 Tscal P_a = EOS::pressure_from_cs(cs_out * cs_out, rho_a);
294
295 P[i] = P_a;
296 cs[i] = cs_out;
297 });
298 });
299
300 } else if (
301 SolverEOS_LocallyIsothermalLP07 *eos_config
302 = std::get_if<SolverEOS_LocallyIsothermalLP07>(&solver_config.eos_config.config)) {
303
304 using EOS = shamphys::EOS_LocallyIsothermal<Tscal>;
305
306 storage.merged_patchdata_ghost.get().for_each([&](u64 id, PatchDataLayer &mpdat) {
307 auto &mfield = storage.merged_xyzh.get().get(id);
308
309 sham::DeviceBuffer<Tvec> &buf_xyz = mfield.template get_field_buf_ref<Tvec>(0);
310
311 sham::DeviceBuffer<Tscal> &buf_P
312 = shambase::get_check_ref(storage.pressure).get_field(id).get_buf();
313 sham::DeviceBuffer<Tscal> &buf_cs
314 = shambase::get_check_ref(storage.soundspeed).get_field(id).get_buf();
315 sham::DeviceBuffer<Tscal> &buf_uint = mpdat.get_field_buf_ref<Tscal>(iuint_interf);
316 auto rho_getter = rho_getter_gen(mpdat);
317
318 Tscal cs0 = eos_config->cs0;
319 Tscal r0sq = eos_config->r0 * eos_config->r0;
320 Tscal mq = -eos_config->q;
321
322 u32 total_elements
323 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes.get(id);
324
325 sham::kernel_call(
326 q,
327 sham::MultiRef{rho_getter, buf_uint, buf_xyz},
328 sham::MultiRef{buf_P, buf_cs},
329 total_elements,
330 [cs0, r0sq, mq](
331 u32 i,
332 auto rho,
333 const Tscal *__restrict U,
334 const Tvec *__restrict xyz,
335 Tscal *__restrict P,
336 Tscal *__restrict cs) {
337 using namespace shamrock::sph;
338
339 Tvec R = xyz[i];
340 Tscal rho_a = rho(i);
341
342 Tscal Rsq = sycl::dot(R, R);
343 Tscal cs_sq = EOS::soundspeed_sq(cs0 * cs0, Rsq / r0sq, mq);
344 Tscal cs_out = sycl::sqrt(cs_sq);
345
346 Tscal P_a = EOS::pressure_from_cs(cs_sq, rho_a);
347
348 P[i] = P_a;
349 cs[i] = cs_out;
350 });
351 });
352
353 } else if (
354 SolverEOS_LocallyIsothermalFA2014 *eos_config
355 = std::get_if<SolverEOS_LocallyIsothermalFA2014>(&solver_config.eos_config.config)) {
356
357 Tscal _G = solver_config.get_constant_G();
358
359 using EOS = shamphys::EOS_LocallyIsothermal<Tscal>;
360
361 auto &sink_parts = storage.sinks.get();
362 std::vector<Tvec> sink_pos;
363 std::vector<Tscal> sink_mass;
364 u32 sink_cnt = 0;
365
366 for (auto &s : sink_parts) {
367 sink_pos.push_back(s.pos);
368 sink_mass.push_back(s.mass);
369 sink_cnt++;
370 }
371
372 sycl::buffer<Tvec> sink_pos_buf{sink_pos};
373 sycl::buffer<Tscal> sink_mass_buf{sink_mass};
374
375 storage.merged_patchdata_ghost.get().for_each([&](u64 id, PatchDataLayer &mpdat) {
376 auto &mfield = storage.merged_xyzh.get().get(id);
377
378 sham::DeviceBuffer<Tvec> &buf_xyz = mfield.template get_field_buf_ref<Tvec>(0);
379
380 sham::DeviceBuffer<Tscal> &buf_P
381 = shambase::get_check_ref(storage.pressure).get_field(id).get_buf();
382 sham::DeviceBuffer<Tscal> &buf_cs
383 = shambase::get_check_ref(storage.soundspeed).get_field(id).get_buf();
384 sham::DeviceBuffer<Tscal> &buf_uint = mpdat.get_field_buf_ref<Tscal>(iuint_interf);
385 auto rho_getter = rho_getter_gen(mpdat);
386
387 // TODO: Use the complex kernel call when implemented
388
389 sham::EventList depends_list;
390
391 auto P = buf_P.get_write_access(depends_list);
392 auto cs = buf_cs.get_write_access(depends_list);
393 auto rho = rho_getter.get_read_access(depends_list);
394 auto U = buf_uint.get_read_access(depends_list);
395 auto xyz = buf_xyz.get_read_access(depends_list);
396
397 u32 total_elements
398 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes.get(id);
399
400 auto e = q.submit(depends_list, [&](sycl::handler &cgh) {
401 sycl::accessor spos{sink_pos_buf, cgh, sycl::read_only};
402 sycl::accessor smass{sink_mass_buf, cgh, sycl::read_only};
403 u32 scount = sink_cnt;
404
405 Tscal h_over_r = eos_config->h_over_r;
406 Tscal G = _G;
407
408 cgh.parallel_for(sycl::range<1>{total_elements}, [=](sycl::item<1> item) {
409 using namespace shamrock::sph;
410
411 Tvec R = xyz[item];
412 Tscal rho_a = rho(item.get_linear_id());
413
414 Tscal mpotential = 0;
415 for (u32 i = 0; i < scount; i++) {
416 Tvec s_r = spos[i] - R;
417 Tscal s_m = smass[i];
418 Tscal s_r_abs = sycl::length(s_r);
419 mpotential += G * s_m / s_r_abs;
420 }
421
422 Tscal cs_out = h_over_r * sycl::sqrt(mpotential);
423 Tscal P_a = EOS::pressure_from_cs(cs_out * cs_out, rho_a);
424
425 P[item] = P_a;
426 cs[item] = cs_out;
427 });
428 });
429
430 buf_P.complete_event_state(e);
431 buf_cs.complete_event_state(e);
432 rho_getter.complete_event_state(e);
433 buf_uint.complete_event_state(e);
434 buf_xyz.complete_event_state(e);
435 });
436
437 } else if (
438 SolverEOS_LocallyIsothermalFA2014Extended *eos_config
439 = std::get_if<SolverEOS_LocallyIsothermalFA2014Extended>(
440 &solver_config.eos_config.config)) {
441
442 Tscal _cs0 = eos_config->cs0;
443 Tscal _r0 = eos_config->r0;
444 Tscal _q = eos_config->q;
445 u32 n_sinks = eos_config->n_sinks;
446
447 using EOS = shamphys::EOS_LocallyIsothermal<Tscal>;
448
449 auto &sink_parts = storage.sinks.get();
450 std::vector<Tvec> sink_pos;
451 std::vector<Tscal> sink_mass;
452 u32 sink_cnt = 0;
453
454 for (auto &s : sink_parts) {
455 sink_pos.push_back(s.pos);
456 sink_mass.push_back(s.mass);
457 sink_cnt++;
458 if (sink_pos.size() >= n_sinks) { // We only consider the first n_sinks sinks
459 break;
460 }
461 }
462
463 if (sink_cnt == 0) {
464 throw shambase::make_except_with_loc<std::runtime_error>(
465 "No sinks found for the equation of state");
466 }
467
468 sycl::buffer<Tvec> sink_pos_buf{sink_pos};
469 sycl::buffer<Tscal> sink_mass_buf{sink_mass};
470
471 storage.merged_patchdata_ghost.get().for_each([&](u64 id, PatchDataLayer &mpdat) {
472 auto &mfield = storage.merged_xyzh.get().get(id);
473
474 sham::DeviceBuffer<Tvec> &buf_xyz = mfield.template get_field_buf_ref<Tvec>(0);
475
476 sham::DeviceBuffer<Tscal> &buf_P
477 = shambase::get_check_ref(storage.pressure).get_field(id).get_buf();
478 sham::DeviceBuffer<Tscal> &buf_cs
479 = shambase::get_check_ref(storage.soundspeed).get_field(id).get_buf();
480 sham::DeviceBuffer<Tscal> &buf_uint = mpdat.get_field_buf_ref<Tscal>(iuint_interf);
481 auto rho_getter = rho_getter_gen(mpdat);
482
483 // TODO: Use the complex kernel call when implemented
484
485 sham::EventList depends_list;
486
487 auto P = buf_P.get_write_access(depends_list);
488 auto cs = buf_cs.get_write_access(depends_list);
489 auto rho = rho_getter.get_read_access(depends_list);
490 auto U = buf_uint.get_read_access(depends_list);
491 auto xyz = buf_xyz.get_read_access(depends_list);
492
493 u32 total_elements
494 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes.get(id);
495
496 auto e = q.submit(depends_list, [&](sycl::handler &cgh) {
497 sycl::accessor spos{sink_pos_buf, cgh, sycl::read_only};
498 sycl::accessor smass{sink_mass_buf, cgh, sycl::read_only};
499 u32 scount = sink_cnt;
500
501 Tscal cs0 = _cs0;
502 Tscal r0 = _r0;
503 Tscal q = _q;
504
505 Tscal inv_r0_q = 1. / sycl::pow(r0, q);
506
507 cgh.parallel_for(sycl::range<1>{total_elements}, [=](sycl::item<1> item) {
508 using namespace shamrock::sph;
509
510 Tvec R = xyz[item];
511 Tscal rho_a = rho(item.get_linear_id());
512
513 Tscal sink_mass_sum = 0;
514 Tscal pot_sum = 0;
515 for (u32 i = 0; i < scount; i++) {
516 Tvec s_r = spos[i] - R;
517 Tscal s_m = smass[i];
518 Tscal s_r_abs = sycl::length(s_r);
519 sink_mass_sum += s_m;
520 pot_sum += s_m / s_r_abs;
521 }
522
523 Tscal cs_out = cs0 * inv_r0_q * sycl::pow(pot_sum / sink_mass_sum, q);
524 Tscal P_a = EOS::pressure_from_cs(cs_out * cs_out, rho_a);
525
526 P[item] = P_a;
527 cs[item] = cs_out;
528 });
529 });
530
531 buf_P.complete_event_state(e);
532 buf_cs.complete_event_state(e);
533 rho_getter.complete_event_state(e);
534 buf_uint.complete_event_state(e);
535 buf_xyz.complete_event_state(e);
536 });
537
538 } else if (
539 SolverEOS_Fermi *eos_config
540 = std::get_if<SolverEOS_Fermi>(&solver_config.eos_config.config)) {
541
542 using EOS = shamphys::EOS_Fermi<Tscal>;
543
544 storage.merged_patchdata_ghost.get().for_each([&](u64 id, PatchDataLayer &mpdat) {
545 sham::DeviceBuffer<Tscal> &buf_P
546 = shambase::get_check_ref(storage.pressure).get_field(id).get_buf();
547 sham::DeviceBuffer<Tscal> &buf_cs
548 = shambase::get_check_ref(storage.soundspeed).get_field(id).get_buf();
549 auto rho_getter = rho_getter_gen(mpdat);
550
551 u32 total_elements
552 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes.get(id);
553
554 using namespace shamunits;
555 auto unit_sys = *solver_config.unit_sys;
556
557 Tscal mass = unit_sys.template to<units::kilogram>();
558 Tscal length = unit_sys.template to<units::metre>();
559 Tscal time = unit_sys.template to<units::second>();
560
561 Tscal pressure_unit = mass / length / (time * time);
562 Tscal density_unit = mass / (length * length * length);
563 Tscal velocity_unit = length / time;
564
565 sham::kernel_call(
566 q,
567 sham::MultiRef{rho_getter},
568 sham::MultiRef{buf_P, buf_cs},
569 total_elements,
570 [mu_e = eos_config->mu_e, density_unit, pressure_unit, velocity_unit](
571 u32 i, auto rho, Tscal *__restrict P, Tscal *__restrict cs) {
572 Tscal rho_a = rho(i);
573 auto const res = EOS::pressure_and_soundspeed(mu_e, rho_a * density_unit);
574 P[i] = res.pressure / pressure_unit;
575 cs[i] = res.soundspeed / velocity_unit;
576 });
577 });
578
579 } else {
580 shambase::throw_unimplemented();
581 }
582}
583
584template<class Tvec, template<class> class SPHKernel>
585void shammodels::sph::modules::ComputeEos<Tvec, SPHKernel>::compute_eos() {
586
587 NamedStackEntry stack_loc{"compute eos"};
588
589 Tscal gpart_mass = solver_config.gpart_mass;
590
591 using namespace shamrock;
592 using namespace shamrock::patch;
593
594 shamrock::patch::PatchDataLayerLayout &ghost_layout
595 = shambase::get_check_ref(storage.ghost_layout.get());
596 u32 ihpart_interf = ghost_layout.get_field_idx<Tscal>("hpart");
597
598 shamrock::SchedulerUtility utility(scheduler());
599
600 shambase::DistributedData<u32> &counts_with_ghosts
601 = shambase::get_check_ref(storage.part_counts_with_ghost).indexes;
602
603 shambase::get_check_ref(storage.pressure).ensure_sizes(counts_with_ghosts);
604 shambase::get_check_ref(storage.soundspeed).ensure_sizes(counts_with_ghosts);
605
606 if (solver_config.dust_config.has_epsilon_field()) {
607
608 u32 iepsilon_interf = ghost_layout.get_field_idx<Tscal>("epsilon");
609 u32 nvar_dust = solver_config.dust_config.get_dust_nvar();
610
611 compute_eos_internal([&](PatchDataLayer &mpdat) {
612 return RhoGetterMonofluid<Tscal>{
613 mpdat.get_field_buf_ref<Tscal>(ihpart_interf),
614 mpdat.get_field_buf_ref<Tscal>(iepsilon_interf),
615 nvar_dust,
616 gpart_mass,
617 Kernel::hfactd};
618 });
619 } else if (solver_config.dust_config.has_s_j_field()) {
620
621 u32 is_j_interf = ghost_layout.get_field_idx<Tscal>("s_j");
622 u32 nvar_dust = solver_config.dust_config.get_dust_nvar();
623
624 compute_eos_internal([&](PatchDataLayer &mpdat) {
625 return RhoGetterSJ<Tscal>{
626 mpdat.get_field_buf_ref<Tscal>(ihpart_interf),
627 mpdat.get_field_buf_ref<Tscal>(is_j_interf),
628 nvar_dust,
629 gpart_mass,
630 Kernel::hfactd};
631 });
632 } else {
633 compute_eos_internal([&](PatchDataLayer &mpdat) {
634 return RhoGetterBase<Tscal>{
635 mpdat.get_field_buf_ref<Tscal>(ihpart_interf), gpart_mass, Kernel::hfactd};
636 });
637 }
638}
639
640using namespace shammath;
641template class shammodels::sph::modules::ComputeEos<f64_3, M4>;
642template class shammodels::sph::modules::ComputeEos<f64_3, M6>;
643template class shammodels::sph::modules::ComputeEos<f64_3, M8>;
644
645template class shammodels::sph::modules::ComputeEos<f64_3, C2>;
646template class shammodels::sph::modules::ComputeEos<f64_3, C4>;
647template class shammodels::sph::modules::ComputeEos<f64_3, C6>;
shammodels::gsph::names::common::xyz
constexpr const char * xyz
Position field (3D coordinates).
Definition FieldNames.hpp:33
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
sham::DeviceBuffer
A buffer allocated in USM (Unified Shared Memory).
Definition DeviceBuffer.hpp:106
sham::DeviceBuffer::complete_event_state
void complete_event_state(sycl::event e) const
Complete the event state of the buffer.
Definition DeviceBuffer.hpp:368
sham::DeviceBuffer::get_write_access
T * get_write_access(sham::EventList &depends_list, SourceLocation src_loc=SourceLocation{})
Get a read-write pointer to the buffer's data.
Definition DeviceBuffer.hpp:349
sham::DeviceBuffer::get_read_access
const T * get_read_access(sham::EventList &depends_list, SourceLocation src_loc=SourceLocation{}) const
Get a read-only pointer to the buffer's data.
Definition DeviceBuffer.hpp:331
sham::DeviceQueue
A SYCL queue associated with a device and a context.
Definition DeviceQueue.hpp:32
sham::DeviceQueue::submit
sycl::event submit(Fct &&fct)
Submits a kernel to the SYCL queue.
Definition DeviceQueue.hpp:101
sham::EventList
Class to manage a list of SYCL events.
Definition EventList.hpp:31
shambase::DistributedData
Represents a collection of objects distributed across patches identified by a u64 id.
Definition DistributedData.hpp:43
shammodels::sph::modules::ComputeEos
Module for computing equation of state quantities.
Definition ComputeEos.hpp:34
shammodels::sph::modules::ComputeEos::compute_eos
void compute_eos()
Computes pressure and sound speed from equation of state.
Definition ComputeEos.cpp:585
shamrock::patch::PatchDataLayerLayout::get_field_idx
u32 get_field_idx(const std::string &field_name) const
Get the field id if matching name & type.
Definition PatchDataLayerLayout.hpp:349
shamrock::patch::PatchDataLayer
PatchDataLayer container class, the layout is described in patchdata_layout.
Definition PatchDataLayer.hpp:35
exception.hpp
This header file contains utility functions related to exception handling in the code.
kernel_call.hpp
sham::kernel_call
void kernel_call(sham::DeviceQueue &q, RefIn in, RefOut in_out, u32 n, Functor &&func, SourceLocation &&callsite=SourceLocation{})
Submit a kernel to a SYCL queue.
Definition kernel_call.hpp:514
shambase::get_check_ref
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...
Definition memory.hpp:110
shambase::make_except_with_loc
ExcptTypes make_except_with_loc(std::string message, SourceLocation loc=SourceLocation{})
Create an exception with a message and a location.
Definition exception.hpp:105
shambase::throw_unimplemented
void throw_unimplemented(SourceLocation loc=SourceLocation{})
Throw a std::runtime_error saying that the function is unimplemented.
Definition exception.hpp:153
shammath
namespace for math utility
Definition AABB.hpp:26
shamrock
namespace for the main framework
Definition __init__.py:1
shamunits
namespace containing the units library
Definition units_json.hpp:24
sphkernels.hpp
sph kernels
NamedStackEntry
shambase::details::NamedBasicStackEntry NamedStackEntry
Alias for shambase::details::NamedBasicStackEntry.
Definition stacktrace.hpp:249
sham::MultiRef
A class that references multiple buffers or similar objects.
Definition kernel_call.hpp:234
shamphys::EOS_Adiabatic
Adiabatic equation of state.
Definition eos.hpp:45
shamphys::EOS_Fermi
Fermi Gas EoS.
Definition eos.hpp:196
shamphys::EOS_Isothermal
Isothermal equation of state.
Definition eos.hpp:32
shamphys::EOS_LocallyIsothermal
Locally isothermal equation of state with radial dependence.
Definition eos.hpp:87
shamphys::EOS_Polytropic
Polytropic equation of state.
Definition eos.hpp:66
Generated on for Shamrock by doxygen 1.15.0