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() != nring_rays) {
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 ring_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();
119 Tvec ez = ring_ray.get_ez();
121 constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
123 particle_looper.rtree_for(
125 Tscal rint_cell = hmax[node_id] * Kernel::Rkern;
130 Tvec r_center = ring_ray.center - xyz[id_b];
132 Tscal z_val = sycl::dot(r_center, ez);
133 Tscal x_val = sycl::dot(r_center, ring_ray.e_x);
134 Tscal y_val = sycl::dot(r_center, ring_ray.e_y);
135 Tscal r_val = sycl::sqrt(x_val * x_val + y_val * y_val);
137 Tscal delta_r = r_val - ring_ray.radius;
139 Tscal rab2_ring = z_val * z_val + delta_r * delta_r;
140 Tscal h_b = hpart[id_b];
142 if (rab2_ring > h_b * h_b * Rker2) {
146 Tscal rab = sycl::sqrt(rab2_ring);
148 T val = torender[id_b];
150 Tscal rho_b = shamrock::sph::rho_h(partmass, h_b, Kernel::hfactd);
153 acc += partmass * val * Kernel::Y_3d(rab, h_b, 4) / rho_b;
156 render_field[gid] += acc;
160 shamalgs::collective::reduce_buffer_in_place_sum(output_buf, MPI_COMM_WORLD);