25template<
class Tvec,
template<
class>
class SPHKernel>
29 using Cfg_AV =
typename Config::AVConfig;
31 using None =
typename Cfg_AV::None;
32 using Constant =
typename Cfg_AV::Constant;
33 using VaryingMM97 =
typename Cfg_AV::VaryingMM97;
34 using VaryingCD10 =
typename Cfg_AV::VaryingCD10;
35 using ConstantDisc =
typename Cfg_AV::ConstantDisc;
36 if (None *v = std::get_if<None>(&solver_config.artif_viscosity.config)) {
37 shamlog_debug_ln(
"UpdateViscosity",
"skipping artif viscosity update (No viscosity mode)");
38 }
else if (Constant *v = std::get_if<Constant>(&solver_config.artif_viscosity.config)) {
39 shamlog_debug_ln(
"UpdateViscosity",
"skipping artif viscosity update (Constant mode)");
40 }
else if (VaryingMM97 *v = std::get_if<VaryingMM97>(&solver_config.artif_viscosity.config)) {
41 update_artificial_viscosity_mm97(dt, *v);
42 }
else if (VaryingCD10 *v = std::get_if<VaryingCD10>(&solver_config.artif_viscosity.config)) {
43 update_artificial_viscosity_cd10(dt, *v);
44 }
else if (ConstantDisc *v = std::get_if<ConstantDisc>(&solver_config.artif_viscosity.config)) {
45 shamlog_debug_ln(
"UpdateViscosity",
"skipping artif viscosity update (constant AV)");
51template<
class Tvec,
template<
class>
class SPHKernel>
53 Tscal dt,
typename Config::AVConfig::VaryingMM97 cfg) {
55 shamlog_debug_ln(
"UpdateViscosity",
"Updating alpha viscosity (Morris & Monaghan 1997)");
57 using namespace shamrock::patch;
77 u32 obj_cnt = pdat.get_obj_cnt();
84 auto cs = buf_cs.get_read_access(depends_list);
85 auto h = buf_h.get_read_access(depends_list);
86 auto alpha_AV = buf_alpha_AV.get_read_access(depends_list);
87 auto alpha_AV_updated = buf_alpha_AV_updated.get_write_access(depends_list);
89 auto e = queue.
submit(depends_list, [&, dt](sycl::handler &cgh) {
90 Tscal sigma_decay = cfg.sigma_decay;
91 Tscal alpha_min = cfg.alpha_min;
92 Tscal alpha_max = cfg.alpha_max;
94 cgh.parallel_for(sycl::range<1>{obj_cnt}, [=](sycl::item<1> item) {
95 Tscal cs_a = cs[item];
97 Tscal alpha_a = alpha_AV[item];
98 Tscal divv_a = divv[item];
101 Tscal inv_tau_a = vsig * sigma_decay / h_a;
102 Tscal fact_t = dt * inv_tau_a;
103 Tscal euler_impl_fact = 1 / (1 + fact_t);
105 Tscal source = sham::max<Tscal>(0., -divv_a);
107 Tscal new_alpha = (alpha_a + source * dt + fact_t * alpha_min) * euler_impl_fact;
109 alpha_AV_updated[item] = sham::min(alpha_max, new_alpha);
114 buf_cs.complete_event_state(e);
115 buf_h.complete_event_state(e);
116 buf_alpha_AV.complete_event_state(e);
117 buf_alpha_AV_updated.complete_event_state(e);
121template<
class Tvec,
template<
class>
class SPHKernel>
123 Tscal dt,
typename Config::AVConfig::VaryingCD10 cfg) {
126 shamlog_debug_ln(
"UpdateViscosity",
"Updating alpha viscosity (Cullen & Dehnen 2010)");
128 using namespace shamrock::patch;
151 u32 obj_cnt = pdat.get_obj_cnt();
158 auto curlv = buf_curlv.get_read_access(depends_list);
159 auto dtdivv = buf_dtdivv.get_read_access(depends_list);
160 auto cs = buf_cs.get_read_access(depends_list);
161 auto h = buf_h.get_read_access(depends_list);
162 auto alpha_AV = buf_alpha_AV.get_read_access(depends_list);
163 auto alpha_AV_updated = buf_alpha_AV_updated.get_write_access(depends_list);
165 auto e = queue.
submit(depends_list, [&, dt](sycl::handler &cgh) {
166 Tscal sigma_decay = cfg.sigma_decay;
167 Tscal alpha_min = cfg.alpha_min;
168 Tscal alpha_max = cfg.alpha_max;
170 cgh.parallel_for(sycl::range<1>{obj_cnt}, [=](sycl::item<1> item) {
171 Tscal cs_a = cs[item];
173 Tscal alpha_a = alpha_AV[item];
174 Tscal divv_a = divv[item];
175 Tvec curlv_a = curlv[item];
176 Tscal dtdivv_a = dtdivv[item];
179 Tscal inv_tau_a = vsig * sigma_decay / h_a;
180 Tscal fact_t = dt * inv_tau_a;
181 Tscal euler_impl_fact = 1 / (1 + fact_t);
190 auto xi_lim = [](Tscal divv, Tvec curlv) {
191 auto fac = sham::max(-divv, Tscal{0});
193 auto traceS = sycl::dot(curlv, curlv);
195 return fac / (fac + traceS);
199 Tscal balsara_corec = xi_lim(divv_a, curlv_a);
201 Tscal A_a = balsara_corec * sham::max<Tscal>(-dtdivv_a, 0);
203 Tscal temp = cs_a * cs_a;
206 = sham::min((cs_a > 0) ? 10 * h_a * h_a * A_a / (temp) : alpha_min, alpha_max);
211 Tscal new_alpha = (alpha_a + alpha_loc_a * fact_t) * euler_impl_fact;
213 if (alpha_loc_a > alpha_a) {
214 new_alpha = alpha_loc_a;
217 alpha_AV_updated[item] = new_alpha;
222 buf_curlv.complete_event_state(e);
223 buf_dtdivv.complete_event_state(e);
224 buf_cs.complete_event_state(e);
225 buf_h.complete_event_state(e);
226 buf_alpha_AV.complete_event_state(e);
227 buf_alpha_AV_updated.complete_event_state(e);
std::uint32_t u32
32 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.
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.
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.
void throw_with_loc(std::string message, SourceLocation loc=SourceLocation{})
Throw an exception and append the source location to it.
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
file containing formulas for sph forces
Patch object that contain generic patch information.
u64 id_patch
unique key that identify the patch