equals.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/exception.hpp"
#include "shamalgs/primitives/is_all_true.hpp"
#include "shambackends/DeviceBuffer.hpp"
#include "shambackends/EventList.hpp"
#include "shambackends/kernel_call.hpp"
#include "shambackends/math.hpp"
#include "shambackends/sycl.hpp"
#include "shamsys/NodeInstance.hpp"
 
namespace shamalgs::primitives {
 
    template<class T>
    bool equals(sycl::queue &q, sycl::buffer<T> &buf1, sycl::buffer<T> &buf2, u32 cnt) {
 
        if (buf1.size() < cnt) {
            throw shambase::make_except_with_loc<std::invalid_argument>("buf 1 is larger than cnt");
        }
 
        if (buf2.size() < cnt) {
            throw shambase::make_except_with_loc<std::invalid_argument>("buf 2 is larger than cnt");
        }
 
        sham::DeviceBuffer<u8> res(cnt, shamsys::instance::get_compute_scheduler_ptr());
 
        sham::EventList deps;
        auto out = res.get_write_access(deps);
 
        deps.set_consumed(true);
        auto e = q.submit([&](sycl::handler &cgh) {
            cgh.depends_on(deps.get_events());
            sycl::accessor acc1{buf1, cgh, sycl::read_only};
            sycl::accessor acc2{buf2, cgh, sycl::read_only};
 
            cgh.parallel_for(sycl::range{cnt}, [=](sycl::item<1> item) {
                out[item] = sham::equals(acc1[item], acc2[item]);
            });
        });
 
        res.complete_event_state(e);
 
        return shamalgs::primitives::is_all_true(res, cnt);
    }
 
    template<class T>
    inline bool equals(
        const sham::DeviceScheduler_ptr &dev_sched,
        sham::DeviceBuffer<T> &buf1,
        sham::DeviceBuffer<T> &buf2,
        u32 cnt) {
 
        // kernel call does not support 0 elements
        if (cnt == 0) {
            return true;
        }
 
        // if the buffers are the same early return true
        if (&buf1 == &buf2) {
            return true;
        }
 
        if (buf1.get_size() < cnt) {
            throw shambase::make_except_with_loc<std::invalid_argument>("buf 1 is larger than cnt");
        }
 
        if (buf2.get_size() < cnt) {
            throw shambase::make_except_with_loc<std::invalid_argument>("buf 2 is larger than cnt");
        }
 
        sham::DeviceBuffer<u8> res(cnt, dev_sched);
 
        auto &q = shambase::get_check_ref(dev_sched).get_queue();
 
        sham::kernel_call(
            q,
            sham::MultiRef{buf1, buf2},
            sham::MultiRef{res},
            cnt,
            [](u32 i, const T *__restrict acc1, const T *__restrict acc2, u8 *__restrict out) {
                out[i] = sham::equals(acc1[i], acc2[i]);
            });
 
        return shamalgs::primitives::is_all_true(res, cnt);
    }
 
    template<class T>
    inline bool equals(
        const sham::DeviceScheduler_ptr &q,
        sham::DeviceBuffer<T> &buf1,
        sham::DeviceBuffer<T> &buf2) {
 
        bool same_size = buf1.get_size() == buf2.get_size();
        if (!same_size) {
            return false;
        }
 
        return equals(q, buf1, buf2, buf1.get_size());
    }
 
    template<class T>
    bool equals(sycl::queue &q, sycl::buffer<T> &buf1, sycl::buffer<T> &buf2) {
        bool same_size = buf1.size() == buf2.size();
        if (!same_size) {
            return false;
        }
 
        return equals(q, buf1, buf2, buf1.size());
    }
 
    template<class T>
    bool equals_ptr_s(
        sycl::queue &q,
        const std::unique_ptr<sycl::buffer<T>> &buf1,
        const std::unique_ptr<sycl::buffer<T>> &buf2,
        u32 cnt) {
        bool same_alloc = bool(buf1) == bool(buf2);
 
        if (!same_alloc) {
            return false;
        }
 
        if (!bool(buf1)) {
            return true;
        }
 
        return equals(q, *buf1, *buf2, cnt);
    }
 
    template<class T>
    bool equals_ptr(
        sycl::queue &q,
        const std::unique_ptr<sycl::buffer<T>> &buf1,
        const std::unique_ptr<sycl::buffer<T>> &buf2) {
        bool same_alloc = bool(buf1) == bool(buf2);
 
        if (!same_alloc) {
            return false;
        }
 
        if (!bool(buf1)) {
            return true;
        }
 
        return equals(q, *buf1, *buf2);
    }
} // namespace shamalgs::primitives