Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
TreeCellRanges.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/TreeCellRanges.hpp"
17#include "shamalgs/primitives/equals.hpp"
18#include "shambackends/fmt_bindings/fmt_defs.hpp"
19#include "shamtree/MortonKernels.hpp"
20#include "shamtree/kernels/compute_ranges.hpp"
21
22namespace shamrock::tree {
23
24 template<class u_morton, class pos_t>
25 void TreeCellRanges<u_morton, pos_t>::build1(
26 sycl::queue &queue,
27 TreeReducedMortonCodes<u_morton> &tree_reduced_morton_codes,
28 TreeStructure<u_morton> &tree_struct) {
29 if (!tree_struct.one_cell_mode) {
30
31 shamlog_debug_sycl_ln("RadixTree", "compute_cellvolume");
32
33 buf_pos_min_cell = std::make_unique<sycl::buffer<ipos_t>>(
34 tree_struct.internal_cell_count + tree_reduced_morton_codes.tree_leaf_count);
35 buf_pos_max_cell = std::make_unique<sycl::buffer<ipos_t>>(
36 tree_struct.internal_cell_count + tree_reduced_morton_codes.tree_leaf_count);
37
38 sycl_compute_cell_ranges(
39 queue,
40 tree_reduced_morton_codes.tree_leaf_count,
41 tree_struct.internal_cell_count,
42 tree_reduced_morton_codes.buf_tree_morton,
43 tree_struct.buf_lchild_id,
44 tree_struct.buf_rchild_id,
45 tree_struct.buf_lchild_flag,
46 tree_struct.buf_rchild_flag,
47 tree_struct.buf_endrange,
48 buf_pos_min_cell,
49 buf_pos_max_cell);
50
51 } else {
52 // throw shamrock_exc("one cell mode is not implemented");
53 // TODO do some extensive test on one cell mode
54
55 buf_pos_min_cell = std::make_unique<sycl::buffer<ipos_t>>(
56 tree_struct.internal_cell_count + tree_reduced_morton_codes.tree_leaf_count);
57 buf_pos_max_cell = std::make_unique<sycl::buffer<ipos_t>>(
58 tree_struct.internal_cell_count + tree_reduced_morton_codes.tree_leaf_count);
59
60 {
61
62 sycl::host_accessor pos_min_cell{
63 *buf_pos_min_cell, sycl::write_only, sycl::no_init};
64 sycl::host_accessor pos_max_cell{
65 *buf_pos_max_cell, sycl::write_only, sycl::no_init};
66
67 pos_min_cell[0] = {0, 0, 0};
68 pos_max_cell[0] = {Morton::max_val, Morton::max_val, Morton::max_val};
69
70 pos_min_cell[1] = {0, 0, 0};
71 pos_max_cell[1] = {Morton::max_val, Morton::max_val, Morton::max_val};
72
73 pos_min_cell[2] = {0, 0, 0};
74 pos_max_cell[2] = {0, 0, 0};
75
76 shamlog_debug_sycl_ln("RadixTree", "compute_cellvolume one cell mode");
77 shamlog_debug_sycl_ln(
78 "RadixTree",
79 " -> ",
80 pos_min_cell[0],
81 pos_max_cell[0],
82 pos_min_cell[1],
83 pos_max_cell[1],
84 pos_min_cell[2],
85 pos_max_cell[2],
86 "len =",
87 tree_struct.internal_cell_count + tree_reduced_morton_codes.tree_leaf_count);
88 }
89 }
90 }
91
92 template<class u_morton, class pos_t>
93 void TreeCellRanges<u_morton, pos_t>::build2(
94 sycl::queue &queue, u32 total_count, std::tuple<pos_t, pos_t> bounding_box) {
95
96 buf_pos_min_cell_flt = std::make_unique<sycl::buffer<pos_t>>(total_count);
97 buf_pos_max_cell_flt = std::make_unique<sycl::buffer<pos_t>>(total_count);
98
99 shamlog_debug_sycl_ln("RadixTree", "sycl_convert_cell_range");
100
101 shamrock::sfc::MortonKernels<u_morton, pos_t, dim>::sycl_irange_to_range(
102 queue,
103 total_count,
104 std::get<0>(bounding_box),
105 std::get<1>(bounding_box),
106 buf_pos_min_cell,
107 buf_pos_max_cell,
108 buf_pos_min_cell_flt,
109 buf_pos_max_cell_flt);
110
111 // remove old buf ?
112 }
113
114 template<class u_morton, class pos_t>
115 TreeCellRanges<u_morton, pos_t>::TreeCellRanges(const TreeCellRanges<u_morton, pos_t> &other)
116 : buf_pos_min_cell(
117 shamalgs::memory::duplicate(
118 shamsys::instance::get_compute_queue(),
119 other.buf_pos_min_cell)), // size = total count
120 buf_pos_max_cell(
121 shamalgs::memory::duplicate(
122 shamsys::instance::get_compute_queue(),
123 other.buf_pos_max_cell)), // size = total count
124 buf_pos_min_cell_flt(
125 shamalgs::memory::duplicate(
126 shamsys::instance::get_compute_queue(),
127 other.buf_pos_min_cell_flt)), // size = total count
128 buf_pos_max_cell_flt(
129 shamalgs::memory::duplicate(
130 shamsys::instance::get_compute_queue(),
131 other.buf_pos_max_cell_flt)) // size = total count
132 {}
133
134 template<class u_morton, class pos_t>
135 bool TreeCellRanges<u_morton, pos_t>::operator==(
136 const TreeCellRanges<u_morton, pos_t> &rhs) const {
137 bool cmp = true;
138
139 using namespace shamalgs::primitives;
140
141 cmp = cmp
142 && equals_ptr(
143 shamsys::instance::get_compute_queue(), buf_pos_min_cell, rhs.buf_pos_min_cell);
144 cmp = cmp
145 && equals_ptr(
146 shamsys::instance::get_compute_queue(), buf_pos_max_cell, rhs.buf_pos_max_cell);
147 cmp = cmp
148 && equals_ptr(
149 shamsys::instance::get_compute_queue(),
150 buf_pos_min_cell_flt,
151 rhs.buf_pos_min_cell_flt);
152 cmp = cmp
153 && equals_ptr(
154 shamsys::instance::get_compute_queue(),
155 buf_pos_max_cell_flt,
156 rhs.buf_pos_max_cell_flt);
157
158 return cmp;
159 }
160
161 template class TreeCellRanges<u32, f64_3>;
162 template class TreeCellRanges<u64, f64_3>;
163 template class TreeCellRanges<u32, f32_3>;
164 template class TreeCellRanges<u64, f32_3>;
165 template class TreeCellRanges<u32, i64_3>;
166 template class TreeCellRanges<u64, i64_3>;
167 template class TreeCellRanges<u32, u64_3>;
168 template class TreeCellRanges<u64, u64_3>;
169 template class TreeCellRanges<u32, u32_3>;
170 template class TreeCellRanges<u64, u32_3>;
171
172} // namespace shamrock::tree
shamsys::instance::get_compute_queue
sycl::queue & get_compute_queue(u32 id=0)
Definition NodeInstance.cpp:370
u32
std::uint32_t u32
32 bit unsigned integer
Definition aliases_int.hpp:27
compute_ranges.hpp
equals.hpp
Element-wise equality comparison algorithms for buffers.
fmt_defs.hpp
Define the fmt formatters for sycl::vec.
shamalgs::primitives
namespace for primitive algorithm (e.g. sort, scan, reductions, ...)
Definition ImplControl.hpp:23
shamalgs::primitives::equals_ptr
bool equals_ptr(sycl::queue &q, const std::unique_ptr< sycl::buffer< T > > &buf1, const std::unique_ptr< sycl::buffer< T > > &buf2)
Compare all elements between two unique_ptr-wrapped sycl::buffers.
Definition equals.hpp:386
shamalgs
namespace to contain everything implemented by shamalgs
Definition algorithm.hpp:21
shamsys
namespace for the system handling
Definition CommProtocolException.hpp:21
Generated on Wed May 13 2026 13:21:49 for Shamrock by doxygen 1.9.8