Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
RadixTreeMortonBuilder.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 "shambase/exception.hpp"
17#include "shamalgs/algorithm.hpp"
18#include "shambackends/fmt_bindings/fmt_defs.hpp"
19#include "shambackends/math.hpp"
20#include "shamcomm/logs.hpp"
21#include "shammath/sfc/morton.hpp"
22#include "shamtree/MortonKernels.hpp"
23#include "shamtree/RadixTreeMortonBuilder.hpp"
24#include "shamtree/key_morton_sort.hpp"
25
26template<class morton_t, class pos_t, u32 dim>
27void RadixTreeMortonBuilder<morton_t, pos_t, dim>::build(
28 sycl::queue &queue,
29 std::tuple<pos_t, pos_t> bounding_box,
30 sycl::buffer<pos_t> &pos_buf,
31 u32 cnt_obj,
32 std::unique_ptr<sycl::buffer<morton_t>> &out_buf_morton,
33 std::unique_ptr<sycl::buffer<u32>> &out_buf_particle_index_map) {
34
35 using namespace logger;
36 using namespace shamrock::sfc;
37
38 if (cnt_obj > i32_max - 1) {
39 throw shambase::make_except_with_loc<std::invalid_argument>(
40 "number of element in patch above i32_max-1");
41 }
42
43 debug_sycl_ln("RadixTree", "box dim :", bounding_box);
44
45 u32 morton_len = sham::roundup_pow2_clz(cnt_obj);
46
47 debug_sycl_ln("RadixTree", "morton buffer length :", morton_len);
48 out_buf_morton = std::make_unique<sycl::buffer<morton_t>>(morton_len);
49
50 MortonKernels<morton_t, pos_t, dim>::sycl_xyz_to_morton(
51 queue,
52 cnt_obj,
53 pos_buf,
54 std::get<0>(bounding_box),
55 std::get<1>(bounding_box),
56 out_buf_morton);
57
58 MortonKernels<morton_t, pos_t, dim>::sycl_fill_trailling_buffer(
59 queue, cnt_obj, morton_len, out_buf_morton);
60
61 out_buf_particle_index_map = std::make_unique<sycl::buffer<u32>>(
62 shamalgs::algorithm::gen_buffer_index(queue, morton_len));
63
64 sycl_sort_morton_key_pair(queue, morton_len, out_buf_particle_index_map, out_buf_morton);
65}
66
67template<class morton_t, class pos_t, u32 dim>
68void RadixTreeMortonBuilder<morton_t, pos_t, dim>::build(
69 sham::DeviceScheduler_ptr dev_sched,
70 std::tuple<pos_t, pos_t> bounding_box,
71 sham::DeviceBuffer<pos_t> &pos_buf,
72 u32 cnt_obj,
73 std::unique_ptr<sycl::buffer<morton_t>> &out_buf_morton,
74 std::unique_ptr<sycl::buffer<u32>> &out_buf_particle_index_map) {
75 sycl::queue &queue = dev_sched->get_queue().q;
76
77 using namespace logger;
78 using namespace shamrock::sfc;
79
80 if (cnt_obj > i32_max - 1) {
81 throw shambase::make_except_with_loc<std::invalid_argument>(
82 "number of element in patch above i32_max-1");
83 }
84
85 debug_sycl_ln("RadixTree", "box dim :", bounding_box);
86
87 u32 morton_len = sham::roundup_pow2_clz(cnt_obj);
88
89 debug_sycl_ln("RadixTree", "morton buffer length :", morton_len);
90 out_buf_morton = std::make_unique<sycl::buffer<morton_t>>(morton_len);
91
92 MortonKernels<morton_t, pos_t, dim>::sycl_xyz_to_morton(
93 dev_sched,
94 cnt_obj,
95 pos_buf,
96 std::get<0>(bounding_box),
97 std::get<1>(bounding_box),
98 out_buf_morton);
99
100 MortonKernels<morton_t, pos_t, dim>::sycl_fill_trailling_buffer(
101 queue, cnt_obj, morton_len, out_buf_morton);
102
103 out_buf_particle_index_map = std::make_unique<sycl::buffer<u32>>(
104 shamalgs::algorithm::gen_buffer_index(queue, morton_len));
105
106 sycl_sort_morton_key_pair(queue, morton_len, out_buf_particle_index_map, out_buf_morton);
107}
108
109template<class morton_t, class pos_t, u32 dim>
110void RadixTreeMortonBuilder<morton_t, pos_t, dim>::build_raw(
111 sycl::queue &queue,
112 std::tuple<pos_t, pos_t> bounding_box,
113 sycl::buffer<pos_t> &pos_buf,
114 u32 cnt_obj,
115 std::unique_ptr<sycl::buffer<morton_t>> &out_buf_morton) {
116
117 using namespace logger;
118 using namespace shamrock::sfc;
119
120 if (cnt_obj > i32_max - 1) {
121 throw shambase::make_except_with_loc<std::invalid_argument>(
122 "number of element in patch above i32_max-1");
123 }
124
125 debug_sycl_ln("RadixTree", "box dim :", bounding_box);
126
127 debug_sycl_ln("RadixTree", "morton buffer length :", cnt_obj);
128 out_buf_morton = std::make_unique<sycl::buffer<morton_t>>(cnt_obj);
129
130 MortonKernels<morton_t, pos_t, dim>::sycl_xyz_to_morton(
131 queue,
132 cnt_obj,
133 pos_buf,
134 std::get<0>(bounding_box),
135 std::get<1>(bounding_box),
136 out_buf_morton);
137}
138
139template class RadixTreeMortonBuilder<u32, f32_3, 3>;
140template class RadixTreeMortonBuilder<u64, f32_3, 3>;
141template class RadixTreeMortonBuilder<u32, f64_3, 3>;
142template class RadixTreeMortonBuilder<u64, f64_3, 3>;
143template class RadixTreeMortonBuilder<u32, u32_3, 3>;
144template class RadixTreeMortonBuilder<u64, u32_3, 3>;
145template class RadixTreeMortonBuilder<u32, u64_3, 3>;
146template class RadixTreeMortonBuilder<u64, u64_3, 3>;
147template class RadixTreeMortonBuilder<u32, i64_3, 3>;
148template class RadixTreeMortonBuilder<u64, i64_3, 3>;
149
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
Helper class to build morton codes.
Definition RadixTreeMortonBuilder.hpp:32
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
exception.hpp
This header file contains utility functions related to exception handling in the code.
fmt_defs.hpp
Define the fmt formatters for sycl::vec.
key_morton_sort.hpp
sycl_sort_morton_key_pair
void sycl_sort_morton_key_pair(sycl::queue &queue, u32 morton_count_rounded_pow, std::unique_ptr< sycl::buffer< u32 > > &buf_index, std::unique_ptr< sycl::buffer< u_morton > > &buf_morton)
sort morton code and generate remap table
morton.hpp
Morton curve implementation.
logger
alias namespace to simplify the use of log functions
Definition logs.hpp:319
sham::roundup_pow2_clz
constexpr T roundup_pow2_clz(T v) noexcept
round up to the next power of two 0 is rounded up to 1 as it is not a pow of 2 every input above the ...
Definition math.hpp:805
shamalgs::algorithm::gen_buffer_index
sycl::buffer< u32 > gen_buffer_index(sycl::queue &q, u32 len)
generate a buffer such that for i in [0,len[, buf[i] = i
Definition algorithm.cpp:25
shambase::throw_with_loc
void throw_with_loc(std::string message, SourceLocation loc=SourceLocation{})
Throw an exception and append the source location to it.
Definition exception.hpp:132
i32_max
constexpr i32 i32_max
i32 max value
Definition numeric_limits.hpp:66
shamcomm::logs::debug_sycl_ln
void debug_sycl_ln(std::string module_name, Types... var2)
Prints a log message with multiple arguments followed by a newline.
Definition logs.hpp:133
Generated on Wed May 13 2026 13:21:49 for Shamrock by doxygen 1.9.8