31 template<
class Tgr
idVec>
33 const sham::DeviceScheduler_ptr &dev_sched,
51 const TgridVec *__restrict block_min,
52 const TgridVec *__restrict block_max,
53 u32 *__restrict is_in_ghost) {
66template<
class Tgr
idVec>
68 auto edges = get_edges();
71 auto &sim_box = edges.sim_box.value;
72 auto &patch_data_layers = edges.patch_data_layers;
73 auto &ghost_layers_candidates = edges.ghost_layers_candidates;
74 auto &patch_boxes = edges.patch_boxes;
77 auto &idx_in_ghost = edges.idx_in_ghost;
79 auto dev_sched = shamsys::instance::get_compute_scheduler_ptr();
82 auto paving_function = get_paving(mode, sim_box);
84 auto &patch_data_layers_ref = patch_data_layers.get_const_refs();
87 idx_in_ghost.buffers = ghost_layers_candidates.values.template map<sham::DeviceBuffer<u32>>(
96 auto brecv = patch_boxes.values.get(receiver);
99 = paving_function.f_aabb_inv(brecv, infos.xoff, infos.yoff, infos.zoff);
101 return get_ids_in_ghost(
107 sender_patch.get_obj_cnt());
111template<
class Tgr
idVec>
Header file describing a Node Instance.
std::uint32_t u32
32 bit unsigned integer
std::uint64_t u64
64 bit unsigned integer
Shamrock assertion utility.
A buffer allocated in USM (Unified Shared Memory)
A SYCL queue associated with a device and a context.
void _impl_evaluate_internal()
evaluate the node
virtual std::string _impl_get_tex() const
get the tex of the node
PatchDataLayer container class, the layout is described in patchdata_layout.
This header file contains utility functions related to exception handling in the code.
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.
std::tuple< std::optional< sycl::buffer< u32 > >, u32 > stream_compact(sycl::queue &q, sycl::buffer< u32 > &buf_flags, u32 len)
Stream compaction algorithm.
T & get_check_ref(const std::unique_ptr< T > &ptr, SourceLocation loc=SourceLocation())
Takes a std::unique_ptr and returns a reference to the object it holds. It throws a std::runtime_erro...
A class that references multiple buffers or similar objects.
Axis-Aligned bounding box.
bool is_not_empty() const noexcept
Checks if the AABB is non-empty.
AABB get_intersect(AABB other) const noexcept
Compute the intersection of two AABB.