DeviceBuffer.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/assert.hpp"
#include "shambase/memory.hpp"
#include "shambase/narrowing.hpp"
#include "shambase/type_traits.hpp"
#include "shambackends/DeviceScheduler.hpp"
#include "shambackends/USMPtrHolder.hpp"
#include "shambackends/details/BufferEventHandler.hpp"
#include "shambackends/details/memoryHandle.hpp"
#include "shambackends/sycl_utils.hpp"
#include <memory>
 
namespace sham {
 
    template<class T>
    inline sham::EventList safe_copy(
        sham::DeviceQueue &q, sham::EventList &depends_list, const T *src, T *dest, size_t count) {
        __shamrock_stack_entry();
        u64 max_copy_len = (1 << 30) / sizeof(T); // 1GB to avoid memcpy above i32_max bytes
 
        sham::EventList events;
 
        for (size_t i = 0; i < count; i += max_copy_len) {
            i32 copy_len
                = shambase::narrow_or_throw<i32>(std::min<size_t>(max_copy_len, count - i));
            sycl::event e = q.submit(depends_list, [&](sycl::handler &cgh) {
                cgh.copy(src + i, dest + i, copy_len);
            });
            events.add_event(e);
        }
 
        return events;
    }
 
    template<class T>
    inline sham::EventList safe_fill(
        sham::DeviceQueue &q, sham::EventList &depends_list, T *dest, size_t count, T value) {
        __shamrock_stack_entry();
        u64 max_copy_len = (1 << 30); // 1G elements, this garanteee indexing below i32_max
 
        sham::EventList events;
 
        for (size_t i = 0; i < count; i += max_copy_len) {
            i32 copy_len
                = shambase::narrow_or_throw<i32>(std::min<size_t>(max_copy_len, count - i));
            sycl::event e = q.submit(depends_list, [&](sycl::handler &cgh) {
                cgh.parallel_for(sycl::range<1>(copy_len), [dest, i, value](sycl::item<1> gid) {
                    dest[i + gid.get_linear_id()] = value;
                });
            });
            events.add_event(e);
        }
 
        return events;
    }
 
    template<class T, class Fct>
    inline sham::EventList safe_fill_lambda(
        sham::DeviceQueue &q, sham::EventList &depends_list, T *dest, size_t count, Fct &&fct) {
        __shamrock_stack_entry();
        u64 max_copy_len = (1 << 30); // 1G elements, this garanteee indexing below i32_max
 
        sham::EventList events;
 
        for (size_t i = 0; i < count; i += max_copy_len) {
            i32 copy_len
                = shambase::narrow_or_throw<i32>(std::min<size_t>(max_copy_len, count - i));
            sycl::event e = q.submit(depends_list, [&](sycl::handler &cgh) {
                cgh.parallel_for(sycl::range<1>(copy_len), [dest, i, fct](sycl::item<1> gid) {
                    dest[i + gid.get_linear_id()] = fct(i + gid.get_linear_id());
                });
            });
            events.add_event(e);
        }
 
        return events;
    }
 
    template<class T, USMKindTarget target = host, USMKindTarget orgin_target = device>
    class BufferMirror;
 
    template<class T, USMKindTarget target = device>
    class DeviceBuffer {
 
        static size_t upgrade_multiple(size_t sz, size_t mult) {
            if (sz % mult)
                return sz + (mult - sz % mult);
            return sz;
        };
 
        public:
        static std::optional<size_t> get_alignment(const DeviceScheduler_ptr &dev_sched) {
            return upgrade_multiple(
                alignof(T),
                std::max(
                    shambase::get_check_ref(dev_sched)
                        .get_queue()
                        .get_device_prop()
                        .mem_base_addr_align,
                    8_u32));
        }
 
        static size_t alloc_request_size_fct(size_t sz, const DeviceScheduler_ptr &dev_sched) {
            size_t ret = sz * sizeof(T);
 
            auto align = get_alignment(dev_sched);
            if (align) {
                ret = upgrade_multiple(ret, *align);
                SHAM_ASSERT(ret % *align == 0);
            }
 
            SHAM_ASSERT((sz == 0) ? (ret == 0) : (ret >= sz * sizeof(T)));
 
            return ret;
        }
 
