Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_Tuple.H
Go to the documentation of this file.
1
2#ifndef AMREX_TUPLE_H_
3#define AMREX_TUPLE_H_
4#include <AMReX_Config.H>
5
6#include <AMReX_Array.H>
7#include <AMReX_Functional.H>
8#include <AMReX_GpuQualifiers.H>
9#include <AMReX_TypeList.H>
10#include <AMReX_TypeTraits.H>
11
12#include <array>
13#include <functional>
14#include <tuple>
15
16namespace amrex {
17 template <class... Ts>
18 using Tuple = std::tuple<Ts...>;
19}
20
21namespace amrex {
22
24namespace detail {
25
26template <std::size_t I, typename T>
27struct gpu_tuple_element
28{
29 AMREX_GPU_HOST_DEVICE
30 constexpr gpu_tuple_element () requires (std::is_default_constructible_v<T>) {} // NOLINT
31
32 // This overload is what catches arguments the forwarding constructor below
33 // cannot deduce, a braced initializer list above all. Everything else, both
34 // lvalues and rvalues of T, already goes there. Taking this one by value
35 // would therefore add a move for exactly those callers while saving a copy
36 // for nobody.
37 // NOLINTNEXTLINE(modernize-pass-by-value)
38 explicit constexpr gpu_tuple_element (T const& a_value)
39 : m_value(a_value)
40 {}
41
42 template <typename U>
43 requires (std::is_convertible_v<U&&,T>)
44 explicit constexpr gpu_tuple_element (U && a_value) // NOLINT(bugprone-forwarding-reference-overload)
45 : m_value(std::forward<U>(a_value))
46 {}
47
48 T m_value{};
49};
50
51template <std::size_t I, typename... Ts> struct gpu_tuple_impl;
52
53template <std::size_t I, typename Head, typename... Tail>
54struct gpu_tuple_impl<I, Head, Tail...>
55 : public gpu_tuple_impl<I+1, Tail...>,
56 public gpu_tuple_element<I, Head>
57{
58 AMREX_GPU_HOST_DEVICE
59 constexpr gpu_tuple_impl () requires (std::is_default_constructible_v<Head>) {} // NOLINT
60
61 constexpr gpu_tuple_impl (Head const& a_head, Tail const&... a_tail)
62 : gpu_tuple_impl<I+1, Tail...>(a_tail...),
63 gpu_tuple_element<I, Head>(a_head)
64 {}
65
66 template <typename UH, typename... UT>
67 requires (std::is_convertible_v<UH&&,Head>)
68 constexpr gpu_tuple_impl (UH&& a_head, UT &&... a_tail)
69 : gpu_tuple_impl<I+1, Tail...>(std::forward<UT>(a_tail)...),
70 gpu_tuple_element<I, Head>(std::forward<UH>(a_head))
71 {}
72};
73
74template <std::size_t I, typename Head>
75struct gpu_tuple_impl<I, Head>
76 : public gpu_tuple_element<I, Head>
77{
78
79 AMREX_GPU_HOST_DEVICE
80 constexpr gpu_tuple_impl () requires (std::is_default_constructible_v<Head>) {} // NOLINT
81
82 explicit constexpr gpu_tuple_impl (Head const& a_head)
83 : gpu_tuple_element<I, Head>(a_head)
84 {}
85
86 template <typename U>
87 requires (std::is_convertible_v<U&&,Head>)
88 explicit constexpr gpu_tuple_impl (U&& a_head) // NOLINT(bugprone-forwarding-reference-overload)
89 : gpu_tuple_element<I, Head>(std::forward<U>(a_head))
90 {}
91};
92
93} // detail
95
101template <typename... Ts>
102class GpuTuple
103 : public detail::gpu_tuple_impl<0, Ts...>
104{
105public:
106 AMREX_GPU_HOST_DEVICE // Some versions of nvcc require this in debug build
107 constexpr GpuTuple () = default;
108
109 constexpr GpuTuple (Ts const&... args)
110 : detail::gpu_tuple_impl<0, Ts...>(args...)
111 {}
112
113 template <typename... Us>
114 requires (sizeof...(Us) == sizeof...(Ts))
115 constexpr GpuTuple (Us&&... args)
116 : detail::gpu_tuple_impl<0, Ts...>(std::forward<Us>(args)...)
117 {}
118
119 template <typename... Us>
120 AMREX_GPU_HOST_DEVICE
121 inline GpuTuple<Ts...>&
122 operator= (GpuTuple<Us...> const& rhs);
123
124
125 template <typename... Us>
126 AMREX_GPU_HOST_DEVICE
127 inline GpuTuple<Ts...>&
128 operator= (GpuTuple<Us...> && rhs);
129};
130
131// GpuTupleSize
132
133template <typename T> struct GpuTupleSize;
134
135template <typename... Ts>
136struct GpuTupleSize<GpuTuple<Ts...> >
137 : public std::integral_constant<std::size_t, sizeof...(Ts)> {};
138
139// GpuTupleElement
140
141template <std::size_t I, typename T> struct GpuTupleElement;
142
143template <std::size_t I, typename Head, typename... Tail>
144struct GpuTupleElement<I, GpuTuple<Head, Tail...> >
145 : GpuTupleElement<I-1, GpuTuple<Tail...> > {};
146
147template <typename Head, typename... Tail>
148struct GpuTupleElement<0, GpuTuple<Head, Tail...> > {
149 using type = Head;
150};
151
152// get
153
155namespace detail {
156
157template <std::size_t I, typename... Ts>
158AMREX_GPU_HOST_DEVICE
159constexpr
160typename GpuTupleElement<I, GpuTuple<Ts...> >::type&
161get_impl (detail::gpu_tuple_element
162 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type>& te) noexcept
163{
164 return te.m_value;
165}
166
167template <std::size_t I, typename... Ts>
168AMREX_GPU_HOST_DEVICE
169constexpr
170typename GpuTupleElement<I, GpuTuple<Ts...> >::type const&
171get_impl (detail::gpu_tuple_element
172 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type> const& te) noexcept
173{
174 return te.m_value;
175}
176
177template <std::size_t I, typename... Ts>
178AMREX_GPU_HOST_DEVICE
179constexpr
180typename GpuTupleElement<I, GpuTuple<Ts...> >::type &&
181get_impl (detail::gpu_tuple_element
182 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type> && te) noexcept
183{
184 return std::move(te).m_value;
185}
186
187} // detail
189
190template <std::size_t I, typename... Ts>
191AMREX_GPU_HOST_DEVICE
192constexpr
193typename GpuTupleElement<I, GpuTuple<Ts...> >::type&
194get (GpuTuple<Ts...>& tup) noexcept
195{
196 return detail::get_impl<I,Ts...>(tup);
197}
198
199template <std::size_t I, typename... Ts>
200AMREX_GPU_HOST_DEVICE
201constexpr
202typename GpuTupleElement<I, GpuTuple<Ts...> >::type const&
203get (GpuTuple<Ts...> const& tup) noexcept
204{
205 return detail::get_impl<I,Ts...>(tup);
206}
207
208template <std::size_t I, typename... Ts>
209AMREX_GPU_HOST_DEVICE
210constexpr
211typename GpuTupleElement<I, GpuTuple<Ts...> >::type &&
212get (GpuTuple<Ts...> && tup) noexcept
213{
214 return detail::get_impl<I,Ts...>(std::move(tup));
215}
216
218namespace detail {
219 template <std::size_t I, std::size_t N, typename TP1, typename TP2>
220 requires (I<N-1)
221 AMREX_GPU_HOST_DEVICE
222 void
223 tuple_copy (TP1 & a, TP2 && b)
224 {
225 (amrex::get<I>(a) = amrex::get<I>(std::forward<TP2>(b)),
226 tuple_copy<I+1,N>(a,std::forward<TP2>(b)));
227 }
228
229 template <std::size_t I, std::size_t N, typename TP1, typename TP2>
230 requires (I==N-1)
231 AMREX_GPU_HOST_DEVICE
232 void
233 tuple_copy (TP1 & a, TP2 && b)
234 {
235 amrex::get<I>(a) = amrex::get<I>(std::forward<TP2>(b));
236 }
237}
239
240template <typename... Ts>
241template <typename... Us>
242AMREX_GPU_HOST_DEVICE inline GpuTuple<Ts...>&
243GpuTuple<Ts...>::operator= (GpuTuple<Us...> const& rhs)
244{
245 static_assert(sizeof...(Us) == sizeof...(Ts)); // We use static_assert instead of constraint to work around a Clang issue.
246 detail::tuple_copy<0,sizeof...(Ts)>(*this, rhs);
247 return *this;
248}
249
250template <typename... Ts>
251template <typename... Us>
252AMREX_GPU_HOST_DEVICE inline GpuTuple<Ts...>&
253GpuTuple<Ts...>::operator= (GpuTuple<Us...> && rhs)
254{
255 static_assert(sizeof...(Us) == sizeof...(Ts)); // We use static_assert instead of constraint to work around a Clang issue.
256 detail::tuple_copy<0,sizeof...(Ts)>(*this, std::move(rhs));
257 return *this;
258}
259
260// makeTuple
261
263namespace detail {
264 template <typename T> struct unwrap { using type = T; };
265 template <typename T> struct unwrap<std::reference_wrapper<T> > { using type = T&; };
266 template <typename T>
267 using tuple_decay_t = typename unwrap<std::decay_t<T>>::type;
268}
270
271template <typename... Ts>
272AMREX_GPU_HOST_DEVICE
273constexpr
274GpuTuple<detail::tuple_decay_t<Ts>...>
275makeTuple (Ts &&... args)
276{
277 return GpuTuple<detail::tuple_decay_t<Ts>...>(std::forward<Ts>(args)...);
278}
279
281namespace detail {
282 template <typename...> struct tuple_cat_result {};
283
284 template <typename... Ts>
285 struct tuple_cat_result<GpuTuple<Ts...> >
286 {
287 using type = GpuTuple<Ts...>;
288 };
289
290 template <typename... T1s, typename... T2s, typename... TPs>
291 struct tuple_cat_result<GpuTuple<T1s...>,GpuTuple<T2s...>,TPs...>
292 {
293 using type = typename tuple_cat_result<GpuTuple<T1s..., T2s...>, TPs...>::type;
294 };
295
296 // NOLINTNEXTLINE(misc-confusable-identifiers)
297 template <typename R, typename TP1, typename TP2, std::size_t... N1, std::size_t... N2>
298 AMREX_GPU_HOST_DEVICE constexpr R
299 make_tuple (TP1 const& a, TP2 const& b,
300 std::index_sequence<N1...> const& /*n1*/, std::index_sequence<N2...> const& /*n2*/)
301 {
302 return R(amrex::get<N1>(a)..., amrex::get<N2>(b)...);
303 }
304}
306
307// TupleCat
308
309template <typename TP>
310AMREX_GPU_HOST_DEVICE
311constexpr auto
312TupleCat (TP && a) -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type
313{
314 using ReturnType = typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type;
315 return ReturnType(std::forward<TP>(a));
316}
317
318template <typename TP1, typename TP2>
319AMREX_GPU_HOST_DEVICE
320constexpr auto
321TupleCat (TP1 && a, TP2 && b) -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
322 detail::tuple_decay_t<TP2> >::type
323{
324 using ReturnType = typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
325 detail::tuple_decay_t<TP2> >::type;
326 return detail::make_tuple<ReturnType>
327 (std::forward<TP1>(a), std::forward<TP2>(b),
328 std::make_index_sequence<GpuTupleSize<std::decay_t<TP1>>::value>(),
329 std::make_index_sequence<GpuTupleSize<std::decay_t<TP2>>::value>());
330}
331
332template <typename TP1, typename TP2, typename... TPs>
333AMREX_GPU_HOST_DEVICE
334constexpr auto
335TupleCat (TP1&& a, TP2&& b, TPs&&... args)
336 -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
337 detail::tuple_decay_t<TP2>,
338 detail::tuple_decay_t<TPs>...>::type
339{
340 return TupleCat(TupleCat(std::forward<TP1>(a),std::forward<TP2>(b)),
341 std::forward<TPs>(args)...);
342}
343
344// TupleSplit
345
347namespace detail {
348
349 template<std::size_t...Is>
350 struct SplitIndexList {
351 template<std::size_t J>
352 AMREX_GPU_HOST_DEVICE
353 static constexpr std::size_t get () noexcept {
354 std::size_t arr[sizeof...(Is)] = {Is...};
355 return arr[J];
356 }
357
358 template<std::size_t J>
359 AMREX_GPU_HOST_DEVICE
360 static constexpr std::size_t get_exclusive_sum () noexcept {
361 std::size_t arr[sizeof...(Is)] = {Is...};
362 std::size_t sum = 0;
363 for (std::size_t k=0; k<J; ++k) {
364 sum += arr[k];
365 }
366 return sum;
367 }
368 };
369
370 template <std::size_t start, typename... Args, std::size_t... Is>
371 AMREX_GPU_HOST_DEVICE
372 constexpr auto
373 GetSubTuple (const GpuTuple<Args...>& tup, std::index_sequence<Is...>) noexcept
374 {
375 return makeTuple(amrex::get<start+Is>(tup)...);
376 }
377
378 template <typename... Args, std::size_t... Is, typename SIL>
379 AMREX_GPU_HOST_DEVICE
380 constexpr auto
381 TupleSplitImp (const GpuTuple<Args...>& tup, std::index_sequence<Is...>, SIL) noexcept
382 {
383 return makeTuple(
384 GetSubTuple<(SIL::template get_exclusive_sum<Is>())>(
385 tup,
386 std::make_index_sequence<SIL::template get<Is>()>()
387 )...
388 );
389 }
390}
392
397template <std::size_t... Is, typename... Args>
398AMREX_GPU_HOST_DEVICE
399constexpr auto
400TupleSplit (const GpuTuple<Args...>& tup) noexcept
401{
402 static_assert((0 + ... + Is) == sizeof...(Args), "Incorrect total size in TupleSplit");
403 return detail::TupleSplitImp(
404 tup,
405 std::make_index_sequence<sizeof...(Is)>(),
406 detail::SplitIndexList<Is...>()
407 );
408}
409
410// Apply
411
413namespace detail {
414
415 template <typename F, typename... Args>
416 AMREX_GPU_HOST_DEVICE
417 auto INVOKE (F&& f, Args&&... args) -> decltype(f(std::forward<Args>(args)...));
418
419 template <typename V, typename F, typename... Args> struct invoke_result {};
420
421 template <typename F, typename... Args>
422 struct invoke_result<decltype(void(INVOKE(std::declval<F>(), std::declval<Args>()...))),
423 F, Args...>
424 {
425 using type = decltype(INVOKE(std::declval<F>(), std::declval<Args>()...));
426 };
427
428 template <typename F, typename...> struct apply_result {};
429
430 template <typename F, typename... Ts>
431 struct apply_result<F, GpuTuple<Ts...> >
432 {
433 using type = typename invoke_result<void, F, Ts...>::type;
434 };
435
436 template <typename F, typename TP, std::size_t... N>
437 AMREX_GPU_HOST_DEVICE
438 constexpr auto
439 apply_impl (F&& f, TP&& t, std::index_sequence<N...> /*is*/)
440 -> typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
441 {
442 return std::forward<F>(f)(amrex::get<N>(std::forward<TP>(t))...);
443 }
444}
446
447template <typename F, typename TP>
448AMREX_GPU_HOST_DEVICE
449constexpr auto
450Apply (F&& f, TP&& t) -> typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
451{
452 return detail::apply_impl(std::forward<F>(f), std::forward<TP>(t),
453 std::make_index_sequence<GpuTupleSize<std::decay_t<TP>>::value>());
454}
455
456// Tie
457
458template <typename... Args>
459AMREX_GPU_HOST_DEVICE
460constexpr GpuTuple<Args&...>
461Tie (Args&... args) noexcept
462{
463 return GpuTuple<Args&...>(args...);
464}
465
466// ForwardAsTuple
467
468template <typename... Ts>
469AMREX_GPU_HOST_DEVICE
470constexpr GpuTuple<Ts&&...>
471ForwardAsTuple (Ts&&... args) noexcept
472{
473 return GpuTuple<Ts&&...>(std::forward<Ts>(args)...);
474}
475
476// MakeZeroTuple
477
482template <typename... Ts>
483AMREX_GPU_HOST_DEVICE
484constexpr GpuTuple<Ts...>
485MakeZeroTuple (GpuTuple<Ts...>) noexcept
486{
487 return GpuTuple<Ts...>(static_cast<Ts>(0)...);
488}
489
491namespace detail {
492 template <typename T, std::size_t... I>
493 AMREX_GPU_HOST_DEVICE constexpr
494 auto tuple_to_array_helper (T const& tup, std::index_sequence<I...>) {
495 return GpuArray<typename GpuTupleElement<0,T>::type, sizeof...(I)>{amrex::get<I>(tup)...};
496 }
497}
499
500template <typename T>
501AMREX_GPU_HOST_DEVICE constexpr
502auto tupleToArray (GpuTuple<T> const& tup)
503{
504 return GpuArray<T,1>{amrex::get<0>(tup)};
505}
506
508template <typename T, typename T2, typename... Ts>
509requires (Same<T,T2,Ts...>::value)
510AMREX_GPU_HOST_DEVICE constexpr
511auto tupleToArray (GpuTuple<T,T2,Ts...> const& tup)
512{
513 return detail::tuple_to_array_helper(tup, std::index_sequence_for<T,T2,Ts...>{});
514}
515
516} // namespace amrex
517
518// Spcialize std::tuple_size for GpuTuple. Used by structured bindings.
519template<typename... Ts>
520struct std::tuple_size<amrex::GpuTuple<Ts...>> {
521 static constexpr std::size_t value = sizeof...(Ts);
522};
523
524// Spcialize std::tuple_element for GpuTuple. Used by structured bindings.
525template<typename T, typename... Ts>
526struct std::tuple_element<std::size_t{0}, amrex::GpuTuple<T, Ts...>> {
527 using type = T;
528};
529
530template<std::size_t s, typename T, typename... Ts>
531struct std::tuple_element<s, amrex::GpuTuple<T, Ts...>> {
532 using type = typename std::tuple_element_t<s-1, amrex::GpuTuple<Ts...>>;
533};
534
535#endif /*AMREX_TUPLE_H_*/
AMReX_Array.H
Fixed-size array types for use on GPU and CPU.
AMReX_Functional.H
Function object (functor) types for use with reductions and transforms.
AMReX_GpuQualifiers.H
AMREX_GPU_HOST_DEVICE
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
AMReX_TypeList.H
AMReX_TypeTraits.H
amrex::GpuTuple
GPU-compatible tuple.
Definition AMReX_Tuple.H:104
amrex::GpuTuple::operator=
__host__ __device__ GpuTuple< Ts... > & operator=(GpuTuple< Us... > const &rhs)
Definition AMReX_Tuple.H:243
amrex::GpuTuple::GpuTuple
__host__ __device__ constexpr GpuTuple()=default
amrex::GpuTuple::GpuTuple
constexpr GpuTuple(Ts const &... args)
Definition AMReX_Tuple.H:109
amrex
Definition AMReX_Amr.cpp:50
amrex::Tuple
std::tuple< Ts... > Tuple
Definition AMReX_Tuple.H:18
amrex::MakeZeroTuple
__host__ __device__ constexpr GpuTuple< Ts... > MakeZeroTuple(GpuTuple< Ts... >) noexcept
Return a GpuTuple containing all zeros. Note that a default-constructed GpuTuple can have uninitializ...
Definition AMReX_Tuple.H:485
amrex::Tie
__host__ __device__ constexpr GpuTuple< Args &... > Tie(Args &... args) noexcept
Definition AMReX_Tuple.H:461
amrex::makeTuple
__host__ __device__ constexpr GpuTuple< detail::tuple_decay_t< Ts >... > makeTuple(Ts &&... args)
Definition AMReX_Tuple.H:275
amrex::TupleSplit
__host__ __device__ constexpr auto TupleSplit(const GpuTuple< Args... > &tup) noexcept
Returns a GpuTuple of GpuTuples obtained by splitting the input GpuTuple according to the sizes speci...
Definition AMReX_Tuple.H:400
amrex::ForwardAsTuple
__host__ __device__ constexpr GpuTuple< Ts &&... > ForwardAsTuple(Ts &&... args) noexcept
Definition AMReX_Tuple.H:471
amrex::Apply
__host__ __device__ constexpr auto Apply(F &&f, IntVectND< dim > const &iv)
Definition AMReX_IntVect.H:1350
amrex::tupleToArray
__host__ __device__ constexpr auto tupleToArray(GpuTuple< T > const &tup)
Definition AMReX_Tuple.H:502
amrex::TupleCat
__host__ __device__ constexpr auto TupleCat(TP &&a) -> typename detail::tuple_cat_result< detail::tuple_decay_t< TP > >::type
Definition AMReX_Tuple.H:312
amrex::get
__host__ __device__ constexpr int get(IntVectND< dim > const &iv) noexcept
Get I'th element of IntVectND<dim>
Definition AMReX_IntVect.H:1334
amrex::GpuArray
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:52
amrex::GpuTupleElement
Definition AMReX_Tuple.H:141
amrex::GpuTupleSize
Definition AMReX_Tuple.H:133
std::tuple_element< s, amrex::GpuTuple< T, Ts... > >::type
typename std::tuple_element_t< s-1, amrex::GpuTuple< Ts... > > type
Definition AMReX_Tuple.H:532
std::tuple_element< std::size_t{0}, amrex::GpuTuple< T, Ts... > >::type
T type
Definition AMReX_Tuple.H:527