algorithm.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 "shamalgs/primitives/sort_by_key_pow2_len.hpp"
#include "shambackends/DeviceBuffer.hpp"
#include "shambackends/DeviceQueue.hpp"
#include "shambackends/sycl.hpp"
#include "shambackends/typeAliasVec.hpp"

namespace shamalgs::algorithm {

    template<class Tkey, class Tval>
    void sort_by_key_pow2_len(
        sycl::queue &q, sycl::buffer<Tkey> &buf_key, sycl::buffer<Tval> &buf_values, u32 len) {
        shamalgs::primitives::sort_by_key_pow2_len(q, buf_key, buf_values, len);
    }
 
    template<class Tkey, class Tval>
    void sort_by_key_pow2_len(
        const sham::DeviceScheduler_ptr &sched,
        sham::DeviceBuffer<Tkey> &buf_key,
        sham::DeviceBuffer<Tval> &buf_values,
        u32 len) {
        shamalgs::primitives::sort_by_key_pow2_len(sched, buf_key, buf_values, len);
    }

    template<class Fct>
    inline sycl::buffer<typename std::invoke_result_t<Fct, u32>> gen_buffer_device(
        sycl::queue &q, u32 len, Fct &&func) {
 
        using ret_t = typename std::invoke_result_t<Fct, u32>;
 
        sycl::buffer<ret_t> ret(len);
 
        q.submit([&](sycl::handler &cgh) {
            sycl::accessor out{ret, cgh, sycl::write_only, sycl::no_init};
 
            cgh.parallel_for(sycl::range<1>(len), [=](sycl::item<1> item) {
                out[item] = func(item.get_linear_id());
            });
        });
 
        return std::move(ret);
    }

    template<class T>
    sycl::buffer<T> index_remap(
        sycl::queue &q, sycl::buffer<T> &source_buf, sycl::buffer<u32> &index_map, u32 len);

    template<class T>
    sycl::buffer<T> index_remap_nvar(
        sycl::queue &q,
        sycl::buffer<T> &source_buf,
        sycl::buffer<u32> &index_map,
        u32 len,
        u32 nvar);
 
    template<class T>
    void index_remap(
        const sham::DeviceScheduler_ptr &sched,
        sham::DeviceBuffer<T> &source,
        sham::DeviceBuffer<T> &dest,
        sham::DeviceBuffer<u32> &index_map,
        u32 len);
 
    template<class T>
    void index_remap_nvar(
        const sham::DeviceScheduler_ptr &sched,
        sham::DeviceBuffer<T> &source,
        sham::DeviceBuffer<T> &dest,
        sham::DeviceBuffer<u32> &index_map,
        u32 len,
        u32 nvar);
 
    template<class T>
    sham::DeviceBuffer<T> index_remap(
        const sham::DeviceScheduler_ptr &sched_ptr,
        sham::DeviceBuffer<T> &source,
        sham::DeviceBuffer<u32> &index_map,
        u32 len) {
 
        sham::DeviceBuffer<T> dest(len, sched_ptr);
        index_remap<T>(sched_ptr, source, dest, index_map, len);
        return dest;
    }
 
    template<class T>
    sham::DeviceBuffer<T> index_remap_nvar(
        const sham::DeviceScheduler_ptr &sched_ptr,
        sham::DeviceBuffer<T> &source,
        sham::DeviceBuffer<u32> &index_map,
        u32 len,
        u32 nvar) {
 
        sham::DeviceBuffer<T> dest(len * nvar, sched_ptr);
        index_remap_nvar<T>(sched_ptr, source, dest, index_map, len, nvar);
        return dest;
    }

    sycl::buffer<u32> gen_buffer_index(sycl::queue &q, u32 len);
 
} // namespace shamalgs::algorithm