39template<
class Tvec,
template<
class>
class SPHKernel>
43 Cfg_Riemann cfg_riemann = solver_config.riemann_config;
45 if (Iterative *v = std::get_if<Iterative>(&cfg_riemann.config)) {
46 update_derivs_iterative(*v);
47 }
else if (HLLC *v = std::get_if<HLLC>(&cfg_riemann.config)) {
48 update_derivs_hllc(*v);
54template<
class Tvec,
template<
class>
class SPHKernel>
61 using namespace shamrock::patch;
72 const bool has_uint = solver_config.has_field_uint();
73 const u32 iuint = has_uint ? pdl.
get_field_idx<Tscal>(gsph::names::newtonian::uint) : 0;
74 const u32 iduint = has_uint ? pdl.
get_field_idx<Tscal>(gsph::names::newtonian::duint) : 0;
79 u32 ihpart_interf = ghost_layout.
get_field_idx<Tscal>(gsph::names::common::hpart);
80 u32 ivxyz_interf = ghost_layout.
get_field_idx<Tvec>(gsph::names::newtonian::vxyz);
81 u32 iomega_interf = ghost_layout.
get_field_idx<Tscal>(gsph::names::newtonian::omega);
82 u32 idensity_interf = ghost_layout.
get_field_idx<Tscal>(gsph::names::newtonian::density);
84 = has_uint ? ghost_layout.
get_field_idx<Tscal>(gsph::names::newtonian::uint) : 0;
87 auto &merged_xyzh = storage.merged_xyzh.get();
102 = merged_xyzh.get(cur_p.
id_patch).template get_field_buf_ref<Tvec>(0);
125 auto vxyz = buf_vxyz.get_read_access(depends_list);
126 auto hpart = buf_hpart.get_read_access(depends_list);
127 auto omega_acc = buf_omega.get_read_access(depends_list);
129 auto pressure_acc = buf_pressure.get_read_access(depends_list);
131 auto ploop_ptrs = pcache.get_read_access(depends_list);
135 Tscal *duint_acc =
nullptr;
137 buf_duint_ptr = &pdat.get_field_buf_ref<Tscal>(iduint);
141 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
142 const Tscal pmass = solver_config.gpart_mass;
143 const Tscal gamma = solver_config.get_eos_gamma();
144 const Tscal tol = cfg.tol;
145 const u32 max_iter = cfg.max_iter;
146 const bool do_energy = has_uint;
151 constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
153 shambase::parallel_for(cgh, pdat.get_obj_cnt(),
"GSPH derivs iterative", [=](
u64 gid) {
154 u32 id_a = (u32) gid;
156 using namespace shamrock::sph;
159 Tvec sum_axyz = {0, 0, 0};
163 const Tscal h_a = hpart[id_a];
164 const Tvec xyz_a = xyz[id_a];
165 const Tvec vxyz_a = vxyz[id_a];
166 const Tscal omega_a = omega_acc[id_a];
169 const Tscal rho_a = sycl::max(density_acc[id_a], Tscal(1e-30));
172 const Tscal P_a = sycl::max(pressure_acc[id_a], Tscal(1e-30));
173 const Tscal cs_a = sycl::max(cs_acc[id_a], Tscal(1e-10));
176 particle_looper.for_each_object(id_a, [&](
u32 id_b) {
181 const Tvec dr = xyz_a - xyz[id_b];
182 const Tscal rab2 = sycl::dot(dr, dr);
183 const Tscal h_b = hpart[id_b];
186 if (rab2 > h_a * h_a * Rker2 && rab2 > h_b * h_b * Rker2) {
190 const Tscal rab = sycl::sqrt(rab2);
191 const Tvec vxyz_b = vxyz[id_b];
192 const Tscal omega_b = omega_acc[id_b];
195 const Tscal rho_b = sycl::max(density_acc[id_b], Tscal(1e-30));
198 const Tscal P_b = sycl::max(pressure_acc[id_b], Tscal(1e-30));
199 const Tscal cs_b = sycl::max(cs_acc[id_b], Tscal(1e-10));
203 const Tvec r_ab_unit = dr * rab_inv;
206 const Tscal u_a_proj = sycl::dot(vxyz_a, r_ab_unit);
207 const Tscal u_b_proj = sycl::dot(vxyz_b, r_ab_unit);
211 auto riemann_result = riemann::iterative_solver<Tscal>(
221 const Tscal p_star = riemann_result.p_star;
222 const Tscal v_star = riemann_result.v_star;
225 const Tscal Fab_a = Kernel::dW_3d(rab, h_a);
226 const Tscal Fab_b = Kernel::dW_3d(rab, h_b);
229 shammodels::gsph::add_gsph_force_contribution<Tvec, Tscal>(
246 axyz[id_a] = sum_axyz;
247 if (duint_acc !=
nullptr) {
248 duint_acc[id_a] = sum_du_a;
256 buf_vxyz.complete_event_state(e);
257 buf_hpart.complete_event_state(e);
258 buf_omega.complete_event_state(e);
260 buf_pressure.complete_event_state(e);
263 if (has_uint && buf_duint_ptr) {
269 pcache.complete_event_state(resulting_events);
273template<
class Tvec,
template<
class>
class SPHKernel>
279 using namespace shamrock::patch;
289 const bool has_uint = solver_config.has_field_uint();
290 const u32 iuint = has_uint ? pdl.
get_field_idx<Tscal>(gsph::names::newtonian::uint) : 0;
291 const u32 iduint = has_uint ? pdl.
get_field_idx<Tscal>(gsph::names::newtonian::duint) : 0;
296 u32 ihpart_interf = ghost_layout.
get_field_idx<Tscal>(gsph::names::common::hpart);
297 u32 ivxyz_interf = ghost_layout.
get_field_idx<Tvec>(gsph::names::newtonian::vxyz);
298 u32 iomega_interf = ghost_layout.
get_field_idx<Tscal>(gsph::names::newtonian::omega);
299 u32 idensity_interf = ghost_layout.
get_field_idx<Tscal>(gsph::names::newtonian::density);
301 = has_uint ? ghost_layout.
get_field_idx<Tscal>(gsph::names::newtonian::uint) : 0;
303 auto &merged_xyzh = storage.merged_xyzh.get();
316 = merged_xyzh.get(cur_p.
id_patch).template get_field_buf_ref<Tvec>(0);
336 auto vxyz = buf_vxyz.get_read_access(depends_list);
337 auto hpart = buf_hpart.get_read_access(depends_list);
338 auto omega_acc = buf_omega.get_read_access(depends_list);
340 auto pressure_acc = buf_pressure.get_read_access(depends_list);
342 auto ploop_ptrs = pcache.get_read_access(depends_list);
345 Tscal *duint_acc =
nullptr;
347 buf_duint_ptr = &pdat.get_field_buf_ref<Tscal>(iduint);
351 auto e = q.
submit(depends_list, [&](sycl::handler &cgh) {
352 const Tscal pmass = solver_config.gpart_mass;
353 const Tscal gamma = solver_config.get_eos_gamma();
354 const bool do_energy = has_uint;
358 constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
360 shambase::parallel_for(cgh, pdat.get_obj_cnt(),
"GSPH derivs HLLC", [=](
u64 gid) {
361 u32 id_a = (u32) gid;
363 using namespace shamrock::sph;
365 Tvec sum_axyz = {0, 0, 0};
368 const Tscal h_a =
hpart[id_a];
369 const Tvec xyz_a =
xyz[id_a];
370 const Tvec vxyz_a =
vxyz[id_a];
371 const Tscal omega_a = omega_acc[id_a];
374 const Tscal rho_a = sycl::max(density_acc[id_a], Tscal(1e-30));
377 const Tscal P_a = sycl::max(pressure_acc[id_a], Tscal(1e-30));
378 const Tscal cs_a = sycl::max(cs_acc[id_a], Tscal(1e-10));
380 particle_looper.for_each_object(id_a, [&](
u32 id_b) {
384 const Tvec dr = xyz_a -
xyz[id_b];
385 const Tscal rab2 = sycl::dot(dr, dr);
386 const Tscal h_b =
hpart[id_b];
388 if (rab2 > h_a * h_a * Rker2 && rab2 > h_b * h_b * Rker2) {
392 const Tscal rab = sycl::sqrt(rab2);
393 const Tvec vxyz_b =
vxyz[id_b];
394 const Tscal omega_b = omega_acc[id_b];
397 const Tscal rho_b = sycl::max(density_acc[id_b], Tscal(1e-30));
400 const Tscal P_b = sycl::max(pressure_acc[id_b], Tscal(1e-30));
401 const Tscal cs_b = sycl::max(cs_acc[id_b], Tscal(1e-10));
404 const Tvec r_ab_unit = dr * rab_inv;
407 const Tscal u_a_proj = sycl::dot(vxyz_a, r_ab_unit);
408 const Tscal u_b_proj = sycl::dot(vxyz_b, r_ab_unit);
412 auto riemann_result = riemann::hllc_solver<Tscal>(
420 const Tscal p_star = riemann_result.p_star;
421 const Tscal v_star = riemann_result.v_star;
423 const Tscal Fab_a = Kernel::dW_3d(rab, h_a);
424 const Tscal Fab_b = Kernel::dW_3d(rab, h_b);
427 shammodels::gsph::add_gsph_force_contribution<Tvec, Tscal>(
443 axyz[id_a] = sum_axyz;
444 if (duint_acc !=
nullptr) {
445 duint_acc[id_a] = sum_du_a;
452 buf_vxyz.complete_event_state(e);
453 buf_hpart.complete_event_state(e);
454 buf_omega.complete_event_state(e);
456 buf_pressure.complete_event_state(e);
459 if (has_uint && buf_duint_ptr) {
465 pcache.complete_event_state(resulting_events);
Constants for field names in GSPH solver, organized by physics mode.
constexpr const char * axyz
3-acceleration field
constexpr const char * vxyz
3-velocity field
constexpr const char * xyz
Position field (3D coordinates)
constexpr const char * hpart
Smoothing length field.
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.
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::event submit(Fct &&fct)
Submits a kernel to the SYCL queue.
DeviceQueue & get_queue(u32 id=0)
Get a reference to a DeviceQueue.
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.
T & get(u64 id)
Returns a reference to an object in the collection.
GSPH derivative update module.
void update_derivs()
Update all derivatives using GSPH Riemann solver approach.
u32 get_field_idx(const std::string &field_name) const
Get the field id if matching name & type.
PatchDataLayer container class, the layout is described in patchdata_layout.
PatchDataField< T > & get_field(u64 id) const
Get the underlying PatchDataField at the given id.
GSPH force computation using Riemann solver results.
GSPH derivative update module.
Iterative Riemann solver for GSPH (van Leer 1997)
T inv_sat_positive(T v, T minvsat=T{1e-9}, T satval=T{0.}) noexcept
inverse saturated (positive numbers only)
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...
void throw_unimplemented(SourceLocation loc=SourceLocation{})
Throw a std::runtime_error saying that the function is unimplemented.
namespace for math utility
namespace for the main framework
Patch object that contain generic patch information.
u64 id_patch
unique key that identify the patch