32 template<
class Tmorton,
class Tvec, u32 dim>
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;
48 pos_t bounding_box_min, pos_t bounding_box_max) {
49 return MortonConvert::get_transform(bounding_box_min, bounding_box_max);
51 inline static ipos_t to_morton_grid(pos_t pos,
CoordTransform transform) {
52 return MortonConvert::to_morton_grid(pos, transform);
55 inline static pos_t to_real_space(ipos_t pos,
CoordTransform transform) {
56 return MortonConvert::to_real_space(pos, transform);
62 template<
class Tmorton,
class Tvec, u32 dim>
64 const sham::DeviceScheduler_ptr &dev_sched,
72 = std::forward<sham::DeviceBuffer<Tmorton>>(cache_buf_morton_codes);
74 morton_codes.
resize(morton_count);
76 if (morton_count < cnt_obj) {
78 "MortonCodeSet: morton_count < cnt_obj\n morton_count: {}, cnt_obj: {}",
85 auto transform = Utils::get_transform(bounding_box.
lower, bounding_box.
upper);
88 dev_sched->get_queue(),
93 bb = bounding_box](
u32 i,
const Tvec *__restrict pos, Tmorton *__restrict morton) {
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());
101 if (morton_count > cnt_obj) {
103 dev_sched->get_queue(),
106 morton_count - cnt_obj,
108 cnt_obj = cnt_obj](
u32 i, Tmorton *__restrict morton) {
109 morton[cnt_obj + i] = err_code;
114 std::move(bounding_box),
116 std::move(morton_count),
117 std::move(morton_codes));
120 template<
class Tmorton,
class Tvec, u32 dim>
122 const sham::DeviceScheduler_ptr &dev_sched,
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));
140 const sham::DeviceScheduler_ptr &dev_sched,
147 const sham::DeviceScheduler_ptr &dev_sched,
154 const sham::DeviceScheduler_ptr &dev_sched,
162 const sham::DeviceScheduler_ptr &dev_sched,
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.
T lower
Lower bound of the AABB.
T upper
Upper bound of the AABB.