Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
TreeMortonCodes.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#include "shamtree/TreeMortonCodes.hpp"
17#include "shamalgs/primitives/equals.hpp"
18#include "shamtree/RadixTreeMortonBuilder.hpp"
19
20namespace shamrock::tree {
21
22 template<class u_morton>
23 template<class T>
24 void TreeMortonCodes<u_morton>::build(
25 sycl::queue &queue,
26 shammath::CoordRange<T> coord_range,
27 u32 obj_cnt,
28 sycl::buffer<T> &pos_buf) {
29 StackEntry stack_loc{};
30
31 this->obj_cnt = obj_cnt;
32
33 using TProp = shambase::VectorProperties<T>;
34
35 RadixTreeMortonBuilder<u_morton, T, TProp::dimension>::build(
36 queue,
37 {coord_range.lower, coord_range.upper},
38 pos_buf,
39 obj_cnt,
40 buf_morton,
41 buf_particle_index_map);
42 }
43
44 template<class u_morton>
45 template<class T>
46 void TreeMortonCodes<u_morton>::build(
47 sham::DeviceScheduler_ptr dev_sched,
48 shammath::CoordRange<T> coord_range,
49 u32 obj_cnt,
50 sham::DeviceBuffer<T> &pos_buf) {
51 StackEntry stack_loc{};
52
53 this->obj_cnt = obj_cnt;
54
55 using TProp = shambase::VectorProperties<T>;
56
57 RadixTreeMortonBuilder<u_morton, T, TProp::dimension>::build(
58 dev_sched,
59 {coord_range.lower, coord_range.upper},
60 pos_buf,
61 obj_cnt,
62 buf_morton,
63 buf_particle_index_map);
64 }
65
66 template<class u_morton>
67 template<class T>
68 std::unique_ptr<sycl::buffer<u_morton>> TreeMortonCodes<u_morton>::build_raw(
69 sycl::queue &queue,
70 shammath::CoordRange<T> coord_range,
71 u32 obj_cnt,
72 sycl::buffer<T> &pos_buf) {
73
74 std::unique_ptr<sycl::buffer<u_morton>> buf_morton;
75
76 StackEntry stack_loc{};
77
78 using TProp = shambase::VectorProperties<T>;
79
80 RadixTreeMortonBuilder<u_morton, T, TProp::dimension>::build_raw(
81 queue, {coord_range.lower, coord_range.upper}, pos_buf, obj_cnt, buf_morton);
82
83 return buf_morton;
84 }
85
86 template<class u_morton>
87 bool TreeMortonCodes<u_morton>::operator==(const TreeMortonCodes<u_morton> &rhs) const {
88 bool cmp = true;
89
90 cmp = cmp && (obj_cnt == rhs.obj_cnt);
91
92 using namespace shamalgs::primitives;
93
94 cmp = cmp
95 && equals(
96 shamsys::instance::get_compute_queue(), *buf_morton, *rhs.buf_morton, obj_cnt);
97 cmp = cmp
98 && equals(
99 shamsys::instance::get_compute_queue(),
100 *buf_particle_index_map,
101 *rhs.buf_particle_index_map,
102 obj_cnt);
103
104 return cmp;
105 }
106
107 template class TreeMortonCodes<u32>;
108 template class TreeMortonCodes<u64>;
109
110 // u32, f64_3
111 template void TreeMortonCodes<u32>::build<f64_3>(
112 sycl::queue &queue,
113 shammath::CoordRange<f64_3> coord_range,
114 u32 obj_cnt,
115 sycl::buffer<f64_3> &pos_buf);
116
117 template void TreeMortonCodes<u32>::build<f64_3>(
118 sham::DeviceScheduler_ptr dev_sched,
119 shammath::CoordRange<f64_3> coord_range,
120 u32 obj_cnt,
121 sham::DeviceBuffer<f64_3> &pos_buf);
122
123 template std::unique_ptr<sycl::buffer<u32>> TreeMortonCodes<u32>::build_raw<f64_3>(
124 sycl::queue &queue,
125 shammath::CoordRange<f64_3> coord_range,
126 u32 obj_cnt,
127 sycl::buffer<f64_3> &pos_buf);
128
129 // u32, f32_3
130 template void TreeMortonCodes<u32>::build<f32_3>(
131 sycl::queue &queue,
132 shammath::CoordRange<f32_3> coord_range,
133 u32 obj_cnt,
134 sycl::buffer<f32_3> &pos_buf);
135
136 template void TreeMortonCodes<u32>::build<f32_3>(
137 sham::DeviceScheduler_ptr dev_sched,
138 shammath::CoordRange<f32_3> coord_range,
139 u32 obj_cnt,
140 sham::DeviceBuffer<f32_3> &pos_buf);
141
142 template std::unique_ptr<sycl::buffer<u32>> TreeMortonCodes<u32>::build_raw<f32_3>(
143 sycl::queue &queue,
144 shammath::CoordRange<f32_3> coord_range,
145 u32 obj_cnt,
146 sycl::buffer<f32_3> &pos_buf);
147
148 // u32, u64_3
149 template void TreeMortonCodes<u32>::build<u64_3>(
150 sycl::queue &queue,
151 shammath::CoordRange<u64_3> coord_range,
152 u32 obj_cnt,
153 sycl::buffer<u64_3> &pos_buf);
154
155 template void TreeMortonCodes<u32>::build<u64_3>(
156 sham::DeviceScheduler_ptr dev_sched,
157 shammath::CoordRange<u64_3> coord_range,
158 u32 obj_cnt,
159 sham::DeviceBuffer<u64_3> &pos_buf);
160
161 template std::unique_ptr<sycl::buffer<u32>> TreeMortonCodes<u32>::build_raw<u64_3>(
162 sycl::queue &queue,
163 shammath::CoordRange<u64_3> coord_range,
164 u32 obj_cnt,
165 sycl::buffer<u64_3> &pos_buf);
166
167 // u64, f64_3
168 template void TreeMortonCodes<u64>::build<f64_3>(
169 sycl::queue &queue,
170 shammath::CoordRange<f64_3> coord_range,
171 u32 obj_cnt,
172 sycl::buffer<f64_3> &pos_buf);
173
174 template void TreeMortonCodes<u64>::build<f64_3>(
175 sham::DeviceScheduler_ptr dev_sched,
176 shammath::CoordRange<f64_3> coord_range,
177 u32 obj_cnt,
178 sham::DeviceBuffer<f64_3> &pos_buf);
179
180 template std::unique_ptr<sycl::buffer<u64>> TreeMortonCodes<u64>::build_raw<f64_3>(
181 sycl::queue &queue,
182 shammath::CoordRange<f64_3> coord_range,
183 u32 obj_cnt,
184 sycl::buffer<f64_3> &pos_buf);
185
186 // u64, f32_3
187 template void TreeMortonCodes<u64>::build<f32_3>(
188 sycl::queue &queue,
189 shammath::CoordRange<f32_3> coord_range,
190 u32 obj_cnt,
191 sycl::buffer<f32_3> &pos_buf);
192
193 template void TreeMortonCodes<u64>::build<f32_3>(
194 sham::DeviceScheduler_ptr dev_sched,
195 shammath::CoordRange<f32_3> coord_range,
196 u32 obj_cnt,
197 sham::DeviceBuffer<f32_3> &pos_buf);
198
199 template std::unique_ptr<sycl::buffer<u64>> TreeMortonCodes<u64>::build_raw<f32_3>(
200 sycl::queue &queue,
201 shammath::CoordRange<f32_3> coord_range,
202 u32 obj_cnt,
203 sycl::buffer<f32_3> &pos_buf);
204
205 // u64, u64_3
206 template void TreeMortonCodes<u64>::build<u64_3>(
207 sycl::queue &queue,
208 shammath::CoordRange<u64_3> coord_range,
209 u32 obj_cnt,
210 sycl::buffer<u64_3> &pos_buf);
211
212 template void TreeMortonCodes<u64>::build<u64_3>(
213 sham::DeviceScheduler_ptr dev_sched,
214 shammath::CoordRange<u64_3> coord_range,
215 u32 obj_cnt,
216 sham::DeviceBuffer<u64_3> &pos_buf);
217
218 template std::unique_ptr<sycl::buffer<u64>> TreeMortonCodes<u64>::build_raw<u64_3>(
219 sycl::queue &queue,
220 shammath::CoordRange<u64_3> coord_range,
221 u32 obj_cnt,
222 sycl::buffer<u64_3> &pos_buf);
223
224 // u64, u32_3
225 template void TreeMortonCodes<u64>::build<u32_3>(
226 sycl::queue &queue,
227 shammath::CoordRange<u32_3> coord_range,
228 u32 obj_cnt,
229 sycl::buffer<u32_3> &pos_buf);
230
231 template void TreeMortonCodes<u64>::build<u32_3>(
232 sham::DeviceScheduler_ptr dev_sched,
233 shammath::CoordRange<u32_3> coord_range,
234 u32 obj_cnt,
235 sham::DeviceBuffer<u32_3> &pos_buf);
236
237 template std::unique_ptr<sycl::buffer<u64>> TreeMortonCodes<u64>::build_raw<u32_3>(
238 sycl::queue &queue,
239 shammath::CoordRange<u32_3> coord_range,
240 u32 obj_cnt,
241 sycl::buffer<u32_3> &pos_buf);
242
243 // u64, i64_3
244 template void TreeMortonCodes<u64>::build<i64_3>(
245 sycl::queue &queue,
246 shammath::CoordRange<i64_3> coord_range,
247 u32 obj_cnt,
248 sycl::buffer<i64_3> &pos_buf);
249
250 template void TreeMortonCodes<u64>::build<i64_3>(
251 sham::DeviceScheduler_ptr dev_sched,
252 shammath::CoordRange<i64_3> coord_range,
253 u32 obj_cnt,
254 sham::DeviceBuffer<i64_3> &pos_buf);
255
256 template std::unique_ptr<sycl::buffer<u64>> TreeMortonCodes<u64>::build_raw<i64_3>(
257 sycl::queue &queue,
258 shammath::CoordRange<i64_3> coord_range,
259 u32 obj_cnt,
260 sycl::buffer<i64_3> &pos_buf);
261
262 // u32, i64_3
263 template void TreeMortonCodes<u32>::build<i64_3>(
264 sycl::queue &queue,
265 shammath::CoordRange<i64_3> coord_range,
266 u32 obj_cnt,
267 sycl::buffer<i64_3> &pos_buf);
268
269 template void TreeMortonCodes<u32>::build<i64_3>(
270 sham::DeviceScheduler_ptr dev_sched,
271 shammath::CoordRange<i64_3> coord_range,
272 u32 obj_cnt,
273 sham::DeviceBuffer<i64_3> &pos_buf);
274
275 template std::unique_ptr<sycl::buffer<u32>> TreeMortonCodes<u32>::build_raw<i64_3>(
276 sycl::queue &queue,
277 shammath::CoordRange<i64_3> coord_range,
278 u32 obj_cnt,
279 sycl::buffer<i64_3> &pos_buf);
280
281 // u32, u32_3
282 template void TreeMortonCodes<u32>::build<u32_3>(
283 sycl::queue &queue,
284 shammath::CoordRange<u32_3> coord_range,
285 u32 obj_cnt,
286 sycl::buffer<u32_3> &pos_buf);
287
288 template void TreeMortonCodes<u32>::build<u32_3>(
289 sham::DeviceScheduler_ptr dev_sched,
290 shammath::CoordRange<u32_3> coord_range,
291 u32 obj_cnt,
292 sham::DeviceBuffer<u32_3> &pos_buf);
293
294 template std::unique_ptr<sycl::buffer<u32>> TreeMortonCodes<u32>::build_raw<u32_3>(
295 sycl::queue &queue,
296 shammath::CoordRange<u32_3> coord_range,
297 u32 obj_cnt,
298 sycl::buffer<u32_3> &pos_buf);
299
300} // namespace shamrock::tree
RadixTreeMortonBuilder.hpp
Utility to build morton codes for the radix tree.
u32
std::uint32_t u32
32 bit unsigned integer
Definition aliases_int.hpp:27
RadixTreeMortonBuilder::build_raw
static void build_raw(sycl::queue &queue, std::tuple< pos_t, pos_t > bounding_box, sycl::buffer< pos_t > &pos_buf, u32 cnt_obj, std::unique_ptr< sycl::buffer< morton_t > > &out_buf_morton)
build a raw mrton table from a position buffer (no sorting & index map)
Definition RadixTreeMortonBuilder.cpp:110
RadixTreeMortonBuilder::build
static void build(sycl::queue &queue, std::tuple< pos_t, pos_t > bounding_box, sycl::buffer< pos_t > &pos_buf, u32 cnt_obj, std::unique_ptr< sycl::buffer< morton_t > > &out_buf_morton, std::unique_ptr< sycl::buffer< u32 > > &out_buf_particle_index_map)
build morton code table for the tree
Definition RadixTreeMortonBuilder.cpp:27
sham::DeviceBuffer
A buffer allocated in USM (Unified Shared Memory)
Definition DeviceBuffer.hpp:106
equals.hpp
Element-wise equality comparison algorithms for buffers.
shamalgs::primitives
namespace for primitive algorithm (e.g. sort, scan, reductions, ...)
Definition ImplControl.hpp:23
shamalgs::primitives::equals
bool equals(sycl::queue &q, sycl::buffer< T > &buf1, sycl::buffer< T > &buf2, u32 cnt)
Compare elements between two sycl::buffers for equality.
Definition equals.hpp:77
Generated on Wed May 13 2026 13:21:49 for Shamrock by doxygen 1.9.8