39 auto edges = get_edges();
41 edges.field_axyz_ext.ensure_sizes(edges.sizes.indexes);
43 if (edges.sizes.indexes.get_ids().size() != 1) {
45 "Self gravity MM mode only supports one patch so far, current number "
47 + std::to_string(edges.sizes.indexes.get_ids().size()));
50 Tscal G = edges.constant_G.data;
51 Tscal gpart_mass = edges.gpart_mass.data;
53 Tscal gravitational_softening = epsilon * epsilon;
56 static constexpr u32 mass_moment_terms = MassMoments::num_component;
58 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
63 [&](
u64 id,
const u64 &n) {
67 Tvec bmax = xyz.compute_max();
68 Tvec bmin = xyz.compute_min();
72 auto bvh = RTree::make_empty(dev_sched);
73 bvh.rebuild_from_positions(
74 xyz.get_buf(), xyz.get_obj_cnt(), aabb, reduction_level);
77 auto mass_moments_tree = compute_tree_mass_moments<Tvec, Umorton, mm_order - 1>(
78 bvh, xyz.get_buf(), gpart_mass);
81 auto obj_it = bvh.get_object_iterator();
84 sham::MultiRef{xyz.get_buf(), obj_it, mass_moments_tree.buf_field},
87 [theta_crit = theta_crit, gravitational_softening, gpart_mass, G](
91 const Tscal *mass_moments_scal,
97 auto &tree_traverser = particle_looper.tree_traverser;
98 auto &cell_iterator = particle_looper.cell_iterator;
102 [&](
u32 node_id) ->
bool {
106 tree_traverser.aabb_min[node_id],
107 tree_traverser.aabb_max[node_id]};
111 Tvec delta_B = node_aabb.
delt();
113 Tscal sz_B = sham::max_component(delta_B) / 2;
115 Tvec r = s_B - xyz_i;
116 Tscal r_sq = r.x() * r.x() + r.y() * r.y() + r.z() * r.z();
118 Tscal theta = (r_sq == 0) ? theta_crit * 2
119 : sz_B * sycl::rsqrt(r_sq);
121 return theta > theta_crit;
125 = node_id - tree_traverser.tree_traverser.offset_leaf;
126 cell_iterator.for_each_in_leaf_cell(leaf_id, [&](
u32 j) {
127 Tvec R = xyz[j] - xyz_i;
128 const Tscal r_inv = sycl::rsqrt(
129 R.x() * R.x() + R.y() * R.y() + R.z() * R.z()
130 + gravitational_softening);
131 f_i += gpart_mass * r_inv * r_inv * r_inv * R;
138 Tvec r_fmm = s_B - xyz_i;
140 MassMoments Q_n_B = MassMoments::load(
142 + node_id * MassMoments::num_component,
149 contract_grav_moment_to_force<Tscal, mm_order>(
153 axyz_ext[i] += f_i * G;
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.