PatchDataFieldSpan.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 "shambase/SourceLocation.hpp"
#include "shambase/aliases_int.hpp"
#include "shambase/format.hpp"
#include "shambase/stacktrace.hpp"
#include "shambase/string.hpp"
#include "shambackends/DeviceBuffer.hpp"
 
// forward declare PatchDataField
template<class T>
class PatchDataField;
 
namespace shamrock {
 
    namespace details {
 
        template<class T>
        struct PatchDataFieldSpan_access_rw_dyn_nvar {
            T *ptr;   
            u32 nvar; 
 
            T &operator()(u32 idx, u32 offset) const { return ptr[idx * nvar + offset]; }
 
            T &operator[](u32 idx) const { return ptr[idx]; }
        };
 
        template<class T>
        struct PatchDataFieldSpan_access_ro_dyn_nvar {
            const T *ptr; 
            u32 nvar;     
 
            const T &operator()(u32 idx, u32 offset) const { return ptr[idx * nvar + offset]; }
 
            const T &operator[](u32 idx) const { return ptr[idx]; }
        };
 
        template<class T, u32 nvar>
        struct PatchDataFieldSpan_access_rw_static_nvar {
            T *ptr; 
 
            T &operator()(u32 idx, u32 offset) const { return ptr[idx * nvar + offset]; }
 
            template<typename Dummy = void, typename = std::enable_if_t<nvar == 1, Dummy>>
            T &operator()(u32 idx) const {
                return ptr[idx];
            }
        };
 
        template<class T, u32 nvar>
        struct PatchDataFieldSpan_access_ro_static_nvar {
            const T *ptr; 
 
            const T &operator()(u32 idx, u32 offset) const { return ptr[idx * nvar + offset]; }
 
            template<typename Dummy = void, typename = std::enable_if_t<nvar == 1, Dummy>>
            const T &operator()(u32 idx) const {
                return ptr[idx];
            }
        };
    } // namespace details
 
    template<class T>
    using pdat_span_rw_dyn = details::PatchDataFieldSpan_access_rw_dyn_nvar<T>;
 
    template<class T>
    using pdat_span_ro_dyn = details::PatchDataFieldSpan_access_ro_dyn_nvar<T>;
 
    template<class T, u32 nvar>
    using pdat_span_rw = details::PatchDataFieldSpan_access_rw_static_nvar<T, nvar>;
 
    template<class T, u32 nvar>
    using pdat_span_ro = details::PatchDataFieldSpan_access_ro_static_nvar<T, nvar>;
 
    inline constexpr u32 dynamic_nvar = u32_max;
 
    inline constexpr bool access_t_pointer = true;
    inline constexpr bool access_t_span    = !access_t_pointer;
 
    template<class T, u32 nvar = dynamic_nvar, bool pointer_access = access_t_span>
    class PatchDataFieldSpan {
        public:
        inline static constexpr bool is_nvar_dynamic() { return nvar == dynamic_nvar; }
 
        inline static constexpr bool is_nvar_static() { return nvar != dynamic_nvar; }
 
        inline static constexpr bool is_pointer_access() {
            return pointer_access == access_t_pointer;
        }
        inline static constexpr bool is_span_access() { return pointer_access == access_t_span; }
 
        PatchDataFieldSpan(
            PatchDataField<T> &field_ref,
            u32 start,
            u32 count,
            SourceLocation loc = SourceLocation{})
            : field_ref(field_ref), start(start), count(count) {
 
            StackEntry stack_loc{};
 
            // ensure that the underlying USM pointer can be accessed
            if (field_ref.get_buf().is_empty()) {
                shambase::throw_with_loc<std::invalid_argument>(
                    "PatchDataFieldSpan can not be binded to empty buffer", loc);
            }
 
            if (is_nvar_static()) {
                if (field_ref.get_nvar() != nvar) {
                    shambase::throw_with_loc<std::invalid_argument>(
                        shambase::format(
                            "You are trying to bind a PatchDataFieldSpan with static nvar={} to a "
                            "PatchDataField with nvar={}",
                            nvar,
                            field_ref.get_nvar()),
                        loc);
                }
            }
 
            if (start + count > field_ref.get_obj_cnt()) {
                shambase::throw_with_loc<std::invalid_argument>(
                    shambase::format(
                        "PatchDataFieldSpan out of bounds: {} + {} > {}",
                        start,
                        count,
                        field_ref.get_obj_cnt()),
                    loc);
            }
        }
 
