[Patch] migrate SerialPatchTree & PatchtreeField to sham::DeviceBuffer#1798
[Patch] migrate SerialPatchTree & PatchtreeField to sham::DeviceBuffer#1798TApplencourt wants to merge 15 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>
|
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: |
|
pre-commit.ci autofix |
for more information, see https://pre-commit.ci
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.
| auto tree = sptree.serial_tree_buf->template mirror_to<sham::host>(); | ||
| auto lpid = sptree.linked_patch_ids_buf->template mirror_to<sham::host>(); |
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.
| auto tree = sptree.serial_tree_buf->template mirror_to<sham::host>(); | ||
| auto lpid = sptree.linked_patch_ids_buf->template mirror_to<sham::host>(); |
| 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(); |
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(); |
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>
|
pre-commit.ci autofix |
for more information, see https://pre-commit.ci
Workflow reportworkflow report corresponding to commit 08d1a1b Pre-commit check reportSome failures were detected in base source checks checks. Suggested changesDetailed changes :diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp
index 203a79a8..1d691a2e 100644
--- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp
+++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp
@@ -256,10 +256,9 @@ sham::DeviceBuffer<u64> SerialPatchTree<vec>::compute_patch_owner(
sham::DeviceScheduler_ptr dev_sched, sham::DeviceBuffer<vec> &position_buffer, u32 len) {
sham::DeviceBuffer<u64> new_owned_id(len, dev_sched);
-
- if(len == 0){
- return new_owned_id;
- }
+ if (len == 0) {
+ return new_owned_id;
+ }
using namespace shamrock::patch;
|
for more information, see https://pre-commit.ci
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