From 2fb03d1fc1bf6ae52f31686121d6d4bd5fb68767 Mon Sep 17 00:00:00 2001 From: tapplencourt Date: Fri, 8 May 2026 14:17:37 +0000 Subject: [PATCH 01/10] [Patch] migrate SerialPatchTree & PatchtreeField to sham::DeviceBuffer Continues the sycl::buffer -> sham::DeviceBuffer migration (#490, #672, #1461). Swap the underlying storage in: - PatchtreeField: internal_buf is now std::optional>. allocate() now requires a DeviceScheduler_ptr. - SerialPatchTree: serial_tree_buf and linked_patch_ids_buf become std::optional>. attach_buf/detach_buf API preserved for caller compatibility. - compute_patch_owner: returns DeviceBuffer, 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 instead of sycl::host_accessor. Propagate the DistributedData> 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 --- .../shammodels/gsph/modules/GSPHUtilities.hpp | 2 +- .../gsph/src/modules/GSPHGhostHandler.cpp | 10 +- .../include/shammodels/sph/SPHUtilities.hpp | 2 +- src/shammodels/sph/src/BasicSPHGhosts.cpp | 27 +-- src/shammodels/sph/src/modules/SPHSetup.cpp | 4 +- .../legacy/patch/utility/patch_reduc_tree.hpp | 65 ------ .../include/shamrock/patch/PatchField.hpp | 9 +- .../scheduler/ReattributeDataUtility.hpp | 15 +- .../shamrock/scheduler/SerialPatchTree.hpp | 219 +++++------------- 9 files changed, 89 insertions(+), 264 deletions(-) delete mode 100644 src/shamrock/include/shamrock/legacy/patch/utility/patch_reduc_tree.hpp diff --git a/src/shammodels/gsph/include/shammodels/gsph/modules/GSPHUtilities.hpp b/src/shammodels/gsph/include/shammodels/gsph/modules/GSPHUtilities.hpp index f96ae01d2..9ea0f3e07 100644 --- a/src/shammodels/gsph/include/shammodels/gsph/modules/GSPHUtilities.hpp +++ b/src/shammodels/gsph/include/shammodels/gsph/modules/GSPHUtilities.hpp @@ -77,7 +77,7 @@ namespace shammodels::gsph { PatchtreeField interactR_mpi_tree = sptree.make_patch_tree_field( sched, - shamsys::instance::get_compute_queue(), + shamsys::instance::get_compute_scheduler_ptr(), interactR_patch, [](Tscal h0, Tscal h1, Tscal h2, Tscal h3, Tscal h4, Tscal h5, Tscal h6, Tscal h7) { return sham::max_8points(h0, h1, h2, h3, h4, h5, h6, h7); diff --git a/src/shammodels/gsph/src/modules/GSPHGhostHandler.cpp b/src/shammodels/gsph/src/modules/GSPHGhostHandler.cpp index 070fc6f85..381f91946 100644 --- a/src/shammodels/gsph/src/modules/GSPHGhostHandler.cpp +++ b/src/shammodels/gsph/src/modules/GSPHGhostHandler.cpp @@ -15,6 +15,7 @@ */ #include "shambase/exception.hpp" +#include "shambackends/BufferMirror.hpp" #include "shamcomm/collectives.hpp" #include "shammodels/gsph/modules/GSPHGhostHandler.hpp" #include @@ -133,8 +134,7 @@ auto GSPHGhostHandler::find_interfaces( using BCShearingPeriodic = typename CfgClass::ShearingPeriodic; if (BCPeriodic *cfg = std::get_if(&ghost_config)) { - sycl::host_accessor acc_tf{ - shambase::get_check_ref(int_range_max_tree.internal_buf), sycl::read_only}; + auto acc_tf = int_range_max_tree.internal_buf->template mirror_to(); for (i32 xoff = -repetition_x; xoff <= repetition_x; xoff++) { for (i32 yoff = -repetition_y; yoff <= repetition_y; yoff++) { @@ -188,8 +188,7 @@ auto GSPHGhostHandler::find_interfaces( } } } else if (BCShearingPeriodic *cfg = std::get_if(&ghost_config)) { - sycl::host_accessor acc_tf{ - shambase::get_check_ref(int_range_max_tree.internal_buf), sycl::read_only}; + auto acc_tf = int_range_max_tree.internal_buf->template mirror_to(); for_each_patch_shift(*cfg, bsize, [&](i32_3 ioff, ShiftInfo shift) { i32 xoff = ioff.x(); @@ -242,8 +241,7 @@ auto GSPHGhostHandler::find_interfaces( }); } else { - sycl::host_accessor acc_tf{ - shambase::get_check_ref(int_range_max_tree.internal_buf), sycl::read_only}; + auto acc_tf = int_range_max_tree.internal_buf->template mirror_to(); // sender translation vec periodic_offset = vec{0, 0, 0}; diff --git a/src/shammodels/sph/include/shammodels/sph/SPHUtilities.hpp b/src/shammodels/sph/include/shammodels/sph/SPHUtilities.hpp index af5f17a9b..2b8d0c835 100644 --- a/src/shammodels/sph/include/shammodels/sph/SPHUtilities.hpp +++ b/src/shammodels/sph/include/shammodels/sph/SPHUtilities.hpp @@ -94,7 +94,7 @@ namespace shammodels::sph { PatchtreeField interactR_mpi_tree = sptree.make_patch_tree_field( sched, - shamsys::instance::get_compute_queue(), + shamsys::instance::get_compute_scheduler_ptr(), interactR_patch, [](flt h0, flt h1, flt h2, flt h3, flt h4, flt h5, flt h6, flt h7) { return sham::max_8points(h0, h1, h2, h3, h4, h5, h6, h7); diff --git a/src/shammodels/sph/src/BasicSPHGhosts.cpp b/src/shammodels/sph/src/BasicSPHGhosts.cpp index 6b6357761..0972dcae5 100644 --- a/src/shammodels/sph/src/BasicSPHGhosts.cpp +++ b/src/shammodels/sph/src/BasicSPHGhosts.cpp @@ -154,6 +154,7 @@ int main(){ #include "shambase/exception.hpp" #include "shambase/time.hpp" #include "shamalgs/collective/reduction.hpp" +#include "shambackends/BufferMirror.hpp" #include "shamcomm/collectives.hpp" #include "shamcomm/worldInfo.hpp" #include "shammodels/sph/BasicSPHGhosts.hpp" @@ -290,8 +291,7 @@ auto BasicSPHGhostHandler::find_interfaces( base_timer.start(); if (BCPeriodic *cfg = std::get_if(&ghost_config)) { - sycl::host_accessor acc_tf{ - shambase::get_check_ref(int_range_max_tree.internal_buf), sycl::read_only}; + auto acc_tf = int_range_max_tree.internal_buf->template mirror_to(); for (i32 xoff = -repetition_x; xoff <= repetition_x; xoff++) { for (i32 yoff = -repetition_y; yoff <= repetition_y; yoff++) { @@ -300,10 +300,8 @@ auto BasicSPHGhostHandler::find_interfaces( // sender translation vec periodic_offset = vec{xoff * bsize.x(), yoff * bsize.y(), zoff * bsize.z()}; - sycl::host_accessor tree{ - shambase::get_check_ref(sptree.serial_tree_buf), sycl::read_only}; - sycl::host_accessor lpid{ - shambase::get_check_ref(sptree.linked_patch_ids_buf), sycl::read_only}; + auto tree = sptree.serial_tree_buf->template mirror_to(); + auto lpid = sptree.linked_patch_ids_buf->template mirror_to(); #pragma omp parallel for for (u32 i = 0; i < sched.patch_list.local.size(); i++) { @@ -359,8 +357,7 @@ auto BasicSPHGhostHandler::find_interfaces( } } } else if (BCShearingPeriodic *cfg = std::get_if(&ghost_config)) { - sycl::host_accessor acc_tf{ - shambase::get_check_ref(int_range_max_tree.internal_buf), sycl::read_only}; + auto acc_tf = int_range_max_tree.internal_buf->template mirror_to(); for_each_patch_shift(*cfg, bsize, [&](i32_3 ioff, ShiftInfo shift) { i32 xoff = ioff.x(); @@ -369,10 +366,8 @@ auto BasicSPHGhostHandler::find_interfaces( vec offset = shift.shift; - sycl::host_accessor tree{ - shambase::get_check_ref(sptree.serial_tree_buf), sycl::read_only}; - sycl::host_accessor lpid{ - shambase::get_check_ref(sptree.linked_patch_ids_buf), sycl::read_only}; + auto tree = sptree.serial_tree_buf->template mirror_to(); + auto lpid = sptree.linked_patch_ids_buf->template mirror_to(); #pragma omp parallel for for (u32 i = 0; i < sched.patch_list.local.size(); i++) { @@ -429,14 +424,12 @@ auto BasicSPHGhostHandler::find_interfaces( }); } else { - sycl::host_accessor acc_tf{ - shambase::get_check_ref(int_range_max_tree.internal_buf), sycl::read_only}; + auto acc_tf = int_range_max_tree.internal_buf->template mirror_to(); // sender translation vec periodic_offset = vec{0, 0, 0}; - sycl::host_accessor tree{shambase::get_check_ref(sptree.serial_tree_buf), sycl::read_only}; - sycl::host_accessor lpid{ - shambase::get_check_ref(sptree.linked_patch_ids_buf), sycl::read_only}; + auto tree = sptree.serial_tree_buf->template mirror_to(); + auto lpid = sptree.linked_patch_ids_buf->template mirror_to(); #pragma omp parallel for for (u32 i = 0; i < sched.patch_list.local.size(); i++) { diff --git a/src/shammodels/sph/src/modules/SPHSetup.cpp b/src/shammodels/sph/src/modules/SPHSetup.cpp index f50dac7a7..406b5e518 100644 --- a/src/shammodels/sph/src/modules/SPHSetup.cpp +++ b/src/shammodels/sph/src/modules/SPHSetup.cpp @@ -508,7 +508,7 @@ void shammodels::sph::modules::SPHSetup::apply_setup_new( shambase::throw_unimplemented(); } - sycl::buffer new_id_buf = sptree.compute_patch_owner( + sham::DeviceBuffer new_id_buf = sptree.compute_patch_owner( shamsys::instance::get_compute_scheduler_ptr(), pos_field.get_buf(), pos_field.get_obj_cnt()); @@ -516,7 +516,7 @@ void shammodels::sph::modules::SPHSetup::apply_setup_new( std::unordered_map> index_per_ranks; bool err_id_in_newid = false; { - sycl::host_accessor nid{new_id_buf, sycl::read_only}; + std::vector nid = new_id_buf.copy_to_stdvec(); for (u32 i = 0; i < pos_field.get_obj_cnt(); i++) { u64 patch_id = nid[i]; bool err = patch_id == u64_max; diff --git a/src/shamrock/include/shamrock/legacy/patch/utility/patch_reduc_tree.hpp b/src/shamrock/include/shamrock/legacy/patch/utility/patch_reduc_tree.hpp deleted file mode 100644 index 3f8b6b4b9..000000000 --- a/src/shamrock/include/shamrock/legacy/patch/utility/patch_reduc_tree.hpp +++ /dev/null @@ -1,65 +0,0 @@ -// -------------------------------------------------------// -// -// SHAMROCK code for hydrodynamics -// Copyright (c) 2021-2026 Timothée David--Cléris -// SPDX-License-Identifier: CeCILL Free Software License Agreement v2.1 -// Shamrock is licensed under the CeCILL 2.1 License, see LICENSE for more information -// -// -------------------------------------------------------// - -#pragma once -/** - * @file patch_reduc_tree.hpp - * @author Timothée David--Cléris (tim.shamrock@proton.me) - * @brief - * - */ - -#include "shambase/exception.hpp" -#include "shamrock/scheduler/PatchTree.hpp" -#include "shamsys/legacy/sycl_handler.hpp" -#include -#include - -template -class PatchFieldReduction { - public: - std::vector tree_field; - sycl::buffer *tree_field_buf = nullptr; - - inline void attach_buf() { - if (tree_field_buf != nullptr) - throw shambase::make_except_with_loc( - "tree_field_buf is already allocated"); - tree_field_buf = new sycl::buffer(tree_field.data(), tree_field.size()); - } - - inline void detach_buf() { - if (tree_field_buf == nullptr) - throw shambase::make_except_with_loc( - "tree_field_buf wasn't allocated"); - delete tree_field_buf; - tree_field_buf = nullptr; - } - - // inline void octtree_reduction( - // sycl::queue & queue, - // SerialPatchTree & sptree, - // SchedulerMPI & sched){ - - // std::unordered_map & idp_to_gid = sched.patch_list.id_patch_to_global_idx; - - // sycl::range<1> range{sptree.get_element_count()}; - - // for (u32 level = 0; level < sptree.get_level_count(); level ++) { - // queue.submit([&](sycl::handler &cgh) { - - // cgh.parallel_for(range, [=](sycl::item<1> item) { - // u64 i = (u64)item.get_id(0); - - // }); - // }); - // } - - // } -}; diff --git a/src/shamrock/include/shamrock/patch/PatchField.hpp b/src/shamrock/include/shamrock/patch/PatchField.hpp index ce010143e..f68b3edcb 100644 --- a/src/shamrock/include/shamrock/patch/PatchField.hpp +++ b/src/shamrock/include/shamrock/patch/PatchField.hpp @@ -16,8 +16,11 @@ */ #include "shambase/DistributedData.hpp" +#include "shambackends/DeviceBuffer.hpp" +#include "shambackends/DeviceScheduler.hpp" #include "shambackends/sycl.hpp" #include +#include namespace shamrock::patch { template @@ -33,10 +36,12 @@ namespace shamrock::patch { template class PatchtreeField { public: - std::unique_ptr> internal_buf; + std::optional> internal_buf; inline void reset() { internal_buf.reset(); } - inline void allocate(u32 size) { internal_buf = std::make_unique>(size); } + inline void allocate(u32 size, sham::DeviceScheduler_ptr sched) { + internal_buf.emplace(size, sched); + } }; } // namespace shamrock::patch diff --git a/src/shamrock/include/shamrock/scheduler/ReattributeDataUtility.hpp b/src/shamrock/include/shamrock/scheduler/ReattributeDataUtility.hpp index 18d9b396b..4aa1f0c44 100644 --- a/src/shamrock/include/shamrock/scheduler/ReattributeDataUtility.hpp +++ b/src/shamrock/include/shamrock/scheduler/ReattributeDataUtility.hpp @@ -59,12 +59,12 @@ namespace shamrock { * bound). */ template - shambase::DistributedData> compute_new_pid( + shambase::DistributedData> compute_new_pid( SerialPatchTree &sptree, u32 ipos) { StackEntry stack_loc{}; - shambase::DistributedData> newid_buf_map; + shambase::DistributedData> newid_buf_map; sched.patch_data.for_each_patchdata([&](u64 id, shamrock::patch::PatchDataLayer &pdat) { if (!pdat.is_empty()) { @@ -84,7 +84,7 @@ namespace shamrock { bool err_id_in_newid = false; { - sycl::host_accessor nid{newid_buf_map.get(id), sycl::read_only}; + std::vector nid = newid_buf_map.get(id).copy_to_stdvec(); for (u32 i = 0; i < pdat.get_obj_cnt(); i++) { bool err = nid[i] == u64_max; err_id_in_newid = err_id_in_newid || (err); @@ -114,7 +114,7 @@ namespace shamrock { * @return A shared distributed data object containing the extracted patch data. */ inline shambase::DistributedDataShared extract_elements( - shambase::DistributedData> new_pid) { + shambase::DistributedData> &&new_pid) { shambase::DistributedDataShared part_exchange; StackEntry stack_loc{}; @@ -128,7 +128,7 @@ namespace shamrock { histogram_extract[current_pid] = 0; if (!pdat.is_empty()) { - sycl::host_accessor nid{new_pid.get(current_pid), sycl::read_only}; + std::vector nid = new_pid.get(current_pid).copy_to_stdvec(); if (false) { @@ -222,9 +222,10 @@ namespace shamrock { u32 ipos = sched.pdl_old().get_field_idx(position_field); - DistributedData> new_pid = compute_new_pid(sptree, ipos); + DistributedData> new_pid = compute_new_pid(sptree, ipos); - DistributedDataShared part_exchange = extract_elements(new_pid); + DistributedDataShared part_exchange + = extract_elements(std::move(new_pid)); part_exchange.for_each([](u64 sender, u64 receiver, PatchDataLayer &pdat) { shamlog_debug_ln("ReattributeDataUtility", sender, receiver, pdat.get_obj_cnt()); diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index 8449755de..78b266472 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -20,14 +20,16 @@ #include "shambase/memory.hpp" #include "shambase/stacktrace.hpp" -#include "shamrock/legacy/patch/utility/patch_field.hpp" -#include "shamrock/legacy/patch/utility/patch_reduc_tree.hpp" +#include "shambackends/BufferMirror.hpp" +#include "shambackends/DeviceBuffer.hpp" +#include "shambackends/kernel_call.hpp" #include "shamrock/patch/PatchField.hpp" #include "shamrock/scheduler/PatchScheduler.hpp" #include "shamrock/scheduler/PatchTree.hpp" #include "shamsys/legacy/log.hpp" #include "shamsys/legacy/sycl_handler.hpp" #include +#include #include #include @@ -38,30 +40,32 @@ class SerialPatchTree { using PatchTree = shamrock::scheduler::PatchTree; - // TODO use unique pointer instead u32 root_count = 0; - std::unique_ptr> serial_tree_buf; - std::unique_ptr> linked_patch_ids_buf; + std::optional> serial_tree_buf; + std::optional> linked_patch_ids_buf; inline void attach_buf() { - if (bool(serial_tree_buf)) + if (serial_tree_buf.has_value()) throw shambase::make_except_with_loc( "serial_tree_buf is already allocated"); - if (bool(linked_patch_ids_buf)) + if (linked_patch_ids_buf.has_value()) throw shambase::make_except_with_loc( "linked_patch_ids_buf is already allocated"); - serial_tree_buf - = std::make_unique>(serial_tree.data(), serial_tree.size()); - linked_patch_ids_buf - = std::make_unique>(linked_patch_ids.data(), linked_patch_ids.size()); + auto dev_sched = shamsys::instance::get_compute_scheduler_ptr(); + + serial_tree_buf.emplace(serial_tree.size(), dev_sched); + serial_tree_buf->copy_from_stdvec(serial_tree); + + linked_patch_ids_buf.emplace(linked_patch_ids.size(), dev_sched); + linked_patch_ids_buf->copy_from_stdvec(linked_patch_ids); } inline void detach_buf() { - if (!bool(serial_tree_buf)) + if (!serial_tree_buf.has_value()) throw shambase::make_except_with_loc( "serial_tree_buf wasn't allocated"); - if (!bool(linked_patch_ids_buf)) + if (!linked_patch_ids_buf.has_value()) throw shambase::make_except_with_loc( "linked_patch_ids_buf wasn't allocated"); @@ -149,8 +153,8 @@ class SerialPatchTree { std::function found_case) { StackEntry stack_loc{false}; - sycl::host_accessor tree{shambase::get_check_ref(serial_tree_buf), sycl::read_only}; - sycl::host_accessor lpid{shambase::get_check_ref(linked_patch_ids_buf), sycl::read_only}; + auto tree = serial_tree_buf->template mirror_to(); + auto lpid = linked_patch_ids_buf->template mirror_to(); host_for_each_leafs_internal(interact_cd, found_case, tree, lpid); } @@ -174,109 +178,18 @@ class SerialPatchTree { sched.patch_tree, sched.get_patch_transform()); } - template - inline PatchFieldReduction reduce_field( - sycl::queue &queue, PatchScheduler &sched, legacy::PatchField &pfield) { - - PatchFieldReduction predfield; - - std::cout << "resize to " << get_element_count() << std::endl; - predfield.tree_field.resize(get_element_count()); - - { - sycl::host_accessor lpid{*linked_patch_ids_buf, sycl::read_only}; - - // init reduction - std::unordered_map &idp_to_gid = sched.patch_list.id_patch_to_global_idx; - for (u64 idx = 0; idx < get_element_count(); idx++) { - predfield.tree_field[idx] - = (lpid[idx] != u64_max) ? pfield.global_values[idp_to_gid[lpid[idx]]] : type(); - - // std::cout << " el " << idx << " " << predfield.tree_field[idx] << std::endl; - } - } - - // std::cout << "predfield.attach_buf();" << std::endl; - - predfield.attach_buf(); - - sycl::range<1> range{get_element_count()}; - - u32 end_loop = get_level_count(); - - for (u32 level = 0; level < end_loop; level++) { - - // { - // auto f = predfield.tree_field_buf->template - // get_access(); std::cout << "["; for (u64 idx = 0; idx < - // get_element_count() ; idx ++) { - // std::cout << f[idx] << ","; - // } - // std::cout << std::endl; - // } - - std::cout << "queue submit : " << level << " " << end_loop << " " << (level < end_loop) - << std::endl; - queue.submit([&](sycl::handler &cgh) { - auto tree - = this->serial_tree_buf->template get_access(cgh); - - auto f - = predfield.tree_field_buf->template get_access( - cgh); - - cgh.parallel_for(range, [=](sycl::item<1> item) { - u64 i = (u64) item.get_id(0); - - u64 idx_c0 = tree[i].childs_id0; - u64 idx_c1 = tree[i].childs_id1; - u64 idx_c2 = tree[i].childs_id2; - u64 idx_c3 = tree[i].childs_id3; - u64 idx_c4 = tree[i].childs_id4; - u64 idx_c5 = tree[i].childs_id5; - u64 idx_c6 = tree[i].childs_id6; - u64 idx_c7 = tree[i].childs_id7; - - if (idx_c0 != u64_max) { - f[i] = reduc_func::reduce( - f[idx_c0], - f[idx_c1], - f[idx_c2], - f[idx_c3], - f[idx_c4], - f[idx_c5], - f[idx_c6], - f[idx_c7]); - } - }); - }); - } - // { - // auto f = predfield.tree_field_buf->template get_access(); - // std::cout << "["; - // for (u64 idx = 0; idx < get_element_count() ; idx ++) { - // std::cout << f[idx] << ","; - // } - // std::cout << std::endl; - // } - - return predfield; - } - template inline shamrock::patch::PatchtreeField make_patch_tree_field( PatchScheduler &sched, - sycl::queue &queue, + sham::DeviceScheduler_ptr dev_sched, shamrock::patch::PatchField pfield, Func &&reducer) { shamrock::patch::PatchtreeField ptfield; - ptfield.allocate(get_element_count()); + ptfield.allocate(get_element_count(), dev_sched); { - sycl::host_accessor lpid{ - shambase::get_check_ref(linked_patch_ids_buf), sycl::read_only}; - sycl::host_accessor tree_field{ - shambase::get_check_ref(ptfield.internal_buf), sycl::write_only, sycl::no_init}; + auto lpid = linked_patch_ids_buf->template mirror_to(); + auto tree_field = ptfield.internal_buf->template mirror_to(); // init reduction std::unordered_map &idp_to_gid = sched.patch_list.id_patch_to_global_idx; @@ -285,18 +198,17 @@ class SerialPatchTree { } } - sycl::range<1> range{get_element_count()}; - u32 end_loop = get_level_count(); + auto &q = shambase::get_check_ref(dev_sched).get_queue(); + u32 end_loop = get_level_count(); + u32 elem_cnt = get_element_count(); for (u32 level = 0; level < end_loop; level++) { - queue.submit([&](sycl::handler &cgh) { - sycl::accessor tree{shambase::get_check_ref(serial_tree_buf), cgh, sycl::read_only}; - sycl::accessor f{ - shambase::get_check_ref(ptfield.internal_buf), cgh, sycl::read_write}; - - cgh.parallel_for(range, [=](sycl::item<1> item) { - u64 i = (u64) item.get_id(0); - + sham::kernel_call( + q, + sham::MultiRef{*serial_tree_buf}, + sham::MultiRef{*ptfield.internal_buf}, + elem_cnt, + [reducer](u32 i, const PtNode *tree, T *f) { std::array n = tree[i].childs_id; if (n[0] != u64_max) { @@ -304,7 +216,6 @@ class SerialPatchTree { f[n[0]], f[n[1]], f[n[2]], f[n[3]], f[n[4]], f[n[5]], f[n[6]], f[n[7]]); } }); - }); } return ptfield; } @@ -332,41 +243,40 @@ class SerialPatchTree { } } - sycl::buffer compute_patch_owner( + sham::DeviceBuffer compute_patch_owner( sham::DeviceScheduler_ptr dev_sched, sham::DeviceBuffer &position_buffer, u32 len); }; template -sycl::buffer SerialPatchTree::compute_patch_owner( +sham::DeviceBuffer SerialPatchTree::compute_patch_owner( sham::DeviceScheduler_ptr dev_sched, sham::DeviceBuffer &position_buffer, u32 len) { - sycl::buffer new_owned_id(len); + sham::DeviceBuffer new_owned_id(len, dev_sched); using namespace shamrock::patch; - sycl::buffer roots = shamalgs::vec_to_buf(roots_ids); - - auto &q = dev_sched->get_queue(); - - sham::EventList depends_list; - auto pos = position_buffer.get_read_access(depends_list); - - auto e = q.submit(depends_list, [&](sycl::handler &cgh) { - sycl::accessor tnode{shambase::get_check_ref(serial_tree_buf), cgh, sycl::read_only}; - sycl::accessor linked_node_id{ - shambase::get_check_ref(linked_patch_ids_buf), cgh, sycl::read_only}; - sycl::accessor roots_id{roots, cgh, sycl::read_only}; - sycl::accessor new_id{new_owned_id, cgh, sycl::write_only, sycl::no_init}; - - u32 root_cnt = roots_id.size(); - auto max_lev = get_level_count(); - - using PtNode = shamrock::scheduler::SerialPatchNode; - - cgh.parallel_for(sycl::range(len), [=](sycl::item<1> item) { - u32 i = (u32) item.get_id(0); - + sham::DeviceBuffer roots(roots_ids.size(), dev_sched); + roots.copy_from_stdvec(roots_ids); + + auto &q = dev_sched->get_queue(); + u32 root_cnt = roots_ids.size(); + auto max_lev = get_level_count(); + + using PtNode = shamrock::scheduler::SerialPatchNode; + + sham::kernel_call( + q, + sham::MultiRef{position_buffer, *serial_tree_buf, *linked_patch_ids_buf, roots}, + sham::MultiRef{new_owned_id}, + len, + [root_cnt, max_lev]( + u32 i, + const vec *pos, + const PtNode *tnode, + const u64 *linked_node_id, + const u64 *roots_id, + u64 *new_id) { auto xyz = pos[i]; u64 current_node = 0; @@ -445,25 +355,8 @@ sycl::buffer SerialPatchTree::compute_patch_owner( } } - if constexpr (false) { - PtNode cur_node = tnode[current_node]; - if (xyz[0] == 0 && xyz[1] == 0 && xyz[2] == 0) { - logger::raw( - shambase::format( - "{:5} ({}) -> {} [{} {}]\n", - i, - Patch::is_in_patch_converted(xyz, cur_node.box_min, cur_node.box_max), - xyz, - cur_node.box_min, - cur_node.box_max)); - } - } - new_id[i] = result_node; }); - }); - - position_buffer.complete_event_state(e); return new_owned_id; } From 8e0470ee6aaef400d790db55b01212ada9168df0 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Fri, 8 May 2026 14:32:25 +0000 Subject: [PATCH 02/10] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- .../include/shamrock/scheduler/SerialPatchTree.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index 78b266472..d388f36e2 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -198,9 +198,9 @@ class SerialPatchTree { } } - auto &q = shambase::get_check_ref(dev_sched).get_queue(); - u32 end_loop = get_level_count(); - u32 elem_cnt = get_element_count(); + auto &q = shambase::get_check_ref(dev_sched).get_queue(); + u32 end_loop = get_level_count(); + u32 elem_cnt = get_element_count(); for (u32 level = 0; level < end_loop; level++) { sham::kernel_call( @@ -259,9 +259,9 @@ sham::DeviceBuffer SerialPatchTree::compute_patch_owner( sham::DeviceBuffer roots(roots_ids.size(), dev_sched); roots.copy_from_stdvec(roots_ids); - auto &q = dev_sched->get_queue(); - u32 root_cnt = roots_ids.size(); - auto max_lev = get_level_count(); + auto &q = dev_sched->get_queue(); + u32 root_cnt = roots_ids.size(); + auto max_lev = get_level_count(); using PtNode = shamrock::scheduler::SerialPatchNode; From f44883e202e22d9ab94b4134870b42411b2b70d9 Mon Sep 17 00:00:00 2001 From: Thomas Applencourt Date: Fri, 8 May 2026 16:37:37 +0200 Subject: [PATCH 03/10] Update src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index d388f36e2..447dd288b 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -153,8 +153,8 @@ class SerialPatchTree { std::function found_case) { StackEntry stack_loc{false}; - auto tree = serial_tree_buf->template mirror_to(); - auto lpid = linked_patch_ids_buf->template mirror_to(); + auto tree = shambase::get_check_ref(serial_tree_buf).template mirror_to(); + auto lpid = shambase::get_check_ref(linked_patch_ids_buf).template mirror_to(); host_for_each_leafs_internal(interact_cd, found_case, tree, lpid); } From 663d24b45248691eda3f99af8e6e4541b9682841 Mon Sep 17 00:00:00 2001 From: Thomas Applencourt Date: Fri, 8 May 2026 16:37:49 +0200 Subject: [PATCH 04/10] Update src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index 447dd288b..f7dc1dc45 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -188,8 +188,8 @@ class SerialPatchTree { ptfield.allocate(get_element_count(), dev_sched); { - auto lpid = linked_patch_ids_buf->template mirror_to(); - auto tree_field = ptfield.internal_buf->template mirror_to(); + auto lpid = shambase::get_check_ref(linked_patch_ids_buf).template mirror_to(); + auto tree_field = shambase::get_check_ref(ptfield.internal_buf).template mirror_to(); // init reduction std::unordered_map &idp_to_gid = sched.patch_list.id_patch_to_global_idx; From 039c48efd5848998aa0406114a791556fd93d697 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?David--Cl=C3=A9ris=20Timoth=C3=A9e?= Date: Fri, 8 May 2026 23:33:31 +0200 Subject: [PATCH 05/10] Update src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index f7dc1dc45..d3041b755 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -205,8 +205,8 @@ class SerialPatchTree { for (u32 level = 0; level < end_loop; level++) { sham::kernel_call( q, - sham::MultiRef{*serial_tree_buf}, - sham::MultiRef{*ptfield.internal_buf}, + sham::MultiRef{shambase::get_check_ref(serial_tree_buf)}, + sham::MultiRef{shambase::get_check_ref(ptfield.internal_buf)}, elem_cnt, [reducer](u32 i, const PtNode *tree, T *f) { std::array n = tree[i].childs_id; From 8b5fdbe79d79fd065b913b7ed9faebd5de156e18 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?David--Cl=C3=A9ris=20Timoth=C3=A9e?= Date: Fri, 8 May 2026 23:33:46 +0200 Subject: [PATCH 06/10] Update src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index d3041b755..c7d4aa030 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -259,7 +259,7 @@ sham::DeviceBuffer SerialPatchTree::compute_patch_owner( sham::DeviceBuffer roots(roots_ids.size(), dev_sched); roots.copy_from_stdvec(roots_ids); - auto &q = dev_sched->get_queue(); + auto &q = shambase::get_check_ref(dev_sched).get_queue(); u32 root_cnt = roots_ids.size(); auto max_lev = get_level_count(); From ff85c067f622b6a5e2a56536f09f0bf8fca3ff36 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?David--Cl=C3=A9ris=20Timoth=C3=A9e?= Date: Fri, 8 May 2026 23:33:56 +0200 Subject: [PATCH 07/10] Update src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index c7d4aa030..97b819fe1 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -267,7 +267,7 @@ sham::DeviceBuffer SerialPatchTree::compute_patch_owner( sham::kernel_call( q, - sham::MultiRef{position_buffer, *serial_tree_buf, *linked_patch_ids_buf, roots}, + sham::MultiRef{position_buffer, shambase::get_check_ref(serial_tree_buf), shambase::get_check_ref(linked_patch_ids_buf), roots}, sham::MultiRef{new_owned_id}, len, [root_cnt, max_lev]( From 170dba25302d7e7d27533931443518bc9dc23bd6 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Fri, 8 May 2026 21:35:29 +0000 Subject: [PATCH 08/10] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- .../include/shamrock/scheduler/SerialPatchTree.hpp | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index 97b819fe1..f47863caa 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -188,8 +188,10 @@ class SerialPatchTree { ptfield.allocate(get_element_count(), dev_sched); { - auto lpid = shambase::get_check_ref(linked_patch_ids_buf).template mirror_to(); - auto tree_field = shambase::get_check_ref(ptfield.internal_buf).template mirror_to(); + auto lpid + = shambase::get_check_ref(linked_patch_ids_buf).template mirror_to(); + auto tree_field + = shambase::get_check_ref(ptfield.internal_buf).template mirror_to(); // init reduction std::unordered_map &idp_to_gid = sched.patch_list.id_patch_to_global_idx; @@ -259,7 +261,7 @@ sham::DeviceBuffer SerialPatchTree::compute_patch_owner( sham::DeviceBuffer roots(roots_ids.size(), dev_sched); roots.copy_from_stdvec(roots_ids); - auto &q = shambase::get_check_ref(dev_sched).get_queue(); + auto &q = shambase::get_check_ref(dev_sched).get_queue(); u32 root_cnt = roots_ids.size(); auto max_lev = get_level_count(); @@ -267,7 +269,11 @@ sham::DeviceBuffer SerialPatchTree::compute_patch_owner( sham::kernel_call( q, - sham::MultiRef{position_buffer, shambase::get_check_ref(serial_tree_buf), shambase::get_check_ref(linked_patch_ids_buf), roots}, + sham::MultiRef{ + position_buffer, + shambase::get_check_ref(serial_tree_buf), + shambase::get_check_ref(linked_patch_ids_buf), + roots}, sham::MultiRef{new_owned_id}, len, [root_cnt, max_lev]( From f1e29d2df3c9c2d29b5065e37eb22530a7c023e6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?David--Cl=C3=A9ris=20Timoth=C3=A9e?= Date: Fri, 29 May 2026 19:47:24 +0200 Subject: [PATCH 09/10] Update src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp --- src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index f47863caa..203a79a88 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -256,6 +256,11 @@ sham::DeviceBuffer SerialPatchTree::compute_patch_owner( sham::DeviceScheduler_ptr dev_sched, sham::DeviceBuffer &position_buffer, u32 len) { sham::DeviceBuffer new_owned_id(len, dev_sched); + + if(len == 0){ + return new_owned_id; + } + using namespace shamrock::patch; sham::DeviceBuffer roots(roots_ids.size(), dev_sched); From 9f9cff8253f7d97208d7e534b9a52fdb318bd660 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Timoth=C3=A9e=20David--Cl=C3=A9ris?= Date: Fri, 29 May 2026 20:13:40 +0200 Subject: [PATCH 10/10] whoops --- .../include/shamrock/scheduler/SerialPatchTree.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp index 240695900..c0b753828 100644 --- a/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp +++ b/src/shamrock/include/shamrock/scheduler/SerialPatchTree.hpp @@ -21,7 +21,6 @@ #include "shambase/memory.hpp" #include "shambase/stacktrace.hpp" #include "shamrock/legacy/patch/utility/patch_field.hpp" -#include "shamrock/legacy/patch/utility/patch_reduc_tree.hpp" #include "shamrock/patch/PatchField.hpp" #include "shamrock/scheduler/PatchScheduler.hpp" #include "shamrock/scheduler/PatchTree.hpp" @@ -215,7 +214,7 @@ class SerialPatchTree { f[n[0]], f[n[1]], f[n[2]], f[n[3]], f[n[4]], f[n[5]], f[n[6]], f[n[7]]); } }); - }); + }); } return ptfield; } @@ -271,9 +270,9 @@ sycl::buffer SerialPatchTree::compute_patch_owner( sycl::accessor new_id{new_owned_id, cgh, sycl::write_only, sycl::no_init}; u32 root_cnt = roots_id.size(); - auto max_lev = get_level_count(); + auto max_lev = get_level_count(); - using PtNode = shamrock::scheduler::SerialPatchNode; + using PtNode = shamrock::scheduler::SerialPatchNode; cgh.parallel_for(sycl::range(len), [=](sycl::item<1> item) { u32 i = (u32) item.get_id(0);