Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
SlopeLimitedGradient.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
18#include "shammodels/ramses/modules/SlopeLimitedGradient.hpp"
19#include "shammodels/common/amr/NeighGraph.hpp"
20#include "shammodels/ramses/modules/SlopeLimitedGradientUtilities.hpp"
21#include "shamrock/patch/PatchDataField.hpp"
22#include <type_traits>
23
24namespace {
25
26 template<class Tvec, class TgridVec, SlopeMode mode>
27 class KernelSlopeLimScalGrad {
28
29 using Edges = typename shammodels::basegodunov::modules::
30 SlopeLimitedScalarGradient<Tvec, TgridVec>::Edges;
31 using Tscal = shambase::VecComponent<Tvec>;
32 using OrientedAMRGraph = shammodels::basegodunov::modules::OrientedAMRGraph<Tvec, TgridVec>;
33 using AMRGraph = shammodels::basegodunov::modules::AMRGraph;
34
35 public:
36 inline static void kernel(Edges &edges, u32 block_size, u32 var_per_cell) {
37
38 edges.cell_neigh_graph.graph.for_each(
39 [&](u64 id, const OrientedAMRGraph &oriented_cell_graph) {
40 auto &field_span = edges.span_field.get_spans().get(id);
41 auto &field_grad_span = edges.span_grad_field.get_spans().get(id);
42 auto &cell_sizes_span = edges.spans_block_cell_sizes.get_spans().get(id);
43
44 AMRGraph &graph_neigh_xp
45 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::xp]);
46 AMRGraph &graph_neigh_xm
47 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::xm]);
48 AMRGraph &graph_neigh_yp
49 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::yp]);
50 AMRGraph &graph_neigh_ym
51 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::ym]);
52 AMRGraph &graph_neigh_zp
53 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::zp]);
54 AMRGraph &graph_neigh_zm
55 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::zm]);
56
57 sham::EventList depends_list;
58
59 auto cell_sizes = cell_sizes_span.get_read_access(depends_list);
60 auto field = field_span.get_read_access(depends_list);
61 auto field_grad = field_grad_span.get_write_access(depends_list);
62
63 auto graph_iter_xp = graph_neigh_xp.get_read_access(depends_list);
64 auto graph_iter_xm = graph_neigh_xm.get_read_access(depends_list);
65 auto graph_iter_yp = graph_neigh_yp.get_read_access(depends_list);
66 auto graph_iter_ym = graph_neigh_ym.get_read_access(depends_list);
67 auto graph_iter_zp = graph_neigh_zp.get_read_access(depends_list);
68 auto graph_iter_zm = graph_neigh_zm.get_read_access(depends_list);
69
70 sham::DeviceQueue &q = shamsys::instance::get_compute_scheduler().get_queue();
71 auto e = q.submit(depends_list, [&](sycl::handler &cgh) {
72 u32 cell_count = (edges.sizes.indexes.get(id)) * block_size;
73
74 shambase::parallel_for(
75 cgh, cell_count * var_per_cell, "compute_grad_rho", [=](u64 gid) {
76 const u32 tmp_gid = (u32) gid;
77
78 const u32 cell_global_id = tmp_gid / var_per_cell;
79 const u32 var_off_loc = tmp_gid % var_per_cell;
80
81 const u32 block_id = cell_global_id / block_size;
82 const u32 cell_loc_id = cell_global_id % block_size;
83
84 Tscal delta_cell = cell_sizes[block_id];
85
86 auto result = get_3d_grad<Tscal, Tvec, mode>(
87 cell_global_id,
88 delta_cell,
89 graph_iter_xp,
90 graph_iter_xm,
91 graph_iter_yp,
92 graph_iter_ym,
93 graph_iter_zp,
94 graph_iter_zm,
95 [=](u32 id) {
96 return field[var_per_cell * id + var_off_loc];
97 });
98
99 field_grad[var_per_cell * cell_global_id + var_off_loc]
100 = {result[0], result[1], result[2]};
101 });
102 });
103
104 cell_sizes_span.complete_event_state(e);
105 field_span.complete_event_state(e);
106 field_grad_span.complete_event_state(e);
107
108 graph_neigh_xp.complete_event_state(e);
109 graph_neigh_xm.complete_event_state(e);
110 graph_neigh_yp.complete_event_state(e);
111 graph_neigh_ym.complete_event_state(e);
112 graph_neigh_zp.complete_event_state(e);
113 graph_neigh_zm.complete_event_state(e);
114 });
115 }
116 };
117
118 template<class Tvec, class TgridVec, SlopeMode mode>
119 class KernelSlopeLimVecGrad {
120
121 using Edges = typename shammodels::basegodunov::modules::
122 SlopeLimitedVectorGradient<Tvec, TgridVec>::Edges;
123 using Tscal = shambase::VecComponent<Tvec>;
124 using OrientedAMRGraph = shammodels::basegodunov::modules::OrientedAMRGraph<Tvec, TgridVec>;
125 using AMRGraph = shammodels::basegodunov::modules::AMRGraph;
126
127 public:
128 inline static void kernel(Edges &edges, u32 block_size, u32 var_per_cell) {
129
130 edges.cell_neigh_graph.graph.for_each(
131 [&](u64 id, const OrientedAMRGraph &oriented_cell_graph) {
132 auto &field_span = edges.span_field.get_spans().get(id);
133 auto &field_dx_span = edges.span_dx_field.get_spans().get(id);
134 auto &field_dy_span = edges.span_dy_field.get_spans().get(id);
135 auto &field_dz_span = edges.span_dz_field.get_spans().get(id);
136 auto &cell_sizes_span = edges.spans_block_cell_sizes.get_spans().get(id);
137
138 AMRGraph &graph_neigh_xp
139 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::xp]);
140 AMRGraph &graph_neigh_xm
141 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::xm]);
142 AMRGraph &graph_neigh_yp
143 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::yp]);
144 AMRGraph &graph_neigh_ym
145 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::ym]);
146 AMRGraph &graph_neigh_zp
147 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::zp]);
148 AMRGraph &graph_neigh_zm
149 = shambase::get_check_ref(oriented_cell_graph.graph_links[Direction::zm]);
150
151 sham::EventList depends_list;
152
153 auto cell_sizes = cell_sizes_span.get_read_access(depends_list);
154 auto field = field_span.get_read_access(depends_list);
155 auto field_dx = field_dx_span.get_write_access(depends_list);
156 auto field_dy = field_dy_span.get_write_access(depends_list);
157 auto field_dz = field_dz_span.get_write_access(depends_list);
158
159 auto graph_iter_xp = graph_neigh_xp.get_read_access(depends_list);
160 auto graph_iter_xm = graph_neigh_xm.get_read_access(depends_list);
161 auto graph_iter_yp = graph_neigh_yp.get_read_access(depends_list);
162 auto graph_iter_ym = graph_neigh_ym.get_read_access(depends_list);
163 auto graph_iter_zp = graph_neigh_zp.get_read_access(depends_list);
164 auto graph_iter_zm = graph_neigh_zm.get_read_access(depends_list);
165
166 sham::DeviceQueue &q = shamsys::instance::get_compute_scheduler().get_queue();
167 auto e = q.submit(depends_list, [&](sycl::handler &cgh) {
168 u32 cell_count = (edges.sizes.indexes.get(id)) * block_size;
169
170 shambase::parallel_for(
171 cgh, cell_count * var_per_cell, "compute_grad_rho", [=](u64 gid) {
172 const u32 tmp_gid = (u32) gid;
173
174 const u32 cell_global_id = tmp_gid / var_per_cell;
175 const u32 var_off_loc = tmp_gid % var_per_cell;
176
177 const u32 block_id = cell_global_id / block_size;
178 const u32 cell_loc_id = cell_global_id % block_size;
179
180 Tscal delta_cell = cell_sizes[block_id];
181
182 auto result = get_3d_grad<Tvec, Tvec, mode>(
183 cell_global_id,
184 delta_cell,
185 graph_iter_xp,
186 graph_iter_xm,
187 graph_iter_yp,
188 graph_iter_ym,
189 graph_iter_zp,
190 graph_iter_zm,
191 [=](u32 id) {
192 return field[var_per_cell * id + var_off_loc];
193 });
194
195 field_dx[var_per_cell * cell_global_id + var_off_loc] = result[0];
196 field_dy[var_per_cell * cell_global_id + var_off_loc] = result[1];
197 field_dz[var_per_cell * cell_global_id + var_off_loc] = result[2];
198 });
199 });
200
201 cell_sizes_span.complete_event_state(e);
202 field_span.complete_event_state(e);
203 field_dx_span.complete_event_state(e);
204 field_dy_span.complete_event_state(e);
205 field_dz_span.complete_event_state(e);
206
207 graph_neigh_xp.complete_event_state(e);
208 graph_neigh_xm.complete_event_state(e);
209 graph_neigh_yp.complete_event_state(e);
210 graph_neigh_ym.complete_event_state(e);
211 graph_neigh_zp.complete_event_state(e);
212 graph_neigh_zm.complete_event_state(e);
213 });
214 }
215 };
216} // namespace
217
218namespace shammodels::basegodunov::modules {
219
220 template<class Tvec, class TgridVec>
221 void SlopeLimitedScalarGradient<Tvec, TgridVec>::_impl_evaluate_internal() {
222 StackEntry stack_loc{};
223 auto edges = get_edges();
224
225 edges.spans_block_cell_sizes.check_sizes(edges.sizes.indexes);
226 edges.span_field.check_sizes(edges.sizes.indexes);
227
228 edges.span_grad_field.ensure_sizes(edges.sizes.indexes);
229
230 if (mode == SlopeMode::None) {
231 using Kern = KernelSlopeLimScalGrad<Tvec, TgridVec, None>;
232 Kern::kernel(edges, block_size, var_per_cell);
233 } else if (mode == SlopeMode::VanLeer_f) {
234 using Kern = KernelSlopeLimScalGrad<Tvec, TgridVec, VanLeer_f>;
235 Kern::kernel(edges, block_size, var_per_cell);
236 } else if (mode == SlopeMode::VanLeer_std) {
237 using Kern = KernelSlopeLimScalGrad<Tvec, TgridVec, VanLeer_std>;
238 Kern::kernel(edges, block_size, var_per_cell);
239 } else if (mode == SlopeMode::VanLeer_sym) {
240 using Kern = KernelSlopeLimScalGrad<Tvec, TgridVec, VanLeer_sym>;
241 Kern::kernel(edges, block_size, var_per_cell);
242 } else if (mode == SlopeMode::Minmod) {
243 using Kern = KernelSlopeLimScalGrad<Tvec, TgridVec, Minmod>;
244 Kern::kernel(edges, block_size, var_per_cell);
245 } else {
246 shambase::throw_unimplemented();
247 }
248 }
249
250 template<class Tvec, class TgridVec>
251 std::string SlopeLimitedScalarGradient<Tvec, TgridVec>::_impl_get_tex() const {
252
253 std::string sizes = get_ro_edge_base(0).get_tex_symbol();
254 std::string cell_neigh_graph = get_ro_edge_base(1).get_tex_symbol();
255 std::string spans_block_cell_sizes = get_ro_edge_base(2).get_tex_symbol();
256 std::string span_field = get_ro_edge_base(3).get_tex_symbol();
257 std::string span_grad_field = get_rw_edge_base(0).get_tex_symbol();
258
259 std::string tex = R"tex(
260 Slope limited gradient (Scalar)
261 )tex";
262
263 shambase::replace_all(tex, "{sizes}", sizes);
264 shambase::replace_all(tex, "{cell_neigh_graph}", cell_neigh_graph);
265 shambase::replace_all(tex, "{spans_block_cell_sizes}", spans_block_cell_sizes);
266 shambase::replace_all(tex, "{span_field}", span_field);
267 shambase::replace_all(tex, "{span_grad_field}", span_grad_field);
268
269 return tex;
270 }
271 template<class Tvec, class TgridVec>
272 void SlopeLimitedVectorGradient<Tvec, TgridVec>::_impl_evaluate_internal() {
273 StackEntry stack_loc{};
274 auto edges = get_edges();
275
276 edges.spans_block_cell_sizes.check_sizes(edges.sizes.indexes);
277 edges.span_field.check_sizes(edges.sizes.indexes);
278
279 edges.span_dx_field.ensure_sizes(edges.sizes.indexes);
280 edges.span_dy_field.ensure_sizes(edges.sizes.indexes);
281 edges.span_dz_field.ensure_sizes(edges.sizes.indexes);
282
283 if (mode == SlopeMode::None) {
284 using Kern = KernelSlopeLimVecGrad<Tvec, TgridVec, None>;
285 Kern::kernel(edges, block_size, var_per_cell);
286 } else if (mode == SlopeMode::VanLeer_f) {
287 using Kern = KernelSlopeLimVecGrad<Tvec, TgridVec, VanLeer_f>;
288 Kern::kernel(edges, block_size, var_per_cell);
289 } else if (mode == SlopeMode::VanLeer_std) {
290 using Kern = KernelSlopeLimVecGrad<Tvec, TgridVec, VanLeer_std>;
291 Kern::kernel(edges, block_size, var_per_cell);
292 } else if (mode == SlopeMode::VanLeer_sym) {
293 using Kern = KernelSlopeLimVecGrad<Tvec, TgridVec, VanLeer_sym>;
294 Kern::kernel(edges, block_size, var_per_cell);
295 } else if (mode == SlopeMode::Minmod) {
296 using Kern = KernelSlopeLimVecGrad<Tvec, TgridVec, Minmod>;
297 Kern::kernel(edges, block_size, var_per_cell);
298 } else {
299 shambase::throw_unimplemented();
300 }
301 }
302
303 template<class Tvec, class TgridVec>
304 std::string SlopeLimitedVectorGradient<Tvec, TgridVec>::_impl_get_tex() const {
305
306 std::string sizes = get_ro_edge_base(0).get_tex_symbol();
307 std::string cell_neigh_graph = get_ro_edge_base(1).get_tex_symbol();
308 std::string spans_block_cell_sizes = get_ro_edge_base(2).get_tex_symbol();
309 std::string span_field = get_ro_edge_base(3).get_tex_symbol();
310 std::string span_dx_field = get_rw_edge_base(0).get_tex_symbol();
311 std::string span_dy_field = get_rw_edge_base(0).get_tex_symbol();
312 std::string span_dz_field = get_rw_edge_base(0).get_tex_symbol();
313
314 std::string tex = R"tex(
315 Slope limited gradient (Vector)
316 )tex";
317
318 shambase::replace_all(tex, "{sizes}", sizes);
319 shambase::replace_all(tex, "{cell_neigh_graph}", cell_neigh_graph);
320 shambase::replace_all(tex, "{spans_block_cell_sizes}", spans_block_cell_sizes);
321 shambase::replace_all(tex, "{span_field}", span_field);
322 shambase::replace_all(tex, "{span_dx_field}", span_dx_field);
323 shambase::replace_all(tex, "{span_dy_field}", span_dy_field);
324 shambase::replace_all(tex, "{span_dz_field}", span_dz_field);
325
326 return tex;
327 }
328
329} // namespace shammodels::basegodunov::modules
330
331template class shammodels::basegodunov::modules::SlopeLimitedScalarGradient<f64_3, i64_3>;
332template class shammodels::basegodunov::modules::SlopeLimitedVectorGradient<f64_3, i64_3>;
u32
std::uint32_t u32
32 bit unsigned integer
Definition aliases_int.hpp:27
u64
std::uint64_t u64
64 bit unsigned integer
Definition aliases_int.hpp:26
sham::DeviceQueue
A SYCL queue associated with a device and a context.
Definition DeviceQueue.hpp:32
sham::DeviceQueue::submit
sycl::event submit(Fct &&fct)
Submits a kernel to the SYCL queue.
Definition DeviceQueue.hpp:101
sham::DeviceScheduler::get_queue
DeviceQueue & get_queue(u32 id=0)
Get a reference to a DeviceQueue.
Definition DeviceScheduler.cpp:57
sham::EventList
Class to manage a list of SYCL events.
Definition EventList.hpp:31
shammodels::basegodunov::modules::SlopeLimitedScalarGradient::_impl_get_tex
virtual std::string _impl_get_tex() const
get the tex of the node
Definition SlopeLimitedGradient.cpp:251
shammodels::basegodunov::modules::SlopeLimitedVectorGradient::_impl_get_tex
virtual std::string _impl_get_tex() const
get the tex of the node
Definition SlopeLimitedGradient.cpp:304
shambase::replace_all
void replace_all(std::string &inout, std::string_view what, std::string_view with)
replace all occurence of a search string with another
Definition string.hpp:183
shambase::get_check_ref
T & get_check_ref(const std::unique_ptr< T > &ptr, SourceLocation loc=SourceLocation())
Takes a std::unique_ptr and returns a reference to the object it holds. It throws a std::runtime_erro...
Definition memory.hpp:110
shambase::throw_unimplemented
void throw_unimplemented(SourceLocation loc=SourceLocation{})
Throw a std::runtime_error saying that the function is unimplemented.
Definition exception.hpp:153
shammodels::basegodunov::modules
namespace for the basegodunov model modules
Definition AMRCellInfos.hpp:20
Generated on Wed May 13 2026 13:21:49 for Shamrock by doxygen 1.9.8