35 auto edges = get_edges();
37 auto &part_counts = edges.part_counts.indexes;
39 edges.positions.check_sizes(part_counts);
40 edges.h_part.check_sizes(part_counts);
41 edges.field_data.check_sizes(part_counts);
47 if (output_buf.
get_size() != npoints) {
50 output_buf.
fill(sham::VectorProperties<T>::get_zero());
55 Tscal partmass = edges.gpart_mass.data;
56 u32 tree_reduction_level = edges.tree_reduction_level.data;
57 sham::DeviceQueue &queue = shamsys::instance::get_compute_scheduler().get_queue();
58 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
60 part_counts.for_each([&](
u64 id,
u32 count) {
70 Tvec bmax = pos.compute_max();
71 Tvec bmin = pos.compute_min();
75 Tscal infty = std::numeric_limits<Tscal>::infinity();
77 aabb.
lower[0] = std::nextafter(aabb.
lower[0], -infty);
78 aabb.
lower[1] = std::nextafter(aabb.
lower[1], -infty);
79 aabb.
lower[2] = std::nextafter(aabb.
lower[2], -infty);
80 aabb.
upper[0] = std::nextafter(aabb.
upper[0], infty);
81 aabb.
upper[1] = std::nextafter(aabb.
upper[1], infty);
82 aabb.
upper[2] = std::nextafter(aabb.
upper[2], infty);
84 u32 obj_cnt = pos.get_obj_cnt();
86 Tree tree = Tree::make_empty(dev_sched);
87 tree.rebuild_from_positions(pos.get_buf(), obj_cnt, aabb, tree_reduction_level);
89 auto &hpart_span = edges.h_part.get_spans().get(
id);
90 auto &field_span = edges.field_data.get_spans().get(
id);
91 auto &buf_hpart = hpart_span.field_ref.get_buf();
92 auto &buf_field = field_span.field_ref.get_buf();
94 auto hmax_tree = shamtree::compute_tree_field_max_field<Tscal>(
96 tree.reduced_morton_set.get_leaf_cell_iterator(),
97 shamtree::new_empty_karras_radix_tree_field<Tscal>(),
100 auto obj_it = tree.get_object_iterator();
110 hmax_tree.buf_field},
114 const Tvec *__restrict pixel_positions,
115 const Tvec *__restrict xyz,
116 const Tscal *__restrict hpart,
117 const T *__restrict torender,
118 auto particle_looper,
119 const Tscal *__restrict hmax,
120 T *__restrict render_field) {
121 Tvec pos_render = pixel_positions[gid];
123 T acc = sham::VectorProperties<T>::get_zero();
125 constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
127 particle_looper.rtree_for(
129 Tscal rint_cell = hmax[node_id] * Kernel::Rkern;
134 Tvec dr = pos_render - xyz[id_b];
135 Tscal rab2 = sycl::dot(dr, dr);
136 Tscal h_b = hpart[id_b];
138 if (rab2 > h_b * h_b * Rker2) {
142 Tscal rab = sycl::sqrt(rab2);
144 T val = torender[id_b];
146 Tscal rho_b = shamrock::sph::rho_h(partmass, h_b, Kernel::hfactd);
148 acc += partmass * val * Kernel::W_3d(rab, h_b) / rho_b;
151 render_field[gid] += acc;
155 shamalgs::collective::reduce_buffer_in_place_sum(output_buf, MPI_COMM_WORLD);
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.