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