        DeviceBuffer(size_t sz, USMPtrHolder<target> &&_hold)
            : hold(std::forward<USMPtrHolder<target>>(_hold)), size(sz),
              events_hndl(std::make_unique<details::BufferEventHandler>()) {}
 
        DeviceBuffer(size_t sz, DeviceScheduler_ptr dev_sched)
            : DeviceBuffer(
                  sz,
                  details::create_usm_ptr<target>(
                      alloc_request_size_fct(sz, dev_sched), dev_sched, get_alignment(dev_sched))) {
        }
 
        DeviceBuffer(sycl::buffer<T> &syclbuf, std::shared_ptr<DeviceScheduler> dev_sched)
            : DeviceBuffer(syclbuf.size(), dev_sched) {
            copy_from_sycl_buffer(syclbuf);
        }
 
        DeviceBuffer(
            sycl::buffer<T> &syclbuf, size_t sz, std::shared_ptr<DeviceScheduler> dev_sched)
            : DeviceBuffer(sz, dev_sched) {
            copy_from_sycl_buffer(syclbuf, sz);
        }
 
        DeviceBuffer(sycl::buffer<T> &&syclbuf, std::shared_ptr<DeviceScheduler> dev_sched)
            : DeviceBuffer(syclbuf, dev_sched) {}
 
        DeviceBuffer(
            sycl::buffer<T> &&syclbuf, size_t sz, std::shared_ptr<DeviceScheduler> dev_sched)
            : DeviceBuffer(syclbuf, sz, dev_sched) {}
 
        DeviceBuffer(const DeviceBuffer &other) = delete;
 
        DeviceBuffer &operator=(const DeviceBuffer &other) = delete;
 
        DeviceBuffer(DeviceBuffer &&other) noexcept
            : hold(std::move(other.hold)), size(other.size),
              events_hndl(std::move(other.events_hndl)) {}
 
        DeviceBuffer &operator=(DeviceBuffer &&other) noexcept {
            std::swap(hold, other.hold);
            std::swap(events_hndl, other.events_hndl);
            size = other.size;
            return *this;
        }
 
        ~DeviceBuffer() {
            if (!bool(events_hndl)) {
                // If this is not allocated it must be a moved object
                if (hold.get_raw_ptr() != nullptr) {
                    shambase::throw_with_loc<std::runtime_error>(
                        "you have an event handler but not pointer, like how ???");
                }
                return;
            }
            // This object is empty, it was probably moved
            if (hold.get_raw_ptr() == nullptr && events_hndl->is_empty()) {
                return;
            }
            // give the ptr holder and event handler to the memory handler
            details::release_usm_ptr(std::move(hold), shambase::extract_pointer(events_hndl));
        }
 
        // Event handling
 
        [[nodiscard]] inline const T *get_read_access(
            sham::EventList &depends_list, SourceLocation src_loc = SourceLocation{}) const {
            shambase::get_check_ref(events_hndl).read_access(depends_list, src_loc);
            return hold.template ptr_cast<T>();
        }
 
        [[nodiscard]] inline T *get_write_access(
            sham::EventList &depends_list, SourceLocation src_loc = SourceLocation{}) {
            shambase::get_check_ref(events_hndl).write_access(depends_list, src_loc);
            return hold.template ptr_cast<T>();
        }
 
        void complete_event_state(sycl::event e) const {
            shambase::get_check_ref(events_hndl).complete_state(e);
        }
 
        void complete_event_state(const std::vector<sycl::event> &e) const {
            shambase::get_check_ref(events_hndl).complete_state(e);
        }
 
        void complete_event_state(sham::EventList &e) const {
            shambase::get_check_ref(events_hndl).complete_state(e);
        }
 
        void synchronize() const { shambase::get_check_ref(events_hndl).wait_all(); }
 
        // Event handling (End)
 
 
        // Queue / Scheduler getters
 
        [[nodiscard]] inline DeviceScheduler &get_dev_scheduler() const {
            return hold.get_dev_scheduler();
        }
 
        [[nodiscard]] inline std::shared_ptr<DeviceScheduler> &get_dev_scheduler_ptr() {
            return hold.get_dev_scheduler_ptr();
        }
 
        [[nodiscard]] inline const std::shared_ptr<DeviceScheduler> &get_dev_scheduler_ptr() const {
            return hold.get_dev_scheduler_ptr();
        }
 
