27 using Tscal = shambase::VecComponent<Tvec>;
28 using Kernel = SPHKernel<Tscal>;
29 static constexpr Tscal hfactd = Kernel::hfactd;
30 static constexpr Tscal Rkern = Kernel::Rkern;
31 static constexpr Tscal Rker2 = Rkern * Rkern;
37 inline void operator()(
39 const Tvec *__restrict xyz,
40 const Tscal *__restrict hpart,
41 const Tvec *__restrict vxyz,
42 const Tscal *__restrict uint,
43 const Tscal *__restrict omega,
44 const Tscal *__restrict pressure,
45 const Tscal *__restrict cs,
46 const Tscal *__restrict alpha_AV,
48 Tvec *__restrict axyz,
49 Tscal *__restrict duint)
const {
51 using namespace shamrock::sph;
55 Tvec xyz_a = xyz[id_a];
56 Tscal h_a = hpart[id_a];
57 Tvec vxyz_a = vxyz[id_a];
58 Tscal u_a = uint[id_a];
59 Tscal omega_a = omega[id_a];
60 Tscal P_a = pressure[id_a];
61 Tscal cs_a = cs[id_a];
62 Tscal alpha_a = alpha_AV[id_a];
64 Tscal rho_a = rho_h(pmass, h_a, hfactd);
65 Tscal rho_a_sq = rho_a * rho_a;
66 Tscal rho_a_inv = 1. / rho_a;
68 Tscal omega_a_rho_a_inv = 1 / (omega_a * rho_a);
70 Tvec force_pressure = Tvec{0, 0, 0};
71 Tscal tmpdU_pressure = Tscal{0};
73 particle_looper.for_each_object(id_a, [&](
u32 id_b) {
74 Tvec dr = xyz_a - xyz[id_b];
75 Tscal rab2 = sycl::dot(dr, dr);
76 Tscal h_b = hpart[id_b];
78 if (rab2 > h_a * h_a * Rker2 && rab2 > h_b * h_b * Rker2) {
82 Tvec vxyz_b = vxyz[id_b];
83 const Tscal u_b = uint[id_b];
84 Tscal P_b = pressure[id_b];
85 Tscal omega_b = omega[id_b];
86 const Tscal alpha_b = alpha_AV[id_b];
87 Tscal cs_b = cs[id_b];
89 Tscal rab = sycl::sqrt(rab2);
91 Tscal rho_b = rho_h(pmass, h_b, hfactd);
93 Tscal Fab_a = Kernel::dW_3d(rab, h_a);
94 Tscal Fab_b = Kernel::dW_3d(rab, h_b);
96 Tvec v_ab = vxyz_a - vxyz_b;
100 Tscal v_ab_r_ab = sycl::dot(v_ab, r_ab_unit);
101 Tscal abs_v_ab_r_ab = sycl::fabs(v_ab_r_ab);
103 Tscal vsig_a = alpha_a * cs_a + beta_AV * abs_v_ab_r_ab;
104 Tscal vsig_b = alpha_b * cs_b + beta_AV * abs_v_ab_r_ab;
106 Tscal vsig_u = shamrock::sph::vsig_u(P_a, P_b, rho_a, rho_b);
108 Tscal qa_ab = shamrock::sph::q_av(rho_a, vsig_a, v_ab_r_ab);
109 Tscal qb_ab = shamrock::sph::q_av(rho_b, vsig_b, v_ab_r_ab);
111 add_to_derivs_sph_artif_visco_cond(
136 axyz[id_a] = force_pressure;
137 duint[id_a] = tmpdU_pressure;
147 auto edges = get_edges();
149 auto &part_counts_with_ghost = edges.part_counts_with_ghost.indexes;
150 auto &part_counts = edges.part_counts.indexes;
153 edges.xyz.check_sizes(part_counts_with_ghost);
154 edges.hpart.check_sizes(part_counts_with_ghost);
155 edges.vxyz.check_sizes(part_counts_with_ghost);
156 edges.uint.check_sizes(part_counts_with_ghost);
157 edges.omega.check_sizes(part_counts_with_ghost);
158 edges.pressure.check_sizes(part_counts_with_ghost);
159 edges.cs.check_sizes(part_counts_with_ghost);
160 edges.alpha_AV.check_sizes(part_counts_with_ghost);
163 edges.axyz.ensure_sizes(part_counts);
164 edges.duint.ensure_sizes(part_counts);
166 const Tscal pmass = edges.gpart_mass.value;
167 const Tscal alpha_u = edges.alpha_u.value;
168 const Tscal beta_AV = edges.beta_AV.value;
174 shamsys::instance::get_compute_scheduler_ptr(),
176 edges.xyz.get_spans(),
177 edges.hpart.get_spans(),
178 edges.vxyz.get_spans(),
179 edges.uint.get_spans(),
180 edges.omega.get_spans(),
181 edges.pressure.get_spans(),
182 edges.cs.get_spans(),
183 edges.alpha_AV.get_spans(),
187 ComputeKernel{pmass, alpha_u, beta_AV});
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.