ErrorChecker.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/DeviceBuffer.hpp"
#include "shambackends/sycl.hpp"
 
namespace shamalgs::atomic {

    struct ErrorCheckerFlags {

        sham::DeviceBuffer<u32> buf_err;

        ErrorCheckerFlags(sham::DeviceScheduler_ptr sched) : buf_err(1, sched) { buf_err.fill(0); }

        struct accessed {
            u32 *ptr; 

            void set_flag_on(u32 flag_val) const {
                sycl::atomic_ref<
                    u32,
                    sycl::memory_order_relaxed,
                    sycl::memory_scope_device,
                    sycl::access::address_space::global_space>
                    atom(*ptr);
                atom |= flag_val;
            }
        };

        accessed get_write_access(sham::EventList &depends_list) {
            return accessed{buf_err.get_write_access(depends_list)};
        }

        void complete_event_state(sycl::event e) { buf_err.complete_event_state(e); }

        u32 get_output() { return buf_err.copy_to_stdvec().at(0); }
    };

    struct ErrorCheckCounter {

        sham::DeviceBuffer<u32> buf_err;

        ErrorCheckCounter(sham::DeviceScheduler_ptr sched, u32 error_counter)
            : buf_err(error_counter, sched) {
            buf_err.fill(0_u32);
        }

        struct accessed {
            u32 *ptr; 

            void set_error(u32 id) const {
                sycl::atomic_ref<
                    u32,
                    sycl::memory_order_relaxed,
                    sycl::memory_scope_device,
                    sycl::access::address_space::global_space>
                    atom(ptr[id]);
                atom.fetch_add(1_u32);
            }
        };

        accessed get_write_access(sham::EventList &depends_list) {
            return accessed{buf_err.get_write_access(depends_list)};
        }

        void complete_event_state(sycl::event e) { buf_err.complete_event_state(e); }

        std::vector<u32> get_outputs() { return buf_err.copy_to_stdvec(); }
    };
 
} // namespace shamalgs::atomic