Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
TreeReducedMortonCodes.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
19
20namespace shamrock::tree {
21
22 template<class u_morton>
23 void TreeReducedMortonCodes<u_morton>::build(
24 sycl::queue &queue,
25 u32 obj_cnt,
26 u32 reduc_level,
27 TreeMortonCodes<u_morton> &morton_codes,
28
29 bool &one_cell_mode) {
30
31 // return a sycl buffer from reduc index map instead
32 shamlog_debug_sycl_ln(
33 "RadixTree", "reduction algorithm"); // TODO put reduction level in class member
34
35 // TODO document that the layout of reduc_index_map is in the end {0 .. ,i .. ,N ,0}
36 // with the trailling 0 to invert the range for the walk in one cell mode
37
39 queue,
40 obj_cnt,
41 morton_codes.buf_morton,
42 reduc_level,
43 buf_reduc_index_map,
44 tree_leaf_count);
45
46 shamlog_debug_sycl_ln(
47 "RadixTree",
48 "reduction results : (before :",
49 obj_cnt,
50 " | after :",
51 tree_leaf_count,
52 ") ratio :",
53 shambase::format_printf("%2.2f", f32(obj_cnt) / f32(tree_leaf_count)));
54
55 if (tree_leaf_count > 1) {
56
57 shamlog_debug_sycl_ln("RadixTree", "sycl_morton_remap_reduction");
58 buf_tree_morton = std::make_unique<sycl::buffer<u_morton>>(tree_leaf_count);
59
61 queue,
62 tree_leaf_count,
63 buf_reduc_index_map,
64 morton_codes.buf_morton,
65 buf_tree_morton);
66
67 one_cell_mode = false;
68
69 } else if (tree_leaf_count == 1) {
70
71 tree_leaf_count = 2;
72 one_cell_mode = true;
73
74 buf_tree_morton = std::make_unique<sycl::buffer<u_morton>>(
75 shamalgs::memory::vector_to_buf(
76 shamsys::instance::get_compute_queue(), std::vector<u_morton>{0, 0})
77 // tree morton = {0,0} is a flag for the one cell mode
78 );
79
80 } else {
81 throw shambase::make_except_with_loc<std::runtime_error>("0 leaf tree cannot exists");
82 }
83 }
84
85 template<class u_morton>
86 bool TreeReducedMortonCodes<u_morton>::operator==(
87 const TreeReducedMortonCodes<u_morton> &rhs) const {
88 bool cmp = true;
89
90 cmp = cmp && (tree_leaf_count == rhs.tree_leaf_count);
91
92 using namespace shamalgs::primitives;
93
94 cmp = cmp && (buf_reduc_index_map->size() == rhs.buf_reduc_index_map->size());
95
96 cmp = cmp
97 && equals(
98 shamsys::instance::get_compute_queue(),
99 *buf_reduc_index_map,
100 *rhs.buf_reduc_index_map,
101 buf_reduc_index_map->size());
102 cmp = cmp
103 && equals(
104 shamsys::instance::get_compute_queue(),
105 *buf_tree_morton,
106 *rhs.buf_tree_morton,
107 tree_leaf_count);
108
109 return cmp;
110 }
111
112 template class TreeReducedMortonCodes<u32>;
113 template class TreeReducedMortonCodes<u64>;
114
115} // namespace shamrock::tree
float f32
Alias for float.
std::uint32_t u32
32 bit unsigned integer
Element-wise equality comparison algorithms for buffers.
namespace for primitive algorithm (e.g. sort, scan, reductions, ...)
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
void throw_with_loc(std::string message, SourceLocation loc=SourceLocation{})
Throw an exception and append the source location to it.
void sycl_morton_remap_reduction(sycl::queue &queue, u32 morton_leaf_count, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map, std::unique_ptr< sycl::buffer< u_morton > > &buf_morton, std::unique_ptr< sycl::buffer< u_morton > > &buf_leaf_morton)
Remaps a Morton tree on device using a reduction index map.
void reduction_alg(sycl::queue &queue, u32 morton_count, std::unique_ptr< sycl::buffer< u_morton > > &buf_morton, u32 reduction_level, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map, u32 &morton_leaf_count)
Reduces a Morton tree on device.