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() != nrays) {
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();
105 rays_buf, pos.get_buf(), buf_hpart, buf_field, obj_it, hmax_tree.buf_field},
110 const Tvec *__restrict xyz,
111 const Tscal *__restrict hpart,
112 const T *__restrict torender,
113 auto particle_looper,
114 const Tscal *__restrict hmax,
115 T *__restrict render_field) {
116 T acc = sham::VectorProperties<T>::get_zero();
120 constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
122 particle_looper.rtree_for(
124 Tscal rint_cell = hmax[node_id] * Kernel::Rkern;
129 Tvec dr = ray.origin - xyz[id_b];
131 dr -= ray.direction * sycl::dot(dr, ray.direction);
133 Tscal rab2 = sycl::dot(dr, dr);
134 Tscal h_b = hpart[id_b];
136 if (rab2 > h_b * h_b * Rker2) {
140 Tscal rab = sycl::sqrt(rab2);
142 T val = torender[id_b];
144 Tscal rho_b = shamrock::sph::rho_h(partmass, h_b, Kernel::hfactd);
146 acc += partmass * val * Kernel::Y_3d(rab, h_b, 4) / rho_b;
149 render_field[gid] += acc;
153 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.