22 template<
class u_morton>
23 void TreeReducedMortonCodes<u_morton>::build(
27 TreeMortonCodes<u_morton> &morton_codes,
29 bool &one_cell_mode) {
32 shamlog_debug_sycl_ln(
33 "RadixTree",
"reduction algorithm");
41 morton_codes.buf_morton,
46 shamlog_debug_sycl_ln(
48 "reduction results : (before :",
55 if (tree_leaf_count > 1) {
57 shamlog_debug_sycl_ln(
"RadixTree",
"sycl_morton_remap_reduction");
58 buf_tree_morton = std::make_unique<sycl::buffer<u_morton>>(tree_leaf_count);
64 morton_codes.buf_morton,
67 one_cell_mode =
false;
69 }
else if (tree_leaf_count == 1) {
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})
85 template<
class u_morton>
86 bool TreeReducedMortonCodes<u_morton>::operator==(
87 const TreeReducedMortonCodes<u_morton> &rhs)
const {
90 cmp = cmp && (tree_leaf_count == rhs.tree_leaf_count);
94 cmp = cmp && (buf_reduc_index_map->size() == rhs.buf_reduc_index_map->size());
98 shamsys::instance::get_compute_queue(),
100 *rhs.buf_reduc_index_map,
101 buf_reduc_index_map->size());
104 shamsys::instance::get_compute_queue(),
106 *rhs.buf_tree_morton,
112 template class TreeReducedMortonCodes<u32>;
113 template class TreeReducedMortonCodes<u64>;
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.
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.