        [[nodiscard]] inline DeviceQueue &get_queue() const {
            return hold.get_dev_scheduler().get_queue();
        }
 
        // Queue / Scheduler getters (END)
 
 
        // Size getters
 
        [[nodiscard]] inline size_t get_size() const { return size; }
 
        [[nodiscard]] inline size_t get_mem_usage() const { return hold.get_bytesize(); }
 
        [[nodiscard]] inline bool is_empty() const { return size == 0; }
 
        // Size getters (END)
 
 
        // Copy fcts
 
        [[nodiscard]] inline std::vector<T> copy_to_stdvec() const {
            std::vector<T> ret(size);
 
            if (size > 0) {
                sham::EventList depends_list;
                const T *ptr = get_read_access(depends_list);
 
                sham::EventList events
                    = safe_copy(get_queue(), depends_list, ptr, ret.data(), size);
 
                events.wait_and_throw();
                complete_event_state(sycl::event{});
            }
 
            return ret;
        }
 
        [[nodiscard]] inline std::vector<T> copy_to_stdvec_idx_range(
            size_t begin, size_t end) const {
 
            if (end > size) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_to_stdvec_idx_range: end > size\n  end = {},\n  size = {}", end, size));
            }
 
            if (begin > end) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_to_stdvec_idx_range: begin >= end\n  begin = {},\n  end = {}",
                    begin,
                    end));
            }
 
            u32 size_cp = end - begin;
            std::vector<T> ret(size_cp);
 
            if (size_cp > 0) {
                sham::EventList depends_list;
                const T *ptr = get_read_access(depends_list);
 
                sham::EventList events
                    = safe_copy(get_queue(), depends_list, ptr + begin, ret.data(), size_cp);
 
                events.wait_and_throw();
                complete_event_state(sycl::event{});
            }
 
            return ret;
        }
 
        template<USMKindTarget dest_target>
        inline void copy_range_offset(
            size_t begin,
            size_t end,
            sham::DeviceBuffer<T, dest_target> &dest,
            size_t dest_offset) const {
 
            if (begin > end) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_range_offset: begin > end\n  begin = {},\n  end = {}", begin, end));
            }
 
            if (end > get_size()) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_range_offset: end index is out of bounds\n  end = {},\n  source buffer "
                    "size = {}",
                    end,
                    get_size()));
            }
 
            if (dest_offset > dest.get_size()) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_range_offset: dest_offset > dest.get_size()\n  dest_offset = {},\n  "
                    "dest.get_size() = {}",
                    dest_offset,
                    dest.get_size()));
            }
 
            if (end - begin > (dest.get_size() - dest_offset)) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_range_offset: end - begin > dest.get_size() - dest_offset\n  end - begin "
                    "= {},\n  "
                    "dest.get_size() - dest_offset = {},\n  dest_offset = {}",
                    end - begin,
                    dest.get_size() - dest_offset,
                    dest_offset));
            }
 
            if (static_cast<const void *>(this) == static_cast<const void *>(&dest)) {
                shambase::throw_with_loc<std::invalid_argument>(
                    "the source and destination buffers must not be the same");
            }
 
            if (begin == end) {
                return;
            }
 
            size_t len = end - begin;
 
            sham::EventList depends_list;
            const T *ptr_src = get_read_access(depends_list) + begin;
            T *ptr_dest      = dest.get_write_access(depends_list) + dest_offset;
 
            sham::EventList events = safe_copy(get_queue(), depends_list, ptr_src, ptr_dest, len);
 
            complete_event_state(events);
            dest.complete_event_state(events);
        }
 
        template<USMKindTarget dest_target>
        inline void copy_range(
            size_t begin, size_t end, sham::DeviceBuffer<T, dest_target> &dest) const {
 
            copy_range_offset(begin, end, dest, 0);
        }
 
        inline void copy_from_stdvec(const std::vector<T> &vec) {
 
            if (size != vec.size()) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_from_stdvec: size mismatch\n  size = {},\n  vec.size() = {}",
                    size,
                    vec.size()));
            }
 
            if (size > 0) {
                sham::EventList depends_list;
                T *ptr = get_write_access(depends_list);
 
                sham::EventList events
                    = safe_copy(get_queue(), depends_list, vec.data(), ptr, size);
 
                events.wait_and_throw();
                complete_event_state(sycl::event{});
            }
        }
 
        inline void copy_from_stdvec(const std::vector<T> &vec, size_t sz) {
 
            if (sz > vec.size() || sz > size) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_from_stdvec: size mismatch (sz > vec.size() || sz > size)\n  size = "
                    "{},\n  vec.size() = {},\n  sz = {}",
                    size,
                    vec.size(),
                    sz));
            }
 
            if (sz > 0) {
                sham::EventList depends_list;
                T *ptr = get_write_access(depends_list);
 
                sham::EventList events = safe_copy(get_queue(), depends_list, vec.data(), ptr, sz);
 
                events.wait_and_throw();
                complete_event_state(sycl::event{});
            }
        }
 
        [[nodiscard]] inline sycl::buffer<T> copy_to_sycl_buffer() const {
            sycl::buffer<T> ret(size);
 
            if (size > 0) {
 
                shambase::narrow_or_throw<i32>(size);
 
                sham::EventList depends_list;
                const T *ptr = get_read_access(depends_list);
 
                sycl::event e = get_queue().submit(depends_list, [&](sycl::handler &cgh) {
                    sycl::accessor acc(ret, cgh, sycl::write_only, sycl::no_init);
                    cgh.copy(ptr, acc);
                });
 
                complete_event_state(e);
            }
 
            return ret;
        }
 
        inline void copy_from_sycl_buffer(sycl::buffer<T> &buf) {
 
            if (size != buf.size()) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_from_sycl_buffer: size mismatch\n  size = {},\n  buf.size() = {}",
                    size,
                    buf.size()));
            }
 
            if (size > 0) {
                shambase::narrow_or_throw<i32>(size);
 
                sham::EventList depends_list;
                T *ptr = get_write_access(depends_list);
 
                sycl::event e = get_queue().submit(depends_list, [&](sycl::handler &cgh) {
                    sycl::accessor acc(buf, cgh, sycl::read_only);
                    cgh.copy(acc, ptr);
                });
 
                complete_event_state(e);
            }
        }
 
        inline void copy_from_sycl_buffer(sycl::buffer<T> &buf, size_t sz) {
 
            if (sz > buf.size() || sz > size) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_from_sycl_buffer: size mismatch (sz > buf.size() || sz > size)\n  size = "
                    "{},\n  buf.size() = {},\n  sz = {}",
                    size,
                    buf.size(),
                    sz));
            }
 
            if (sz > u32_max) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "copy_from_sycl_buffer: size mismatch (sz > u32_max)\n  sz = {}", sz));
            }
 
            if (size > 0) {
                sham::EventList depends_list;
                T *ptr = get_write_access(depends_list);
 
                sycl::event e = get_queue().submit(depends_list, [&](sycl::handler &cgh) {
                    sycl::accessor acc(buf, cgh, sycl::read_only);
 
                    shambase::parallel_for(
                        cgh, shambase::narrow_or_throw<i32>(sz), "copy field", [=](u32 gid) {
                            ptr[gid] = acc[gid];
                        });
                });
 
                complete_event_state(e);
            }
        }
 
        template<USMKindTarget new_target>
        [[nodiscard]] inline DeviceBuffer<T, new_target> copy_to() const {
            DeviceBuffer<T, new_target> ret(size, get_dev_scheduler_ptr());
 
            if (size > 0) {
                sham::EventList depends_list;
                const T *ptr_src = get_read_access(depends_list);
                T *ptr_dest      = ret.get_write_access(depends_list);
 
                sham::EventList events
                    = safe_copy(get_queue(), depends_list, ptr_src, ptr_dest, size);
 
                complete_event_state(events);
                ret.complete_event_state(events);
            }
 
            return ret;
        }
 
        template<USMKindTarget new_target>
        inline void copy_from(const DeviceBuffer<T, new_target> &other, size_t copy_size) {
 
            if (!(copy_size <= get_size() && copy_size <= other.get_size())) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "The size of the copy must be smaller than the size of the buffer involved\n  "
                    "copy_size: {}\n  get_size(): {}\n  other.get_size(): {}",
                    copy_size,
                    get_size(),
                    other.get_size()));
            }
 
            if (copy_size > 0) {
                sham::EventList depends_list;
                T *ptr_dest      = get_write_access(depends_list);
                const T *ptr_src = other.get_read_access(depends_list);
 
                sham::EventList events
                    = safe_copy(get_queue(), depends_list, ptr_src, ptr_dest, copy_size);
 
                complete_event_state(events);
                other.complete_event_state(events);
            }
        }
 
        template<USMKindTarget new_target>
        inline void copy_from(const DeviceBuffer<T, new_target> &other) {
 
            if (get_size() != other.get_size()) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "The other field must be of the same size\n  get_size = {},\n  other.get_size "
                    "= {}",
                    get_size(),
                    other.get_size()));
            }
 
            copy_from(other, get_size());
        }
 
        inline DeviceBuffer<T, target> copy() const { return copy_to<target>(); }
 
        template<USMKindTarget mirror_target>
        inline BufferMirror<T, mirror_target, target> mirror_to() {
            return BufferMirror<T, mirror_target, target>(*this);
        }
 
        // Copy fcts (END)
 
 
        // Filler fcts
 
        inline void fill(T value, std::array<size_t, 2> idx_range) {
 
            size_t start_index = idx_range[0];
            size_t idx_count   = idx_range[1] - start_index;
 
            if (!(start_index + idx_count <= get_size())) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "!(start_index + idx_count <= get_size())\n  start_index = {},\n  idx_count = "
                    "{},\n  get_size() = {}",
                    start_index,
                    idx_count,
                    get_size()));
            }
 
            shambase::narrow_or_throw<i32>(idx_count);
 
            sham::EventList depends_list;
            T *ptr = get_write_access(depends_list);
 
            sham::EventList events
                = safe_fill(get_queue(), depends_list, ptr + start_index, idx_count, value);
 
            complete_event_state(events);
        }
 
        inline void fill(T value, size_t idx_count) { fill(value, {0, idx_count}); }
 
        inline void fill(T value) { fill(value, get_size()); }
 
        template<class Fct>
        inline void fill_lambda(Fct &&fct) {
            if (get_size() == 0) {
                return;
            }
 
            sham::EventList depends_list;
            T *__restrict ptr = get_write_access(depends_list);
 
            sham::EventList events
                = safe_fill_lambda(get_queue(), depends_list, ptr, get_size(), fct);
 
            complete_event_state(events);
        }
 
        // Filler fcts (END)
 
 
        // Getter fcts
 
        T get_val_at_idx(size_t idx) const {
            T ret;
 
            if (idx >= size) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "get_val_at_idx: idx >= size\n  idx = {},\n  size = {}", idx, size));
            }
 
            sham::EventList depends_list;
            const T *ptr = get_read_access(depends_list);
 
            sycl::event e = get_queue().submit(depends_list, [&](sycl::handler &cgh) {
                cgh.copy(ptr + idx, &ret, 1);
            });
 
            e.wait_and_throw();
            complete_event_state(sycl::event{});
 
            return ret;
        }
 
        void set_val_at_idx(size_t idx, T val) {
 
            if (idx >= size) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "set_val_at_idx: idx >= size\n  idx = {},\n  size = {}", idx, size));
            }
 
            sham::EventList depends_list;
            T *ptr = get_write_access(depends_list);
 
            sycl::event e = get_queue().submit(depends_list, [&](sycl::handler &cgh) {
                cgh.copy(&val, ptr + idx, 1);
            });
 
            e.wait_and_throw();
            complete_event_state(sycl::event{});
        }
        // Getter fcts (END)
 
 
        // Size manipulation
 
        inline size_t get_max_alloc_size() const {
            auto &dev_prop = hold.get_dev_scheduler().get_queue().get_device_prop();
 
            if constexpr (target == device) {
                return dev_prop.max_mem_alloc_size_dev;
            } else if constexpr (target == host) {
                return dev_prop.max_mem_alloc_size_host;
            } else if constexpr (target == shared) {
                return sycl::min(dev_prop.max_mem_alloc_size_dev, dev_prop.max_mem_alloc_size_host);
            } else {
                static_assert(
                    shambase::always_false_v<decltype(target)>,
                    "get_max_alloc_size: invalid target");
            }
        }
 
        inline void resize(size_t new_size, bool keep_data = true) {
 
            auto dev_sched = hold.get_dev_scheduler_ptr();
 
            StackEntry __st{};
 
            if (alloc_request_size_fct(new_size, dev_sched) > hold.get_bytesize()) {
                // expand storage
 
                size_t max_alloc_size           = get_max_alloc_size();
                std::optional<size_t> alignment = get_alignment(dev_sched);
                size_t min_size_new_alloc       = alloc_request_size_fct(new_size, dev_sched);
                size_t wanted_size_new_alloc
                    = alloc_request_size_fct(new_size + new_size / 2, dev_sched);
                size_t max_possible_alloc = max_alloc_size;
 
                if (alignment) {
                    max_possible_alloc = max_possible_alloc - (max_possible_alloc % *alignment);
                }
 
                size_t new_storage_size = wanted_size_new_alloc;
                if (new_storage_size > max_alloc_size) {
                    new_storage_size = sycl::max(max_possible_alloc, min_size_new_alloc);
                }
 
                if (new_storage_size > max_alloc_size) {
                    shambase::throw_with_loc<std::runtime_error>(shambase::format(
                        "new_storage_size > max_alloc_size\n"
                        "  new_storage_size      = {}\n"
                        "  max_alloc_size  = {}\n"
                        "  min_size_new_alloc    = {}\n"
                        "  wanted_size_new_alloc = {}",
                        new_storage_size,
                        max_alloc_size,
                        min_size_new_alloc,
                        wanted_size_new_alloc));
                }
 
                DeviceBuffer new_buf(
                    new_size,
                    details::create_usm_ptr<target>(
                        new_storage_size, get_dev_scheduler_ptr(), get_alignment(dev_sched)));
 
                // copy data
                if (keep_data) {
                    new_buf.copy_from(*this, get_size());
                }
 
                // override old buffer
                std::swap(new_buf, *this);
 
            } else if (alloc_request_size_fct(new_size, dev_sched) < hold.get_bytesize() * 0.5) {
                // shrink storage
 
                size_t new_storage_size = alloc_request_size_fct(new_size, dev_sched);
 
                DeviceBuffer new_buf(
                    new_size,
                    details::create_usm_ptr<target>(
                        new_storage_size, get_dev_scheduler_ptr(), get_alignment(dev_sched)));
 
                // copy data
                if (keep_data) {
                    new_buf.copy_from(*this, new_size);
                }
 
                // override old buffer
                std::swap(new_buf, *this);
 
            } else {
                size = new_size;
                // no need to resize
            }
        }
 
        inline void resize_discard_data(size_t new_size) { resize(new_size, false); }
 
        inline void free_alloc() { resize_discard_data(0); }
 
        inline void expand(u32 add_sz) { resize(get_size() + add_sz); }
 
        inline void shrink(u32 sub_sz) {
            if (sub_sz > get_size()) {
                shambase::throw_with_loc<std::invalid_argument>(shambase::format(
                    "shrink called with sub_sz > get_size()\n  sub_sz: {}\n  get_size(): {}",
                    sub_sz,
                    get_size()));
            }
            resize(get_size() - sub_sz);
        }
 
        inline void append(const DeviceBuffer &other) {
            if (this == &other) {
                shambase::throw_with_loc<std::invalid_argument>("cannot append a buffer to itself");
            }
 
            u32 other_size = other.get_size();
            if (other_size == 0) {
                return; // early exit if the other buffer is empty
            }
            u32 old_size = get_size();
 
            // allocate space
            expand(other_size);
 
            sham::EventList depends_list;
            T *ptr             = get_write_access(depends_list);
            const T *other_ptr = other.get_read_access(depends_list);
 
            sham::EventList events
                = safe_copy(get_queue(), depends_list, other_ptr, ptr + old_size, other_size);
 
            complete_event_state(events);
            other.complete_event_state(events);
        }
 
        // Size manipulation (END)
 
        // I'm not sure if enabling this one is a good idea
        inline void reserve(size_t add_sz) {
#if false
            size_t old_sz = get_size();
            resize(old_sz + add_sz);
            size = old_sz;
#endif
        }
 
        private:
        USMPtrHolder<target> hold;
 
        size_t size = 0;
 
        std::unique_ptr<details::BufferEventHandler> events_hndl;
    };
 
} // namespace sham