USMBufferInterop.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/math.hpp"
#include "shambackends/sycl.hpp"
#include "shambackends/sycl_utils.hpp"
#include "shamcomm/logs.hpp"
 
namespace sham {
    template<class T>
    inline std::vector<sycl::event> usmbuffer_memcpy(
        sycl::queue &queue, sycl::buffer<T> &src, T *dest, u64 count) {
 
        u64 offset                  = 0;
        u64 remains                 = count;
        constexpr u64 max_step_size = i32_max / 2;
 
        std::vector<sycl::event> ev_list{};
 
        while (offset < count) {
            u64 stepsize = sham::min(remains, max_step_size);
 
            ev_list.push_back(queue.submit([&, offset](sycl::handler &cgh) {
                sycl::accessor acc{src, cgh, sycl::read_only};
                shambase::parallel_for(cgh, stepsize, "memcpy kernel", [=](u32 gid) {
                    dest[gid + offset] = acc[gid + offset];
                });
            }));
 
            offset += stepsize;
            remains -= stepsize;
        };
 
        return ev_list;
    }
 
    template<class T>
    inline std::vector<sycl::event> usmbuffer_memcpy(
        sycl::queue &queue, const T *src, sycl::buffer<T> &dest, u64 count) {
 
        u64 offset                  = 0;
        u64 remains                 = count;
        constexpr u64 max_step_size = i32_max / 2;
 
        std::vector<sycl::event> ev_list{};
 
        while (offset < count) {
            u64 stepsize = sham::min(remains, max_step_size);
 
            ev_list.push_back(queue.submit([&, offset](sycl::handler &cgh) {
                sycl::accessor acc{dest, cgh, sycl::write_only};
                shambase::parallel_for(cgh, stepsize, "memcpy kernel", [=](u32 gid) {
                    acc[gid + offset] = src[gid + offset];
                });
            }));
 
            offset += stepsize;
            remains -= stepsize;
        };
 
        return ev_list;
    }
 
    template<class T>
    inline std::vector<sycl::event> usmbuffer_memcpy_discard(
        sycl::queue &queue, const T *src, sycl::buffer<T> &dest, u64 count) {
        u64 offset                  = 0;
        u64 remains                 = count;
        constexpr u64 max_step_size = i32_max / 2;
 
        std::vector<sycl::event> ev_list{};
 
        while (offset < count) {
            u64 stepsize = sham::min(remains, max_step_size);
 
            ev_list.push_back(queue.submit([&, offset](sycl::handler &cgh) {
                sycl::accessor acc{dest, cgh, sycl::write_only, sycl::no_init};
                shambase::parallel_for(cgh, stepsize, "memcpy kernel", [=](u32 gid) {
                    acc[gid + offset] = src[gid + offset];
                });
            }));
 
            offset += stepsize;
            remains -= stepsize;
        };
 
        return ev_list;
    }
 
} // namespace sham