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;
38 inline void operator()(
40 const Tvec *__restrict xyz,
41 const Tscal *__restrict hpart,
42 const Tvec *__restrict vxyz,
43 const Tscal *__restrict uint,
44 const Tscal *__restrict omega,
45 const Tscal *__restrict pressure,
46 const Tscal *__restrict cs,
47 const Tscal *__restrict alpha_AV,
48 const Tscal *__restrict s_j,
50 Tvec *__restrict axyz,
51 Tscal *__restrict duint)
const {
53 using namespace shamrock::sph;
57 Tvec xyz_a = xyz[id_a];
58 Tscal h_a = hpart[id_a];
59 Tvec vxyz_a = vxyz[id_a];
60 Tscal u_a = uint[id_a];
61 Tscal omega_a = omega[id_a];
62 Tscal P_a = pressure[id_a];
63 Tscal cs_a = cs[id_a];
64 Tscal alpha_a = alpha_AV[id_a];
66 Tscal rho_a = rho_h(pmass, h_a, hfactd);
67 Tscal rho_a_sq = rho_a * rho_a;
68 Tscal rho_a_inv = 1. / rho_a;
70 Tscal epsilon_sum_a = 0;
71 for (
u32 j = 0; j < ndust; j++) {
72 Tscal s = s_j[id_a * ndust + j];
73 epsilon_sum_a += s * s / rho_a;
76 Tscal omega_a_rho_a_inv = 1 / (omega_a * rho_a);
78 Tvec force_pressure = Tvec{0, 0, 0};
79 Tscal tmpdU_pressure = Tscal{0};
81 particle_looper.for_each_object(id_a, [&](
u32 id_b) {
82 Tvec dr = xyz_a - xyz[id_b];
83 Tscal rab2 = sycl::dot(dr, dr);
84 Tscal h_b = hpart[id_b];
86 if (rab2 > h_a * h_a * Rker2 && rab2 > h_b * h_b * Rker2) {
90 Tvec vxyz_b = vxyz[id_b];
91 const Tscal u_b = uint[id_b];
92 Tscal P_b = pressure[id_b];
93 Tscal omega_b = omega[id_b];
94 const Tscal alpha_b = alpha_AV[id_b];
95 Tscal cs_b = cs[id_b];
97 Tscal rab = sycl::sqrt(rab2);
99 Tscal rho_b = rho_h(pmass, h_b, hfactd);
101 Tscal epsilon_sum_b = 0;
102 for (
u32 j = 0; j < ndust; j++) {
103 Tscal s = s_j[id_b * ndust + j];
104 epsilon_sum_b += s * s / rho_b;
107 Tscal Fab_a = Kernel::dW_3d(rab, h_a);
108 Tscal Fab_b = Kernel::dW_3d(rab, h_b);
110 Tvec v_ab = vxyz_a - vxyz_b;
114 Tscal v_ab_r_ab = sycl::dot(v_ab, r_ab_unit);
115 Tscal abs_v_ab_r_ab = sycl::fabs(v_ab_r_ab);
117 Tscal vsig_a = alpha_a * cs_a + beta_AV * abs_v_ab_r_ab;
118 Tscal vsig_b = alpha_b * cs_b + beta_AV * abs_v_ab_r_ab;
120 Tscal vsig_u = shamrock::sph::vsig_u(P_a, P_b, rho_a, rho_b);
125 add_to_derivs_sph_artif_visco_cond(
150 axyz[id_a] = force_pressure;
151 duint[id_a] = tmpdU_pressure;
161 auto edges = get_edges();
163 auto &part_counts_with_ghost = edges.part_counts_with_ghost.indexes;
164 auto &part_counts = edges.part_counts.indexes;
167 edges.xyz.check_sizes(part_counts_with_ghost);
168 edges.hpart.check_sizes(part_counts_with_ghost);
169 edges.vxyz.check_sizes(part_counts_with_ghost);
170 edges.uint.check_sizes(part_counts_with_ghost);
171 edges.omega.check_sizes(part_counts_with_ghost);
172 edges.pressure.check_sizes(part_counts_with_ghost);
173 edges.cs.check_sizes(part_counts_with_ghost);
174 edges.alpha_AV.check_sizes(part_counts_with_ghost);
175 edges.s_j.check_sizes(part_counts_with_ghost);
178 edges.axyz.ensure_sizes(part_counts);
179 edges.duint.ensure_sizes(part_counts);
181 const Tscal pmass = edges.gpart_mass.value;
182 const Tscal alpha_u = edges.alpha_u.value;
183 const Tscal beta_AV = edges.beta_AV.value;
189 shamsys::instance::get_compute_scheduler_ptr(),
191 edges.xyz.get_spans(),
192 edges.hpart.get_spans(),
193 edges.vxyz.get_spans(),
194 edges.uint.get_spans(),
195 edges.omega.get_spans(),
196 edges.pressure.get_spans(),
197 edges.cs.get_spans(),
198 edges.alpha_AV.get_spans(),
199 edges.s_j.get_spans(),
203 ComputeKernel{pmass, alpha_u, beta_AV, ndust});