Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
MortonCodeSet.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
17#include "shambase/integer.hpp"
18#include "shambase/print.hpp"
19#include "shambase/string.hpp"
22#include "shamcomm/logs.hpp"
25#include <stdexcept>
26#include <string>
27#include <utility>
28
29namespace shamtree {
30
31 namespace details {
32 template<class Tmorton, class Tvec, u32 dim>
34
35 public:
38
39 using pos_t = Tvec;
40 using coord_t = typename shambase::VectorProperties<pos_t>::component_type;
41 using ipos_t = typename Morton::int_vec_repr;
42 using int_t = typename Morton::int_vec_repr_base;
43
44 using CoordTransform
46
47 inline static CoordTransform get_transform(
48 pos_t bounding_box_min, pos_t bounding_box_max) {
49 return MortonConvert::get_transform(bounding_box_min, bounding_box_max);
50 }
51 inline static ipos_t to_morton_grid(pos_t pos, CoordTransform transform) {
52 return MortonConvert::to_morton_grid(pos, transform);
53 }
54
55 inline static pos_t to_real_space(ipos_t pos, CoordTransform transform) {
56 return MortonConvert::to_real_space(pos, transform);
57 }
58 };
59
60 } // namespace details
61
62 template<class Tmorton, class Tvec, u32 dim>
64 const sham::DeviceScheduler_ptr &dev_sched,
65 shammath::AABB<Tvec> bounding_box,
67 u32 cnt_obj,
68 u32 morton_count,
69 sham::DeviceBuffer<Tmorton> &&cache_buf_morton_codes) {
70
72 = std::forward<sham::DeviceBuffer<Tmorton>>(cache_buf_morton_codes);
73
74 morton_codes.resize(morton_count);
75
76 if (morton_count < cnt_obj) {
78 "MortonCodeSet: morton_count < cnt_obj\n morton_count: {}, cnt_obj: {}",
79 morton_count,
80 cnt_obj));
81 }
82
84
85 auto transform = Utils::get_transform(bounding_box.lower, bounding_box.upper);
86
88 dev_sched->get_queue(),
89 sham::MultiRef{pos_buf},
90 sham::MultiRef{morton_codes},
91 cnt_obj,
92 [transf = transform,
93 bb = bounding_box](u32 i, const Tvec *__restrict pos, Tmorton *__restrict morton) {
94 Tvec r = pos[i];
95 auto m = Utils::to_morton_grid(bb.clamp_coord(r), transf);
96 auto mcode = Utils::Morton::icoord_to_morton(m.x(), m.y(), m.z());
97
98 morton[i] = mcode;
99 });
100
101 if (morton_count > cnt_obj) {
103 dev_sched->get_queue(),
105 sham::MultiRef{morton_codes},
106 morton_count - cnt_obj,
108 cnt_obj = cnt_obj](u32 i, Tmorton *__restrict morton) {
109 morton[cnt_obj + i] = err_code;
110 });
111 }
112
114 std::move(bounding_box),
115 std::move(cnt_obj),
116 std::move(morton_count),
117 std::move(morton_codes));
118 }
119
120 template<class Tmorton, class Tvec, u32 dim>
122 const sham::DeviceScheduler_ptr &dev_sched,
123 shammath::AABB<Tvec> bounding_box,
125 u32 cnt_obj,
126 u32 morton_count) {
127
128 sham::DeviceBuffer<Tmorton> morton_codes(morton_count, dev_sched);
129
130 return morton_code_set_from_positions<Tmorton, Tvec, dim>(
131 dev_sched, bounding_box, pos_buf, cnt_obj, morton_count, std::move(morton_codes));
132 }
133
134} // namespace shamtree
135
138
140 const sham::DeviceScheduler_ptr &dev_sched,
141 shammath::AABB<f64_3> bounding_box,
143 u32 cnt_obj,
144 u32 morton_count);
145
147 const sham::DeviceScheduler_ptr &dev_sched,
148 shammath::AABB<f64_3> bounding_box,
150 u32 cnt_obj,
151 u32 morton_count);
152
154 const sham::DeviceScheduler_ptr &dev_sched,
155 shammath::AABB<f64_3> bounding_box,
157 u32 cnt_obj,
158 u32 morton_count,
159 sham::DeviceBuffer<u32> &&cache_buf_morton_codes);
160
162 const sham::DeviceScheduler_ptr &dev_sched,
163 shammath::AABB<f64_3> bounding_box,
165 u32 cnt_obj,
166 u32 morton_count,
167 sham::DeviceBuffer<u64> &&cache_buf_morton_codes);
MortonCodeSet< Tmorton, Tvec, dim > morton_code_set_from_positions(const sham::DeviceScheduler_ptr &dev_sched, shammath::AABB< Tvec > bounding_box, sham::DeviceBuffer< Tvec > &pos_buf, u32 cnt_obj, u32 morton_count, sham::DeviceBuffer< Tmorton > &&cache_buf_morton_codes)
Generate a set of Morton codes from a buffer of positions.
std::uint32_t u32
32 bit unsigned integer
A buffer allocated in USM (Unified Shared Memory)
void resize(size_t new_size, bool keep_data=true)
Resizes the buffer to a given size.
Class representing a set of Morton codes with associated bounding box and position data.
This header file contains utility functions related to exception handling in the code.
Define the fmt formatters for sycl::vec.
Morton curve implementation.
Namespace for internal details of the logs module.
void kernel_call(sham::DeviceQueue &q, RefIn in, RefOut in_out, u32 n, Functor &&func, SourceLocation &&callsite=SourceLocation{})
Submit a kernel to a SYCL queue.
void throw_with_loc(std::string message, SourceLocation loc=SourceLocation{})
Throw an exception and append the source location to it.
A class that references multiple buffers or similar objects.
Axis-Aligned bounding box.
Definition AABB.hpp:99
T lower
Lower bound of the AABB.
Definition AABB.hpp:104
T upper
Upper bound of the AABB.
Definition AABB.hpp:105