        template<
            typename Dummy = void,
            typename       = std::enable_if_t<is_nvar_dynamic() && is_span_access(), Dummy>>
        inline auto get_read_access(
            sham::EventList &depends_list, SourceLocation src_loc = SourceLocation{}) const
            -> details::PatchDataFieldSpan_access_ro_dyn_nvar<T> {
            StackEntry stack_loc{};
            return details::PatchDataFieldSpan_access_ro_dyn_nvar<T>{
                get_buf().get_read_access(depends_list, std::move(src_loc))
                    + start * field_ref.get_nvar(),
                field_ref.get_nvar()};
        }
 
        template<
            typename Dummy = void,
            typename       = std::enable_if_t<is_nvar_dynamic() && is_span_access(), Dummy>>
        inline auto get_write_access(
            sham::EventList &depends_list, SourceLocation src_loc = SourceLocation{})
            -> details::PatchDataFieldSpan_access_rw_dyn_nvar<T> {
            StackEntry stack_loc{};
            return details::PatchDataFieldSpan_access_rw_dyn_nvar<T>{
                get_buf().get_write_access(depends_list, std::move(src_loc))
                    + start * field_ref.get_nvar(),
                field_ref.get_nvar()};
        }
 
        template<
            typename Dummy = void,
            typename       = std::enable_if_t<is_nvar_static() && is_span_access(), Dummy>>
        inline auto get_read_access(
            sham::EventList &depends_list, SourceLocation src_loc = SourceLocation{}) const
            -> details::PatchDataFieldSpan_access_ro_static_nvar<T, nvar> {
            StackEntry stack_loc{};
            return details::PatchDataFieldSpan_access_ro_static_nvar<T, nvar>{
                get_buf().get_read_access(depends_list, std::move(src_loc))
                + start * field_ref.get_nvar()};
        }
 
        template<
            typename Dummy = void,
            typename       = std::enable_if_t<is_nvar_static() && is_span_access(), Dummy>>
        inline auto get_write_access(
            sham::EventList &depends_list, SourceLocation src_loc = SourceLocation{})
            -> details::PatchDataFieldSpan_access_rw_static_nvar<T, nvar> {
            StackEntry stack_loc{};
            return details::PatchDataFieldSpan_access_rw_static_nvar<T, nvar>{
                get_buf().get_write_access(depends_list, std::move(src_loc))
                + start * field_ref.get_nvar()};
        }
 
        template<typename Dummy = void, typename = std::enable_if_t<is_pointer_access(), Dummy>>
        inline auto get_read_access(
            sham::EventList &depends_list, SourceLocation src_loc = SourceLocation{}) const
            -> const T * {
            StackEntry stack_loc{};
            return {
                get_buf().get_read_access(depends_list, std::move(src_loc))
                + start * field_ref.get_nvar()};
        }
 
        template<typename Dummy = void, typename = std::enable_if_t<is_pointer_access(), Dummy>>
        inline auto get_write_access(
            sham::EventList &depends_list, SourceLocation src_loc = SourceLocation{}) -> T * {
            StackEntry stack_loc{};
            return {
                get_buf().get_write_access(depends_list, std::move(src_loc))
                + start * field_ref.get_nvar()};
        }
 
        inline void complete_event_state(sycl::event e) const {
            StackEntry stack_loc{};
            get_buf().complete_event_state(e);
        }
 
        PatchDataField<T> &field_ref;
 
        u32 start;
 
        u32 count;
 
        private:
        inline sham::DeviceBuffer<T> &get_buf() { return field_ref.get_buf(); }
 
        inline const sham::DeviceBuffer<T> &get_buf() const { return field_ref.get_buf(); }
    };
 
    template<class T>
    using PatchDataFieldSpanPointer = PatchDataFieldSpan<T, dynamic_nvar, access_t_pointer>;
 
} // namespace shamrock