[Patch] migrate SerialPatchTree & PatchtreeField to sham::DeviceBuffer#1798
Open
TApplencourt wants to merge 16 commits into
Open
[Patch] migrate SerialPatchTree & PatchtreeField to sham::DeviceBuffer#1798TApplencourt wants to merge 16 commits into
TApplencourt wants to merge 16 commits into
Conversation
Continues the sycl::buffer -> sham::DeviceBuffer migration (Shamrock-code#490, Shamrock-code#672, Shamrock-code#1461). Swap the underlying storage in: - PatchtreeField<T>: internal_buf is now std::optional<DeviceBuffer<T>>. allocate() now requires a DeviceScheduler_ptr. - SerialPatchTree<T>: serial_tree_buf and linked_patch_ids_buf become std::optional<DeviceBuffer<...>>. attach_buf/detach_buf API preserved for caller compatibility. - compute_patch_owner: returns DeviceBuffer<u64>, kernel rewritten with sham::kernel_call + MultiRef. - make_patch_tree_field: takes DeviceScheduler_ptr instead of sycl::queue, reduction loop migrated to sham::kernel_call. - host_for_each_leafs: uses BufferMirror<host> instead of sycl::host_accessor. Propagate the DistributedData<DeviceBuffer<u64>> change through ReattributeDataUtility (compute_new_pid, extract_elements, reatribute_patch_objects). Update callers: BasicSPHGhosts, GSPHGhostHandler, SPHSetup, SPHUtilities, GSPHUtilities. Delete the now-dead reduce_field method and patch_reduc_tree.hpp (PatchFieldReduction had zero callers). Assisted-by: Claude Opus 4.7 <noreply@anthropic.com>
Contributor
|
Thanks @TApplencourt 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: |
Contributor
Author
|
pre-commit.ci autofix |
for more information, see https://pre-commit.ci
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Code Review
This pull request refactors the codebase to replace SYCL buffers and accessors with sham::DeviceBuffer and BufferMirror abstractions, while transitioning from direct SYCL queue submissions to sham::kernel_call. Key changes include updating PatchtreeField and SerialPatchTree to use std::optionalsham::DeviceBuffer and replacing sycl::host_accessor with mirror_tosham::host() or copy_to_stdvec(). Feedback identifies several performance bottlenecks where data transfers or heap allocations occur within loops. Furthermore, the reviewer pointed out multiple safety regressions where std::optional buffers and scheduler pointers are dereferenced without proper validation, which could lead to runtime crashes.
Comment on lines
+303
to
+304
| auto tree = sptree.serial_tree_buf->template mirror_to<sham::host>(); | ||
| auto lpid = sptree.linked_patch_ids_buf->template mirror_to<sham::host>(); |
Contributor
There was a problem hiding this comment.
Creating and destroying BufferMirror objects inside the triple loop over periodic offsets is highly inefficient. Each creation triggers a Device-to-Host copy, and each destruction triggers a Host-to-Device copy (write-back). These mirrors should be moved outside the loops (before line 296) to avoid redundant data transfers.
Comment on lines
+369
to
+370
| auto tree = sptree.serial_tree_buf->template mirror_to<sham::host>(); | ||
| auto lpid = sptree.linked_patch_ids_buf->template mirror_to<sham::host>(); |
Contributor
There was a problem hiding this comment.
Similar to the periodic case, creating mirrors inside the for_each_patch_shift lambda causes repeated data transfers for every shift. These mirrors should be moved outside the for_each_patch_shift call (before line 362).
| bool err_id_in_newid = false; | ||
| { | ||
| sycl::host_accessor nid{newid_buf_map.get(id), sycl::read_only}; | ||
| std::vector<u64> nid = newid_buf_map.get(id).copy_to_stdvec(); |
Contributor
There was a problem hiding this comment.
copy_to_stdvec() performs a heap allocation and a full buffer copy from device to host. Calling it inside for_each_patchdata is inefficient and can become a performance bottleneck as the number of patches increases. Consider if this logic can be moved to a device kernel or if a more efficient host access method is available.
| if (!pdat.is_empty()) { | ||
|
|
||
| sycl::host_accessor nid{new_pid.get(current_pid), sycl::read_only}; | ||
| std::vector<u64> nid = new_pid.get(current_pid).copy_to_stdvec(); |
Contributor
There was a problem hiding this comment.
Similar to the usage in compute_new_pid, calling copy_to_stdvec() inside a loop over patches is inefficient due to repeated allocations and data transfers.
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Member
|
pre-commit.ci autofix |
for more information, see https://pre-commit.ci
tdavidcl
reviewed
May 29, 2026
tdavidcl
added a commit
that referenced
this pull request
May 30, 2026
Contributor
Workflow reportworkflow report corresponding to commit d9ed42a Pre-commit check reportPre-commit check: ✅ Test pipeline can run. Clang-tidy diff reportSuggested changesDetailed changes :diff --git a/src/shamrock/include/shamrock/scheduler/ReattributeDataUtility.hpp b/src/shamrock/include/shamrock/scheduler/ReattributeDataUtility.hpp
index 4aa1f0c44..9459abc90 100644
--- a/src/shamrock/include/shamrock/scheduler/ReattributeDataUtility.hpp
+++ b/src/shamrock/include/shamrock/scheduler/ReattributeDataUtility.hpp
@@ -130,31 +130,7 @@ namespace shamrock {
std::vector<u64> nid = new_pid.get(current_pid).copy_to_stdvec();
- if (false) {
-
- const u32 cnt = pdat.get_obj_cnt();
-
- for (u32 i = cnt - 1; i < cnt; i--) {
- u64 new_pid = nid[i];
- if (current_pid != new_pid) {
-
- if (!part_exchange.has_key(current_pid, new_pid)) {
- part_exchange.add_obj(
- current_pid,
- new_pid,
- PatchDataLayer(sched.get_layout_ptr_old()));
- }
-
- part_exchange.for_each(
- [&](u64 _old_id, u64 _new_id, PatchDataLayer &pdat_int) {
- if (_old_id == current_pid && _new_id == new_pid) {
- pdat.extract_element(i, pdat_int);
- histogram_extract[current_pid]++;
- }
- });
- }
- }
- } else {
+ {
std::vector<u32> keep_ids;
std::unordered_map<u64, std::vector<u32>> extract_indexes;
diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp
index d9fa86142..89acf4573 100644
--- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp
+++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp
@@ -247,14 +247,14 @@ class SerialPatchTree {
}
sham::DeviceBuffer<u64> compute_patch_owner(
- sham::DeviceScheduler_ptr dev_sched,
+ const sham::DeviceScheduler_ptr& dev_sched,
sham::DeviceBuffer<fp_prec_vec> &position_buffer,
u32 len);
};
template<class vec>
sham::DeviceBuffer<u64> SerialPatchTree<vec>::compute_patch_owner(
- sham::DeviceScheduler_ptr dev_sched, sham::DeviceBuffer<vec> &position_buffer, u32 len) {
+ const sham::DeviceScheduler_ptr& dev_sched, sham::DeviceBuffer<vec> &position_buffer, u32 len) {
sham::DeviceBuffer<u64> new_owned_id(len, dev_sched);
if (len == 0) {Detailed changes :- src/shamrock/include/shamrock/patch/PatchField.hpp:24: warning: Compound shamrock::patch::PatchField is not documented.
- src/shamrock/include/shamrock/patch/PatchField.hpp:26: warning: Member field_all (variable) of class shamrock::patch::PatchField is not documented.
+ src/shamrock/include/shamrock/patch/PatchField.hpp:27: warning: Compound shamrock::patch::PatchField is not documented.
- src/shamrock/include/shamrock/patch/PatchField.hpp:28: warning: Member PatchField(shambase::DistributedData< T > &&field_all) (function) of class shamrock::patch::PatchField is not documented.
+ src/shamrock/include/shamrock/patch/PatchField.hpp:29: warning: Member field_all (variable) of class shamrock::patch::PatchField is not documented.
- src/shamrock/include/shamrock/patch/PatchField.hpp:30: warning: Member get(u64 id) (function) of class shamrock::patch::PatchField is not documented.
+ src/shamrock/include/shamrock/patch/PatchField.hpp:31: warning: Member PatchField(shambase::DistributedData< T > &&field_all) (function) of class shamrock::patch::PatchField is not documented.
+ src/shamrock/include/shamrock/patch/PatchField.hpp:33: warning: Member get(u64 id) (function) of class shamrock::patch::PatchField is not documented.
- src/shamrock/include/shamrock/patch/PatchField.hpp:34: warning: Compound shamrock::patch::PatchtreeField is not documented.
- src/shamrock/include/shamrock/patch/PatchField.hpp:36: warning: Member internal_buf (variable) of class shamrock::patch::PatchtreeField is not documented.
+ src/shamrock/include/shamrock/patch/PatchField.hpp:37: warning: Compound shamrock::patch::PatchtreeField is not documented.
- src/shamrock/include/shamrock/patch/PatchField.hpp:38: warning: Member reset() (function) of class shamrock::patch::PatchtreeField is not documented.
+ src/shamrock/include/shamrock/patch/PatchField.hpp:39: warning: Member internal_buf (variable) of class shamrock::patch::PatchtreeField is not documented.
- src/shamrock/include/shamrock/patch/PatchField.hpp:40: warning: Member allocate(u32 size) (function) of class shamrock::patch::PatchtreeField is not documented.
+ src/shamrock/include/shamrock/patch/PatchField.hpp:41: warning: Member reset() (function) of class shamrock::patch::PatchtreeField is not documented.
+ src/shamrock/include/shamrock/patch/PatchField.hpp:43: warning: Member allocate(u32 size, sham::DeviceScheduler_ptr sched) (function) of class shamrock::patch::PatchtreeField is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:103: warning: Member SerialPatchTree(PatchTree &ptree, const shamrock::patch::PatchCoordTransform< fp_prec_vec > box_transform) (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:108: warning: Member SerialPatchTree(PatchTree &ptree, const shamrock::patch::PatchCoordTransform< fp_prec_vec > box_transform) (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:110: warning: Member host_for_each_leafs_internal(std::function< bool(u64, PtNode pnode)> interact_cd, std::function< void(u64, PtNode)> found_case, Acc1 &&tree, Acc2 &&lpid) (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:115: warning: Member host_for_each_leafs_internal(std::function< bool(u64, PtNode pnode)> interact_cd, std::function< void(u64, PtNode)> found_case, Acc1 &&tree, Acc2 &&lpid) (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:147: warning: Member host_for_each_leafs(std::function< bool(u64, PtNode pnode)> interact_cd, std::function< void(u64, PtNode)> found_case) (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:152: warning: Member host_for_each_leafs(std::function< bool(u64, PtNode pnode)> interact_cd, std::function< void(u64, PtNode)> found_case) (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:172: warning: Member build(PatchScheduler &sched) (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:177: warning: Member build(PatchScheduler &sched) (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:178: warning: Member make_patch_tree_field(PatchScheduler &sched, sycl::queue &queue, shamrock::patch::PatchField< T > pfield, Func &&reducer) (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:183: warning: Member make_patch_tree_field(PatchScheduler &sched, sham::DeviceScheduler_ptr dev_sched, shamrock::patch::PatchField< T > pfield, Func &&reducer) (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:223: warning: Member dump_dat() (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:226: warning: Member dump_dat() (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:246: warning: Member compute_patch_owner(sham::DeviceScheduler_ptr dev_sched, sham::DeviceBuffer< fp_prec_vec > &position_buffer, u32 len) (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:249: warning: Member compute_patch_owner(sham::DeviceScheduler_ptr dev_sched, sham::DeviceBuffer< fp_prec_vec > &position_buffer, u32 len) (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:35: warning: Compound SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:37: warning: Member PtNode (typedef) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:38: warning: Compound SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:39: warning: Member PatchTree (typedef) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:40: warning: Member PtNode (typedef) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:42: warning: Member PatchTree (typedef) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:42: warning: Member root_count (variable) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:43: warning: Member serial_tree_buf (variable) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:44: warning: Member linked_patch_ids_buf (variable) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:44: warning: Member root_count (variable) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:45: warning: Member serial_tree_buf (variable) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:46: warning: Member attach_buf() (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:46: warning: Member linked_patch_ids_buf (variable) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:48: warning: Member attach_buf() (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:60: warning: Member detach_buf() (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:65: warning: Member detach_buf() (function) of class SerialPatchTree is not documented.
- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:83: warning: Member print_status() (function) of class SerialPatchTree is not documented.
+ src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp:88: warning: Member print_status() (function) of class SerialPatchTree is not documented. |
Codecov Report❌ Patch coverage is 📢 Thoughts on this report? Let us know! |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Continues the sycl::buffer -> sham::DeviceBuffer migration (#490, #672, #1461). Swap the underlying storage in:
Propagate the DistributedData<DeviceBuffer> change through ReattributeDataUtility (compute_new_pid, extract_elements, reatribute_patch_objects).
Update callers: BasicSPHGhosts, GSPHGhostHandler, SPHSetup, SPHUtilities, GSPHUtilities.
Delete the now-dead reduce_field method and patch_reduc_tree.hpp (PatchFieldReduction had zero callers).
Assisted-by: Claude Opus 4.7 noreply@anthropic.com