Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_ReduceSIMD.H
Go to the documentation of this file.
1#ifndef AMREX_REDUCE_SIMD_H_
2#define AMREX_REDUCE_SIMD_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_Extension.H>
6#include <AMReX_OpenMP.H>
7#include <AMReX_Reduce.H> // provides amrex::SIMDindex via AMReX_Gpu.H
8#include <AMReX_SIMD.H>
9
10#include <concepts>
11#include <cstddef>
12#include <limits>
13#include <type_traits>
14
15namespace amrex {
16
18namespace Reduce::detail {
19
21 template <typename P>
22 inline constexpr bool is_simd_reduce_op =
23 std::is_same_v<P, ReduceOpSum> ||
24 std::is_same_v<P, ReduceOpMin> ||
25 std::is_same_v<P, ReduceOpMax>;
26
27#if defined(AMREX_USE_SIMD) && !defined(AMREX_USE_GPU)
28
30 template <int W, typename T>
31 struct MakeSIMDTuple;
32
33 template <int W, typename... Ts>
34 struct MakeSIMDTuple<W, GpuTuple<Ts...>>
35 {
36 using type = GpuTuple<amrex::simd::stdx::fixed_size_simd<Ts, W>...>;
37 };
38
40 template <typename P, typename T>
42 void simd_init (T& t)
43 {
44 using E = typename T::value_type;
45 if constexpr (std::is_same_v<P, ReduceOpSum>) {
46 t = T(E(0));
47 } else if constexpr (std::is_same_v<P, ReduceOpMin>) {
48 t = T(std::numeric_limits<E>::max());
49 } else if constexpr (std::is_same_v<P, ReduceOpMax>) {
50 t = T(std::numeric_limits<E>::lowest());
51 } else {
52 static_assert(is_simd_reduce_op<P>,
53 "SIMD reductions support only ReduceOpSum, ReduceOpMin and ReduceOpMax");
54 }
55 }
56
58 template <typename P, typename T>
60 void simd_local_update (T& d, T const& s)
61 {
62 namespace stdx = amrex::simd::stdx;
63 if constexpr (std::is_same_v<P, ReduceOpSum>) {
64 d += s;
65 } else if constexpr (std::is_same_v<P, ReduceOpMin>) {
66 d = stdx::min(d, s);
67 } else if constexpr (std::is_same_v<P, ReduceOpMax>) {
68 d = stdx::max(d, s);
69 } else {
70 static_assert(is_simd_reduce_op<P>,
71 "SIMD reductions support only ReduceOpSum, ReduceOpMin and ReduceOpMax");
72 }
73 }
74
76 template <typename P, typename T, typename S>
78 void simd_horizontal_update (T& d, S const& s)
79 {
80 namespace stdx = amrex::simd::stdx;
81 if constexpr (std::is_same_v<P, ReduceOpSum>) {
82 d += stdx::reduce(s);
83 } else if constexpr (std::is_same_v<P, ReduceOpMin>) {
84 d = amrex::min(d, stdx::hmin(s));
85 } else if constexpr (std::is_same_v<P, ReduceOpMax>) {
86 d = amrex::max(d, stdx::hmax(s));
87 } else {
88 static_assert(is_simd_reduce_op<P>,
89 "SIMD reductions support only ReduceOpSum, ReduceOpMin and ReduceOpMax");
90 }
91 }
92
93 // The recursions below mirror for_each_local / for_each_init in AMReX_Reduce.H.
94
95 template <std::size_t I, typename T, typename P>
97 void for_each_simd_init (T& t)
98 {
99 simd_init<P>(amrex::get<I>(t));
100 }
101
102 template <std::size_t I, typename T, typename P, typename P1, typename... Ps>
104 void for_each_simd_init (T& t)
105 {
106 simd_init<P>(amrex::get<I>(t));
107 for_each_simd_init<I+1,T,P1,Ps...>(t);
108 }
109
110 template <std::size_t I, typename T, typename P>
112 void for_each_simd_local (T& d, T const& s)
113 {
114 simd_local_update<P>(amrex::get<I>(d), amrex::get<I>(s));
115 }
116
117 template <std::size_t I, typename T, typename P, typename P1, typename... Ps>
119 void for_each_simd_local (T& d, T const& s)
120 {
121 simd_local_update<P>(amrex::get<I>(d), amrex::get<I>(s));
122 for_each_simd_local<I+1,T,P1,Ps...>(d, s);
123 }
124
125 template <std::size_t I, typename TD, typename TS, typename P>
127 void for_each_horizontal (TD& d, TS const& s)
128 {
129 simd_horizontal_update<P>(amrex::get<I>(d), amrex::get<I>(s));
130 }
131
132 template <std::size_t I, typename TD, typename TS, typename P, typename P1, typename... Ps>
134 void for_each_horizontal (TD& d, TS const& s)
135 {
136 simd_horizontal_update<P>(amrex::get<I>(d), amrex::get<I>(s));
137 for_each_horizontal<I+1,TD,TS,P1,Ps...>(d, s);
138 }
139
140#endif // AMREX_USE_SIMD && !AMREX_USE_GPU
141}
143
144#ifndef AMREX_USE_GPU
145
180template <int WIDTH, std::integral N, typename D, typename F, typename... Ps>
182void evalReduceSIMD (N n, D& reduce_data, F const& f, ReduceOps<Ps...>& /*reduce_ops*/)
183{
184 static_assert(WIDTH >= 1, "evalReduceSIMD: WIDTH must be at least 1");
185 static_assert((Reduce::detail::is_simd_reduce_op<Ps> && ...),
186 "evalReduceSIMD supports only ReduceOpSum, ReduceOpMin and ReduceOpMax");
187
188 using ReduceTuple = typename D::Type;
189 ReduceTuple rr; // scalar accumulator (remainder and horizontal fold target)
190 Reduce::detail::for_each_init<0, ReduceTuple, Ps...>(rr);
191 N i = 0;
192#ifdef AMREX_USE_SIMD
193 if constexpr (WIDTH > 1) {
194 using SimdTuple = typename Reduce::detail::MakeSIMDTuple<WIDTH, ReduceTuple>::type;
195 static_assert(std::is_same_v<SimdTuple,
196 std::decay_t<decltype(f(SIMDindex<WIDTH, N>{0}))>>,
197 "evalReduceSIMD: for SIMDindex<WIDTH> the callable must return the "
198 "SIMD-widened ReduceData tuple (fixed_size_simd of each component type). "
199 "Load particle/array data with amrex::simd::load_1d to get matching types.");
200 SimdTuple sr;
201 Reduce::detail::for_each_simd_init<0, SimdTuple, Ps...>(sr);
202 // vectorized main loop over full SIMD lanes
203 for (; i + WIDTH <= n; i += WIDTH) {
204 Reduce::detail::for_each_simd_local<0, SimdTuple, Ps...>(sr, f(SIMDindex<WIDTH, N>{i}));
205 }
206 Reduce::detail::for_each_horizontal<0, ReduceTuple, SimdTuple, Ps...>(rr, sr);
207 }
208#endif
209 // scalar handling of the remainder (or the whole range without SIMD)
210 for (; i < n; ++i) {
211 Reduce::detail::for_each_local<0, ReduceTuple, Ps...>(rr, f(SIMDindex<1, N>{i}));
212 }
213 Reduce::detail::for_each_local<0, ReduceTuple, Ps...>(
214 reduce_data.reference(OpenMP::get_thread_num()), rr);
215}
216
217#endif // !AMREX_USE_GPU
218}
219
220#endif
Compiler- and backend-specific extension macros (e.g., restrict, SIMD, inline).
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
#define AMREX_ATTRIBUTE_FLATTEN_FOR
Definition AMReX_Extension.H:156
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:31
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:53
constexpr int get_thread_num()
Definition AMReX_OpenMP.H:37
Definition AMReX_SIMD.H:25
Definition AMReX_Amr.cpp:50