Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
SPHInterpolation.cpp
Go to the documentation of this file.
1// -------------------------------------------------------//
2//
3// SHAMROCK code for hydrodynamics
4// Copyright (c) 2021-2026 Timothée David--Cléris <tim.shamrock@proton.me>
5// SPDX-License-Identifier: CeCILL Free Software License Agreement v2.1
6// Shamrock is licensed under the CeCILL 2.1 License, see LICENSE for more information
7//
8// -------------------------------------------------------//
9
16
20#include "shammath/AABB.hpp"
27#include <cmath>
28#include <limits>
29
30template<class Tvec, class T, template<class> class SPHKernel>
32
34
35 auto edges = get_edges();
36
37 auto &part_counts = edges.part_counts.indexes;
38
39 edges.positions.check_sizes(part_counts);
40 edges.h_part.check_sizes(part_counts);
41 edges.field_data.check_sizes(part_counts);
42
43 const sham::DeviceBuffer<Tvec> &interp_points_buf = edges.interp_points.value;
44 sham::DeviceBuffer<T> &output_buf = edges.interpolated_field.value;
45
46 u32 npoints = interp_points_buf.get_size();
47 if (output_buf.get_size() != npoints) {
48 output_buf.resize_discard_data(npoints);
49 }
50 output_buf.fill(sham::VectorProperties<T>::get_zero());
51
52 using u_morton = u32;
54
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();
59
60 part_counts.for_each([&](u64 id, u32 count) {
61 if (count == 0) {
62 return;
63 }
64
65 PatchDataField<Tvec> &pos = edges.positions.get_field(id);
66 if (pos.is_empty()) {
67 return;
68 }
69
70 Tvec bmax = pos.compute_max();
71 Tvec bmin = pos.compute_min();
72
73 shammath::AABB<Tvec> aabb(bmin, bmax);
74
75 Tscal infty = std::numeric_limits<Tscal>::infinity();
76
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);
83
84 u32 obj_cnt = pos.get_obj_cnt();
85
86 Tree tree = Tree::make_empty(dev_sched);
87 tree.rebuild_from_positions(pos.get_buf(), obj_cnt, aabb, tree_reduction_level);
88
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();
93
94 auto hmax_tree = shamtree::compute_tree_field_max_field<Tscal>(
95 tree.structure,
96 tree.reduced_morton_set.get_leaf_cell_iterator(),
97 shamtree::new_empty_karras_radix_tree_field<Tscal>(),
98 buf_hpart);
99
100 auto obj_it = tree.get_object_iterator();
101
103 queue,
105 interp_points_buf,
106 pos.get_buf(),
107 buf_hpart,
108 buf_field,
109 obj_it,
110 hmax_tree.buf_field},
111 sham::MultiRef{output_buf},
112 npoints,
113 [=](u32 gid,
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];
122
123 T acc = sham::VectorProperties<T>::get_zero();
124
125 constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
126
127 particle_looper.rtree_for(
128 [&](u32 node_id, shammath::AABB<Tvec> node_aabb) -> bool {
129 Tscal rint_cell = hmax[node_id] * Kernel::Rkern;
130
131 return node_aabb.expand_all(rint_cell).contains_asymmetric(pos_render);
132 },
133 [&](u32 id_b) {
134 Tvec dr = pos_render - xyz[id_b];
135 Tscal rab2 = sycl::dot(dr, dr);
136 Tscal h_b = hpart[id_b];
137
138 if (rab2 > h_b * h_b * Rker2) {
139 return;
140 }
141
142 Tscal rab = sycl::sqrt(rab2);
143
144 T val = torender[id_b];
145
146 Tscal rho_b = shamrock::sph::rho_h(partmass, h_b, Kernel::hfactd);
147
148 acc += partmass * val * Kernel::W_3d(rab, h_b) / rho_b;
149 });
150
151 render_field[gid] += acc;
152 });
153 });
154
155 shamalgs::collective::reduce_buffer_in_place_sum(output_buf, MPI_COMM_WORLD);
156}
157
158template<class Tvec, class T, template<class> class SPHKernel>
162
163using namespace shammath;
SPH slice interpolation solver graph node.
std::uint32_t u32
32 bit unsigned integer
std::uint64_t u64
64 bit unsigned integer
A buffer allocated in USM (Unified Shared Memory).
void fill(T value, std::array< size_t, 2 > idx_range)
Fill a subpart of the buffer with a given value.
size_t get_size() const
Gets the number of elements in the buffer.
void resize_discard_data(size_t new_size)
same as resize but data will not be copied if reallocation is needed
A SYCL queue associated with a device and a context.
void _impl_evaluate_internal() override
evaluate the node
std::string _impl_get_tex() const override
get the tex of the node
A Compressed Leaf Bounding Volume Hierarchy (CLBVH) for neighborhood queries.
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.
namespace for math utility
Definition AABB.hpp:26
sph kernels
This file contains the definition for the stacktrace related functionality.
#define __shamrock_stack_entry()
Macro to create a stack entry.
A class that references multiple buffers or similar objects.
Definition MultiRef.hpp:33
Axis-Aligned bounding box.
Definition AABB.hpp:99
T lower
Lower bound of the AABB.
Definition AABB.hpp:104
AABB expand_all(Tscal value)
Expand the AABB by a given value on all dimensions.
Definition AABB.hpp:197
bool contains_asymmetric(T point) const noexcept
Check if point is in AABB using half-open interval [lower, upper).
Definition AABB.hpp:256
T upper
Upper bound of the AABB.
Definition AABB.hpp:105