Conversation
Mirrors the TwoStage algorithm, but every tree traversal (leaf-to-leaf search and particle parent-leaf lookup) uses a work-group local_accessor slice as its rtree_for stack instead of a per-work-item std::array, launched over nd_range work-groups of 256. Assisted-by: Claude Code
The shared offload single stage neighbour cache now computes the exact tree depth per patch and sizes its local memory stack as depth + 1. The external stack of rtree_for is now given as an (u32 id) -> u32 & functor with a runtime stack size instead of a raw pointer. Assisted-by: Claude Code
|
Important Draft PR not reviewedDraft PRs are not automatically reviewed by default.
To automatically review draft PRs, update your CodeRabbit configuration: reviews:
auto_review:
drafts: trueThanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
|
Thanks @tdavidcl for opening this PR! You can do multiple things directly here: Once the workflow completes a message will appear displaying informations related to the run. Also the PR gets automatically reviewed by gemini, you can: |
leaf_part_id now uses sham::DeviceBuffer instead of a raw sycl::buffer, and the repeated shamsys::instance::get_compute_scheduler().get_queue() calls in each cache-building lambda are hoisted to a single fetch. Also drops now-redundant wait_and_throw() calls. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Replace the hand-rolled EventList/get_read_access/complete_event_state boilerplate around every SYCL launch in NeighbourCache.cpp with sham::kernel_call (plain launches) and sham::kernel_call_hndl (shared-offload launches needing an nd_range + local_accessor stack). Also swaps shambase::make_range for sham::make_ndrange in those nd_range launches.
Replace the two shamalgs::memory::extract_element(queue, buf, idx) calls with buf.get_val_at_idx(idx), dropping the manual queue fetch.
# Conflicts: # src/shammodels/sph/src/modules/NeighbourCache.cpp
# Conflicts: # src/shammodels/sph/src/modules/NeighbourCache.cpp
Workflow reportworkflow report corresponding to commit 71cd334 Light CI is enabled (the default for pull requests). This will only run the basic tests and not the full tests. Pre-commit check reportPre-commit check: ✅ Test pipeline can run. Clang-tidy diff reportSuggested changesDetailed changes :diff --git a/src/shammodels/sph/src/modules/NeighbourCache.cpp b/src/shammodels/sph/src/modules/NeighbourCache.cpp
index 56f65756..e3d40f8a 100644
--- a/src/shammodels/sph/src/modules/NeighbourCache.cpp
+++ b/src/shammodels/sph/src/modules/NeighbourCache.cpp
@@ -264,7 +264,7 @@ void shammodels::sph::modules::NeighbourCache<Tvec, Tmorton, SPHKernel>::
auto particle_looper,
u32 *__restrict neigh_cnt) {
return [=](sycl::handler &cgh) {
- constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
+ constexpr Tscal rker2 = Kernel::Rkern * Kernel::Rkern;
constexpr u32 group_size = 256;
sycl::local_accessor<u32, 1> stack_local(stack_size * group_size, cgh);
@@ -313,8 +313,8 @@ void shammodels::sph::modules::NeighbourCache<Tvec, Tmorton, SPHKernel>::
Tscal rab2 = sycl::dot(dr, dr);
Tscal rint_b = hpart[id_b] * h_tolerance;
- bool no_interact = rab2 > rint_a * rint_a * Rker2
- && rab2 > rint_b * rint_b * Rker2;
+ bool no_interact = rab2 > rint_a * rint_a * rker2
+ && rab2 > rint_b * rint_b * rker2;
cnt += (no_interact) ? 0 : 1;
});
@@ -341,7 +341,7 @@ void shammodels::sph::modules::NeighbourCache<Tvec, Tmorton, SPHKernel>::
auto particle_looper,
u32 *__restrict neigh) {
return [=](sycl::handler &cgh) {
- constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
+ constexpr Tscal rker2 = Kernel::Rkern * Kernel::Rkern;
constexpr u32 group_size = 256;
sycl::local_accessor<u32, 1> stack_local(stack_size * group_size, cgh);
@@ -390,8 +390,8 @@ void shammodels::sph::modules::NeighbourCache<Tvec, Tmorton, SPHKernel>::
Tscal rab2 = sycl::dot(dr, dr);
Tscal rint_b = hpart[id_b] * h_tolerance;
- bool no_interact = rab2 > rint_a * rint_a * Rker2
- && rab2 > rint_b * rint_b * Rker2;
+ bool no_interact = rab2 > rint_a * rint_a * rker2
+ && rab2 > rint_b * rint_b * rker2;
if (!no_interact) {
neigh[cnt] = id_b;
@@ -1009,7 +1009,7 @@ void shammodels::sph::modules::NeighbourCache<Tvec, Tmorton, SPHKernel>::
u32 offset_leaf = intnode_cnt;
- constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
+ constexpr Tscal rker2 = Kernel::Rkern * Kernel::Rkern;
Tscal rint_a = hpart[id_a] * h_tolerance;
@@ -1028,7 +1028,7 @@ void shammodels::sph::modules::NeighbourCache<Tvec, Tmorton, SPHKernel>::
Tscal rint_b = hpart[id_b] * h_tolerance;
bool no_interact
- = rab2 > rint_a * rint_a * Rker2 && rab2 > rint_b * rint_b * Rker2;
+ = rab2 > rint_a * rint_a * rker2 && rab2 > rint_b * rint_b * rker2;
cnt += (no_interact) ? 0 : 1;
});
@@ -1065,7 +1065,7 @@ void shammodels::sph::modules::NeighbourCache<Tvec, Tmorton, SPHKernel>::
u32 offset_leaf = intnode_cnt;
- constexpr Tscal Rker2 = Kernel::Rkern * Kernel::Rkern;
+ constexpr Tscal rker2 = Kernel::Rkern * Kernel::Rkern;
Tscal rint_a = hpart[id_a] * h_tolerance;
@@ -1084,7 +1084,7 @@ void shammodels::sph::modules::NeighbourCache<Tvec, Tmorton, SPHKernel>::
Tscal rint_b = hpart[id_b] * h_tolerance;
bool no_interact
- = rab2 > rint_a * rint_a * Rker2 && rab2 > rint_b * rint_b * Rker2;
+ = rab2 > rint_a * rint_a * rker2 && rab2 > rint_b * rint_b * rker2;
if (!no_interact) {
neigh[cnt] = id_b;
diff --git a/src/shamtree/include/shamtree/KarrasTreeTraverser.hpp b/src/shamtree/include/shamtree/KarrasTreeTraverser.hpp
index 694ec005..f2d491cc 100644
--- a/src/shamtree/include/shamtree/KarrasTreeTraverser.hpp
+++ b/src/shamtree/include/shamtree/KarrasTreeTraverser.hpp
@@ -148,7 +148,7 @@ struct shamtree::KarrasTreeTraverserAccessed {
Functor2 &&on_found_leaf,
Functor3 &&on_excluded_node) const {
- static constexpr u32 _nindex = 4294967295;
+ static constexpr u32 nindex = 4294967295;
// Init the stack state
u32 stack_cursor = stack_size - 1;
@@ -159,7 +159,7 @@ struct shamtree::KarrasTreeTraverserAccessed {
// Pop the top of the stack
u32 current_node_id = stack(stack_cursor);
- stack(stack_cursor) = _nindex;
+ stack(stack_cursor) = nindex;
stack_cursor++;
// check iteraction creteriaDetailed changes :- src/shammodels/common/include/shammodels/common/config/enum_NeighCacheStrategy.hpp:39: warning: Member from_json(const BasicJsonType &j, NeighCacheStrategy &e) (function) of namespace shammodels is not documented.
- src/shammodels/common/include/shammodels/common/config/enum_NeighCacheStrategy.hpp:39: warning: Member to_json(BasicJsonType &j, const NeighCacheStrategy &e) (function) of namespace shammodels is not documented.
+ src/shammodels/common/include/shammodels/common/config/enum_NeighCacheStrategy.hpp:45: warning: Member from_json(const BasicJsonType &j, NeighCacheStrategy &e) (function) of namespace shammodels is not documented.
+ src/shammodels/common/include/shammodels/common/config/enum_NeighCacheStrategy.hpp:45: warning: Member to_json(BasicJsonType &j, const NeighCacheStrategy &e) (function) of namespace shammodels is not documented.
+ src/shammodels/sph/include/shammodels/sph/modules/NeighbourCache.hpp:46: warning: Member start_neighbors_cache_shared_offload() (function) of class shammodels::sph::modules::NeighbourCache is not documented.
+ src/shammodels/sph/include/shammodels/sph/modules/NeighbourCache.hpp:47: warning: Member start_neighbors_cache_2stages_shared_offload() (function) of class shammodels::sph::modules::NeighbourCache is not documented.
- src/shammodels/sph/src/Solver.cpp:2236: warning: Member map_field_refs(PatchScheduler &sched, u32 field_idx, shamrock::solvergraph::FieldRefs< T > &refs) (function) of file Solver.cpp is not documented.
+ src/shammodels/sph/src/Solver.cpp:2242: warning: Member map_field_refs(PatchScheduler &sched, u32 field_idx, shamrock::solvergraph::FieldRefs< T > &refs) (function) of file Solver.cpp is not documented.
- src/shammodels/sph/src/Solver.cpp:2251: warning: Member map_field_refs_ext(PatchScheduler &sched, shambase::DistributedData< shamrock::patch::PatchDataLayer > &mpdats, u32 field_idx, shamrock::solvergraph::FieldRefs< T > &refs) (function) of file Solver.cpp is not documented.
+ src/shammodels/sph/src/Solver.cpp:2257: warning: Member map_field_refs_ext(PatchScheduler &sched, shambase::DistributedData< shamrock::patch::PatchDataLayer > &mpdats, u32 field_idx, shamrock::solvergraph::FieldRefs< T > &refs) (function) of file Solver.cpp is not documented.
- src/shammodels/sph/src/Solver.cpp:2270: warning: Member map_field_refs_ext(PatchScheduler &sched, shamrock::ComputeField< T > &field_data, shamrock::solvergraph::FieldRefs< T > &refs) (function) of file Solver.cpp is not documented.
+ src/shammodels/sph/src/Solver.cpp:2276: warning: Member map_field_refs_ext(PatchScheduler &sched, shamrock::ComputeField< T > &field_data, shamrock::solvergraph::FieldRefs< T > &refs) (function) of file Solver.cpp is not documented.
- src/shamtree/include/shamtree/CompressedLeafBVH.hpp:125: warning: Member rebuild_from_position_range(sham::DeviceBuffer< Tvec > &min, sham::DeviceBuffer< Tvec > &max, u32 obj_cnt, shammath::AABB< Tvec > &bounding_box, u32 compression_level) (function) of class shamtree::CompressedLeafBVH is not documented.
- src/shamtree/include/shamtree/CompressedLeafBVH.hpp:132: warning: Member rebuild_from_position_range(sham::DeviceBuffer< Tvec > &min, sham::DeviceBuffer< Tvec > &max, shammath::AABB< Tvec > &bounding_box, u32 compression_level) (function) of class shamtree::CompressedLeafBVH is not documented.
+ src/shamtree/include/shamtree/CompressedLeafBVH.hpp:137: warning: Member rebuild_from_position_range(sham::DeviceBuffer< Tvec > &min, sham::DeviceBuffer< Tvec > &max, u32 obj_cnt, shammath::AABB< Tvec > &bounding_box, u32 compression_level) (function) of class shamtree::CompressedLeafBVH is not documented.
- src/shamtree/include/shamtree/CompressedLeafBVH.hpp:138: warning: Member get_traverser() const (function) of class shamtree::CompressedLeafBVH is not documented.
- src/shamtree/include/shamtree/CompressedLeafBVH.hpp:142: warning: Member get_traverser_host() const (function) of class shamtree::CompressedLeafBVH is not documented.
+ src/shamtree/include/shamtree/CompressedLeafBVH.hpp:144: warning: Member rebuild_from_position_range(sham::DeviceBuffer< Tvec > &min, sham::DeviceBuffer< Tvec > &max, shammath::AABB< Tvec > &bounding_box, u32 compression_level) (function) of class shamtree::CompressedLeafBVH is not documented.
+ src/shamtree/include/shamtree/CompressedLeafBVH.hpp:150: warning: Member get_traverser() const (function) of class shamtree::CompressedLeafBVH is not documented.
+ src/shamtree/include/shamtree/CompressedLeafBVH.hpp:154: warning: Member get_traverser_host() const (function) of class shamtree::CompressedLeafBVH is not documented.
- src/shamtree/include/shamtree/CompressedLeafBVH.hpp:164: warning: Member get_object_iterator_host() const (function) of class shamtree::CompressedLeafBVH is not documented.
- src/shamtree/include/shamtree/CompressedLeafBVH.hpp:168: warning: Member get_cell_iterator() const (function) of class shamtree::CompressedLeafBVH is not documented.
- src/shamtree/include/shamtree/CompressedLeafBVH.hpp:173: warning: Member get_cell_iterator_host() const (function) of class shamtree::CompressedLeafBVH is not documented.
+ src/shamtree/include/shamtree/CompressedLeafBVH.hpp:176: warning: Member get_object_iterator_host() const (function) of class shamtree::CompressedLeafBVH is not documented.
+ src/shamtree/include/shamtree/CompressedLeafBVH.hpp:180: warning: Member get_cell_iterator() const (function) of class shamtree::CompressedLeafBVH is not documented.
+ src/shamtree/include/shamtree/CompressedLeafBVH.hpp:185: warning: Member get_cell_iterator_host() const (function) of class shamtree::CompressedLeafBVH is not documented.
- src/shamtree/include/shamtree/KarrasTreeTraverser.hpp:168: warning: Compound shamtree::KarrasTreeTraverserHost is not documented.
+ src/shamtree/include/shamtree/KarrasTreeTraverser.hpp:244: warning: Compound shamtree::KarrasTreeTraverserHost is not documented.
+ src/shamtree/include/shamtree/TreeTraversal.hpp:482: warning: Compound shamrock::tree::ObjectCacheIterator is not documented.
- src/shamtree/include/shamtree/TreeTraversal.hpp:487: warning: Compound shamrock::tree::ObjectCacheIterator is not documented.
+ src/shamtree/include/shamtree/TreeTraversal.hpp:490: warning: Member ObjectCacheIterator(const ObjectCache::ptrs_read &cache) (function) of class shamrock::tree::ObjectCacheIterator is not documented.
- src/shamtree/include/shamtree/TreeTraversal.hpp:495: warning: Member ObjectCacheIterator(const ObjectCache::ptrs_read &cache) (function) of class shamrock::tree::ObjectCacheIterator is not documented.
+ src/shamtree/include/shamtree/TreeTraversal.hpp:498: warning: Member for_each_object(u32 idx, Functor_iter &&func_it) const (function) of class shamrock::tree::ObjectCacheIterator is not documented.
- src/shamtree/include/shamtree/TreeTraversal.hpp:503: warning: Member for_each_object(u32 idx, Functor_iter &&func_it) const (function) of class shamrock::tree::ObjectCacheIterator is not documented.
+ src/shamtree/include/shamtree/TreeTraversal.hpp:510: warning: Member for_each_object_with_id(u32 idx, Functor_iter &&func_it) const (function) of class shamrock::tree::ObjectCacheIterator is not documented.
- src/shamtree/include/shamtree/TreeTraversal.hpp:515: warning: Member for_each_object_with_id(u32 idx, Functor_iter &&func_it) const (function) of class shamrock::tree::ObjectCacheIterator is not documented. |
No description provided.