Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
reduction.cpp
Go to the documentation of this file.
1// -------------------------------------------------------//
2//
3// SHAMROCK code for hydrodynamics
4// Copyright (c) 2021-2026 Timothée David--Cléris <tim.shamrock@proton.me>
5// SPDX-License-Identifier: CeCILL Free Software License Agreement v2.1
6// Shamrock is licensed under the CeCILL 2.1 License, see LICENSE for more information
7//
8// -------------------------------------------------------//
9
16
17#include "shamalgs/details/reduction/reduction.hpp"
18#include "shambase/floats.hpp"
19#include "shambase/memory.hpp"
20#include "shamalgs/details/reduction/fallbackReduction.hpp"
21#include "shamalgs/details/reduction/fallbackReduction_usm.hpp"
22#include "shamalgs/details/reduction/groupReduction.hpp"
23#include "shamalgs/details/reduction/groupReduction_usm.hpp"
24#include "shamalgs/details/reduction/sycl2020reduction.hpp"
25#include "shamalgs/memory.hpp"
26#include "shamalgs/primitives/is_all_true.hpp"
27#include "shambackends/kernel_call.hpp"
28#include "shambackends/math.hpp"
29#include "shambackends/vec.hpp"
30
31namespace shamalgs::reduction {
32
33 template<class T>
34 T sum(sycl::queue &q, sycl::buffer<T> &buf1, u32 start_id, u32 end_id) {
35#ifdef SYCL2020_FEATURE_GROUP_REDUCTION
36 return details::GroupReduction<T, 32>::sum(q, buf1, start_id, end_id);
37#else
38 return details::FallbackReduction<T>::sum(q, buf1, start_id, end_id);
39#endif
40 }
41
42 template<class T>
43 T max(sycl::queue &q, sycl::buffer<T> &buf1, u32 start_id, u32 end_id) {
44#ifdef SYCL2020_FEATURE_GROUP_REDUCTION
45 return details::GroupReduction<T, 32>::max(q, buf1, start_id, end_id);
46#else
47 return details::FallbackReduction<T>::max(q, buf1, start_id, end_id);
48#endif
49 }
50
51 template<class T>
52 T min(sycl::queue &q, sycl::buffer<T> &buf1, u32 start_id, u32 end_id) {
53#ifdef SYCL2020_FEATURE_GROUP_REDUCTION
54 return details::GroupReduction<T, 32>::min(q, buf1, start_id, end_id);
55#else
56 return details::FallbackReduction<T>::min(q, buf1, start_id, end_id);
57#endif
58 }
59
60 template<class T>
61 bool has_nan(sham::DeviceBuffer<T> &buf, u64 cnt) {
62 if constexpr (shambase::VectorProperties<T>::is_float_based) {
63 auto &dev_sched = buf.get_dev_scheduler_ptr();
64
65 // res is filled with 1 if no nan 0 otherwise
66 sham::DeviceBuffer<u8> res(cnt, dev_sched);
67
68 sham::kernel_call(
69 shambase::get_check_ref(dev_sched).get_queue(),
70 sham::MultiRef{buf},
71 sham::MultiRef{res},
72 u32(cnt),
73 [](u32 i, const T *in, u8 *out) {
74 out[i] = !sham::has_nan(in[i]);
75 });
76
77 return !shamalgs::primitives::is_all_true(res, u32(cnt));
78 } else {
79 return false;
80 }
81 }
82
83 template<class T>
84 bool has_inf(sham::DeviceBuffer<T> &buf, u64 cnt) {
85 if constexpr (shambase::VectorProperties<T>::is_float_based) {
86 auto &dev_sched = buf.get_dev_scheduler_ptr();
87
88 // res is filled with 1 if no inf 0 otherwise
89 sham::DeviceBuffer<u8> res(cnt, dev_sched);
90
91 sham::kernel_call(
92 shambase::get_check_ref(dev_sched).get_queue(),
93 sham::MultiRef{buf},
94 sham::MultiRef{res},
95 u32(cnt),
96 [](u32 i, const T *in, u8 *out) {
97 out[i] = !sham::has_inf(in[i]);
98 });
99
100 return !shamalgs::primitives::is_all_true(res, u32(cnt));
101 } else {
102 return false;
103 }
104 }
105
106 template<class T>
107 bool has_nan_or_inf(sham::DeviceBuffer<T> &buf, u64 cnt) {
108 if constexpr (shambase::VectorProperties<T>::is_float_based) {
109 auto &dev_sched = buf.get_dev_scheduler_ptr();
110
111 // res is filled with 1 if no nan or inf 0 otherwise
112 sham::DeviceBuffer<u8> res(cnt, dev_sched);
113
114 sham::kernel_call(
115 shambase::get_check_ref(dev_sched).get_queue(),
116 sham::MultiRef{buf},
117 sham::MultiRef{res},
118 u32(cnt),
119 [](u32 i, const T *in, u8 *out) {
120 out[i] = !sham::has_nan_or_inf(in[i]);
121 });
122
123 return !shamalgs::primitives::is_all_true(res, u32(cnt));
124 } else {
125 return false;
126 }
127 }
128
129#ifndef DOXYGEN
130 #define XMAC_TYPES \
131 X(f32) \
132 X(f32_2) \
133 X(f32_3) \
134 X(f32_4) \
135 X(f32_8) \
136 X(f32_16) \
137 X(f64) \
138 X(f64_2) \
139 X(f64_3) \
140 X(f64_4) \
141 X(f64_8) \
142 X(f64_16) \
143 X(u32) \
144 X(u64) \
145 X(i32) \
146 X(i64) \
147 X(u32_3) \
148 X(u64_3) \
149 X(i64_3) \
150 X(i32_3)
151
152 #define X(_arg_) \
153 template _arg_ sum(sycl::queue &q, sycl::buffer<_arg_> &buf1, u32 start_id, u32 end_id); \
154 template _arg_ max(sycl::queue &q, sycl::buffer<_arg_> &buf1, u32 start_id, u32 end_id); \
155 template _arg_ min(sycl::queue &q, sycl::buffer<_arg_> &buf1, u32 start_id, u32 end_id); \
156 template bool has_nan(sham::DeviceBuffer<_arg_> &buf1, u64 cnt); \
157 template bool has_inf(sham::DeviceBuffer<_arg_> &buf1, u64 cnt); \
158 template bool has_nan_or_inf(sham::DeviceBuffer<_arg_> &buf1, u64 cnt);
159
160 XMAC_TYPES
161 #undef X
162#endif
163
164} // namespace shamalgs::reduction
u8
std::uint8_t u8
8 bit unsigned integer
Definition aliases_int.hpp:29
u32
std::uint32_t u32
32 bit unsigned integer
Definition aliases_int.hpp:27
u64
std::uint64_t u64
64 bit unsigned integer
Definition aliases_int.hpp:26
sham::DeviceBuffer::get_dev_scheduler_ptr
std::shared_ptr< DeviceScheduler > & get_dev_scheduler_ptr()
Gets the Device scheduler pointer corresponding to the held allocation.
Definition DeviceBuffer.hpp:439
fallbackReduction_usm.hpp
groupReduction_usm.hpp
is_all_true.hpp
Boolean reduction algorithm for checking if all elements are non-zero.
kernel_call.hpp
sham::kernel_call
void kernel_call(sham::DeviceQueue &q, RefIn in, RefOut in_out, u32 n, Functor &&func, SourceLocation &&callsite=SourceLocation{})
Submit a kernel to a SYCL queue.
Definition kernel_call.hpp:344
shamalgs::primitives::is_all_true
bool is_all_true(sycl::buffer< T > &buf, u32 cnt)
Check if all elements in a sycl::buffer are non-zero.
Definition is_all_true.cpp:268
shambase::get_check_ref
T & get_check_ref(const std::unique_ptr< T > &ptr, SourceLocation loc=SourceLocation())
Takes a std::unique_ptr and returns a reference to the object it holds. It throws a std::runtime_erro...
Definition memory.hpp:112
memory.hpp
main include file for memory algorithms
Generated on for Shamrock by doxygen 1.15.0