CellIterator.hpp

Go to the documentation of this file.

// -------------------------------------------------------//
//
// SHAMROCK code for hydrodynamics
// Copyright (c) 2021-2026 Timothée David--Cléris <tim.shamrock@proton.me>
// 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
 
#include "shambackends/DeviceBuffer.hpp"
#include "shambackends/DeviceScheduler.hpp"
#include "shambackends/math.hpp"
#include "shamcomm/logs.hpp"
#include <utility>
 
namespace shamtree {
 
    struct CellIteratorAccessed {
        const u32 *sort_index_map;  
        const u32 *reduc_index_map; 
        const u32 *endrange; 
        u32 offset_leaf;     
 
        inline bool is_id_leaf(u32 id) const { return id >= offset_leaf; }
 
        template<class Functor_iter>
        inline void for_each_in_cell(const u32 &cell_id, Functor_iter &&func_it) const {
 
            bool is_leaf = is_id_leaf(cell_id);
 
            // internal cell id or leaf id (hence the sub internal_cell_count if leaf)
            u32 cbeg = (is_leaf) ? cell_id - offset_leaf : cell_id;
 
            // other end of the cell range (either ourself if leaf, or the endrange if internal)
            // this exclude the upper bound as the +1 must be made after the reordering
            u32 cend = ((is_leaf) ? cbeg : endrange[cbeg]);
 
            // tree cell index range
            uint c1 = sham::min(cbeg, cend);
            uint c2 = sham::max(cbeg, cend);
 
            u32 _startrange = reduc_index_map[c1];
            u32 _endrange   = reduc_index_map[c2 + 1]; // <--- this +1
 
            for (unsigned int id_s = _startrange; id_s < _endrange; id_s++) {
 
                // recover old index before morton sort
                uint id_b = sort_index_map[id_s];
 
                // iteration function
                func_it(id_b);
            }
        }
    };
 
    struct CellIterator {
        const sham::DeviceBuffer<u32> &buf_sort_index_map;  
        const sham::DeviceBuffer<u32> &buf_reduc_index_map; 
        const sham::DeviceBuffer<u32> &buf_endrange;        
        u32 offset_leaf; 
 
        using acc = CellIteratorAccessed;
 
        inline acc get_read_access(sham::EventList &deps) const {
            return acc{
                buf_sort_index_map.get_read_access(deps),
                buf_reduc_index_map.get_read_access(deps),
                buf_endrange.get_read_access(deps),
                offset_leaf};
        }
 
        inline void complete_event_state(sycl::event e) const {
            buf_sort_index_map.complete_event_state(e);
            buf_reduc_index_map.complete_event_state(e);
            buf_endrange.complete_event_state(e);
        }
    };
 
    struct CellIteratorHost {
        std::vector<u32> sort_index_map;  
        std::vector<u32> reduc_index_map; 
        std::vector<u32> endrange;        
        u32 offset_leaf; 
 
        using acc = CellIteratorAccessed;
 
        inline acc get_read_access() const {
            return acc{sort_index_map.data(), reduc_index_map.data(), endrange.data(), offset_leaf};
        }
    };
 
} // namespace shamtree