int_chains.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/time.hpp"
#include "shambackends/DeviceBuffer.hpp"
#include "shambackends/DeviceScheduler.hpp"
#include <type_traits>
 
namespace sham::benchmarks {

    enum class IntChainOp {
        Mul, 
        Add, 
    };
 
// Both blocks issue the same number of ALU operations (32 per rotation), so the
// rates they produce are directly comparable.
#define IMAD_4(x, y)                                                                               \
    x = y * x + y;                                                                                 \
    y = x * y + x;                                                                                 \
    x = y * x + y;                                                                                 \
    y = x * y + x;
#define IMAD_16(x, y)                                                                              \
    IMAD_4(x, y);                                                                                  \
    IMAD_4(x, y);                                                                                  \
    IMAD_4(x, y);                                                                                  \
    IMAD_4(x, y);
 
#define IADD_4(x, y)                                                                               \
    x = y + x + y;                                                                                 \
    y = x + y + x;                                                                                 \
    x = y + x + y;                                                                                 \
    y = x + y + x;
#define IADD_16(x, y)                                                                              \
    IADD_4(x, y);                                                                                  \
    IADD_4(x, y);                                                                                  \
    IADD_4(x, y);                                                                                  \
    IADD_4(x, y);

    template<class T, IntChainOp op>
    inline void int_chains(u32 i, int nrotation, T y0, T *__restrict in, T *__restrict out) {
        static_assert(std::is_unsigned_v<T>, "int_chains requires an unsigned type");
 
        T x = in[i];
        T y = y0;
        for (int j = 0; j < nrotation; j++) {
            if constexpr (op == IntChainOp::Mul) {
                IMAD_16(x, y);
            } else {
                IADD_16(x, y);
            }
        }
        out[i] = y;
    }
 
#undef IMAD_4
#undef IMAD_16
#undef IADD_4
#undef IADD_16

    struct int_chains_result {
        std::string func_name; 
        f64 seconds;           
        f64 iops;              
        u32 nrotations;        
    };

    template<class T, IntChainOp op>
    inline int_chains_result int_chains_bench(
        DeviceScheduler_ptr sched, int N, f64 time_threshold) {
 
        sham::DeviceQueue &q = sched->get_queue();
 
        sham::DeviceBuffer<T> x = {size_t(N), sched};
        sham::DeviceBuffer<T> y = {size_t(N), sched};
 
        const T x0 = T{3};
        const T y0 = T{5};
 
        x.fill(x0);
        y.fill(y0);
 
        sham::EventList depends_list;
 
        auto x_ptr = x.get_write_access(depends_list);
        auto y_ptr = y.get_write_access(depends_list);
 
        depends_list.wait();
 
        u32 nrotation = 8;
        double sec    = 0;
 
        auto run_bench = [&q, &N, &x_ptr, &y_ptr, y0](u32 nrotation) -> f64 {
            sham::EventList empty_list{};
 
            shambase::Timer t;
            t.start();
            auto e = q.submit(empty_list, [=](sycl::handler &cgh) {
                cgh.parallel_for(sycl::range<1>{size_t(N)}, [=](sycl::item<1> item) {
                    int_chains<T, op>(item.get_linear_id(), nrotation, y0, x_ptr, y_ptr);
                });
            });
            e.wait();
            t.stop();
 
            return t.elapsed_sec();
        };
 
        // warmup kernel
        run_bench(4);
 
        double ref = run_bench(0);
 
        for (;;) {
 
            sec = run_bench(nrotation);
 
            if (sec >= time_threshold || nrotation >= 256 * 256 * 4) {
                break;
            }
 
            nrotation *= 2;
        }
 
        x.complete_event_state(sycl::event{});
        y.complete_event_state(sycl::event{});
 
        sec -= ref;
 
        u64 op_per_thread = u64(nrotation) * 2_u64 * 16_u64;
        double op_count   = double(N) * double(op_per_thread);
 
        return {
            .func_name  = SourceLocation{}.loc.function_name(),
            .seconds    = sec,
            .iops       = op_count / sec,
            .nrotations = nrotation};
    }
 
} // namespace sham::benchmarks