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 explicit constexpr gpu_tuple_element (T const& a_value)
33 : m_value(a_value)
34 {}
35
36 template <typename U>
37 requires (std::is_convertible_v<U&&,T>)
38 explicit constexpr gpu_tuple_element (U && a_value) // NOLINT(bugprone-forwarding-reference-overload)
39 : m_value(std::forward<U>(a_value))
40 {}
41
42 T m_value{};
43};
44
45template <std::size_t I, typename... Ts> struct gpu_tuple_impl;
46
47template <std::size_t I, typename Head, typename... Tail>
48struct gpu_tuple_impl<I, Head, Tail...>
49 : public gpu_tuple_impl<I+1, Tail...>,
50 public gpu_tuple_element<I, Head>
51{
52 AMREX_GPU_HOST_DEVICE
53 constexpr gpu_tuple_impl () requires (std::is_default_constructible_v<Head>) {} // NOLINT
54
55 constexpr gpu_tuple_impl (Head const& a_head, Tail const&... a_tail)
56 : gpu_tuple_impl<I+1, Tail...>(a_tail...),
57 gpu_tuple_element<I, Head>(a_head)
58 {}
59
60 template <typename UH, typename... UT>
61 requires (std::is_convertible_v<UH&&,Head>)
62 constexpr gpu_tuple_impl (UH&& a_head, UT &&... a_tail)
63 : gpu_tuple_impl<I+1, Tail...>(std::forward<UT>(a_tail)...),
64 gpu_tuple_element<I, Head>(std::forward<UH>(a_head))
65 {}
66};
67
68template <std::size_t I, typename Head>
69struct gpu_tuple_impl<I, Head>
70 : public gpu_tuple_element<I, Head>
71{
72
73 AMREX_GPU_HOST_DEVICE
74 constexpr gpu_tuple_impl () requires (std::is_default_constructible_v<Head>) {} // NOLINT
75
76 explicit constexpr gpu_tuple_impl (Head const& a_head)
77 : gpu_tuple_element<I, Head>(a_head)
78 {}
79
80 template <typename U>
81 requires (std::is_convertible_v<U&&,Head>)
82 explicit constexpr gpu_tuple_impl (U&& a_head) // NOLINT(bugprone-forwarding-reference-overload)
83 : gpu_tuple_element<I, Head>(std::forward<U>(a_head))
84 {}
85};
86
87} // detail
89
95template <typename... Ts>
96class GpuTuple
97 : public detail::gpu_tuple_impl<0, Ts...>
98{
99public:
100 AMREX_GPU_HOST_DEVICE // Some versions of nvcc require this in debug build
101 constexpr GpuTuple () = default;
102
103 constexpr GpuTuple (Ts const&... args)
104 : detail::gpu_tuple_impl<0, Ts...>(args...)
105 {}
106
107 template <typename... Us>
108 requires (sizeof...(Us) == sizeof...(Ts))
109 constexpr GpuTuple (Us&&... args)
110 : detail::gpu_tuple_impl<0, Ts...>(std::forward<Us>(args)...)
111 {}
112
113 template <typename... Us>
114 AMREX_GPU_HOST_DEVICE
115 inline GpuTuple<Ts...>&
116 operator= (GpuTuple<Us...> const& rhs);
117
118
119 template <typename... Us>
120 AMREX_GPU_HOST_DEVICE
121 inline GpuTuple<Ts...>&
122 operator= (GpuTuple<Us...> && rhs);
123};
124
125// GpuTupleSize
126
127template <typename T> struct GpuTupleSize;
128
129template <typename... Ts>
130struct GpuTupleSize<GpuTuple<Ts...> >
131 : public std::integral_constant<std::size_t, sizeof...(Ts)> {};
132
133// GpuTupleElement
134
135template <std::size_t I, typename T> struct GpuTupleElement;
136
137template <std::size_t I, typename Head, typename... Tail>
138struct GpuTupleElement<I, GpuTuple<Head, Tail...> >
139 : GpuTupleElement<I-1, GpuTuple<Tail...> > {};
140
141template <typename Head, typename... Tail>
142struct GpuTupleElement<0, GpuTuple<Head, Tail...> > {
143 using type = Head;
144};
145
146// get
147
149namespace detail {
150
151template <std::size_t I, typename... Ts>
152AMREX_GPU_HOST_DEVICE
153constexpr
154typename GpuTupleElement<I, GpuTuple<Ts...> >::type&
155get_impl (detail::gpu_tuple_element
156 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type>& te) noexcept
157{
158 return te.m_value;
159}
160
161template <std::size_t I, typename... Ts>
162AMREX_GPU_HOST_DEVICE
163constexpr
164typename GpuTupleElement<I, GpuTuple<Ts...> >::type const&
165get_impl (detail::gpu_tuple_element
166 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type> const& te) noexcept
167{
168 return te.m_value;
169}
170
171template <std::size_t I, typename... Ts>
172AMREX_GPU_HOST_DEVICE
173constexpr
174typename GpuTupleElement<I, GpuTuple<Ts...> >::type &&
175get_impl (detail::gpu_tuple_element
176 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type> && te) noexcept
177{
178 return std::move(te).m_value;
179}
180
181} // detail
183
184template <std::size_t I, typename... Ts>
185AMREX_GPU_HOST_DEVICE
186constexpr
187typename GpuTupleElement<I, GpuTuple<Ts...> >::type&
188get (GpuTuple<Ts...>& tup) noexcept
189{
190 return detail::get_impl<I,Ts...>(tup);
191}
192
193template <std::size_t I, typename... Ts>
194AMREX_GPU_HOST_DEVICE
195constexpr
196typename GpuTupleElement<I, GpuTuple<Ts...> >::type const&
197get (GpuTuple<Ts...> const& tup) noexcept
198{
199 return detail::get_impl<I,Ts...>(tup);
200}
201
202template <std::size_t I, typename... Ts>
203AMREX_GPU_HOST_DEVICE
204constexpr
205typename GpuTupleElement<I, GpuTuple<Ts...> >::type &&
206get (GpuTuple<Ts...> && tup) noexcept
207{
208 return detail::get_impl<I,Ts...>(std::move(tup));
209}
210
212namespace detail {
213 template <std::size_t I, std::size_t N, typename TP1, typename TP2>
214 requires (I<N-1)
215 AMREX_GPU_HOST_DEVICE
216 void
217 tuple_copy (TP1 & a, TP2 && b)
218 {
219 (amrex::get<I>(a) = amrex::get<I>(std::forward<TP2>(b)),
220 tuple_copy<I+1,N>(a,std::forward<TP2>(b)));
221 }
222
223 template <std::size_t I, std::size_t N, typename TP1, typename TP2>
224 requires (I==N-1)
225 AMREX_GPU_HOST_DEVICE
226 void
227 tuple_copy (TP1 & a, TP2 && b)
228 {
229 amrex::get<I>(a) = amrex::get<I>(std::forward<TP2>(b));
230 }
231}
233
234template <typename... Ts>
235template <typename... Us>
236AMREX_GPU_HOST_DEVICE inline GpuTuple<Ts...>&
237GpuTuple<Ts...>::operator= (GpuTuple<Us...> const& rhs)
238{
239 static_assert(sizeof...(Us) == sizeof...(Ts)); // We use static_assert instead of constraint to work around a Clang issue.
240 detail::tuple_copy<0,sizeof...(Ts)>(*this, rhs);
241 return *this;
242}
243
244template <typename... Ts>
245template <typename... Us>
246AMREX_GPU_HOST_DEVICE inline GpuTuple<Ts...>&
247GpuTuple<Ts...>::operator= (GpuTuple<Us...> && rhs)
248{
249 static_assert(sizeof...(Us) == sizeof...(Ts)); // We use static_assert instead of constraint to work around a Clang issue.
250 detail::tuple_copy<0,sizeof...(Ts)>(*this, std::move(rhs));
251 return *this;
252}
253
254// makeTuple
255
257namespace detail {
258 template <typename T> struct unwrap { using type = T; };
259 template <typename T> struct unwrap<std::reference_wrapper<T> > { using type = T&; };
260 template <typename T>
261 using tuple_decay_t = typename unwrap<std::decay_t<T>>::type;
262}
264
265template <typename... Ts>
266AMREX_GPU_HOST_DEVICE
267constexpr
268GpuTuple<detail::tuple_decay_t<Ts>...>
269makeTuple (Ts &&... args)
270{
271 return GpuTuple<detail::tuple_decay_t<Ts>...>(std::forward<Ts>(args)...);
272}
273
275namespace detail {
276 template <typename...> struct tuple_cat_result {};
277
278 template <typename... Ts>
279 struct tuple_cat_result<GpuTuple<Ts...> >
280 {
281 using type = GpuTuple<Ts...>;
282 };
283
284 template <typename... T1s, typename... T2s, typename... TPs>
285 struct tuple_cat_result<GpuTuple<T1s...>,GpuTuple<T2s...>,TPs...>
286 {
287 using type = typename tuple_cat_result<GpuTuple<T1s..., T2s...>, TPs...>::type;
288 };
289
290 // NOLINTNEXTLINE(misc-confusable-identifiers)
291 template <typename R, typename TP1, typename TP2, std::size_t... N1, std::size_t... N2>
292 AMREX_GPU_HOST_DEVICE constexpr R
293 make_tuple (TP1 const& a, TP2 const& b,
294 std::index_sequence<N1...> const& /*n1*/, std::index_sequence<N2...> const& /*n2*/)
295 {
296 return R(amrex::get<N1>(a)..., amrex::get<N2>(b)...);
297 }
298}
300
301// TupleCat
302
303template <typename TP>
304AMREX_GPU_HOST_DEVICE
305constexpr auto
306TupleCat (TP && a) -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type
307{
308 using ReturnType = typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type;
309 return ReturnType(std::forward<TP>(a));
310}
311
312template <typename TP1, typename TP2>
313AMREX_GPU_HOST_DEVICE
314constexpr auto
315TupleCat (TP1 && a, TP2 && b) -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
316 detail::tuple_decay_t<TP2> >::type
317{
318 using ReturnType = typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
319 detail::tuple_decay_t<TP2> >::type;
320 return detail::make_tuple<ReturnType>
321 (std::forward<TP1>(a), std::forward<TP2>(b),
322 std::make_index_sequence<GpuTupleSize<std::decay_t<TP1>>::value>(),
323 std::make_index_sequence<GpuTupleSize<std::decay_t<TP2>>::value>());
324}
325
326template <typename TP1, typename TP2, typename... TPs>
327AMREX_GPU_HOST_DEVICE
328constexpr auto
329TupleCat (TP1&& a, TP2&& b, TPs&&... args)
330 -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
331 detail::tuple_decay_t<TP2>,
332 detail::tuple_decay_t<TPs>...>::type
333{
334 return TupleCat(TupleCat(std::forward<TP1>(a),std::forward<TP2>(b)),
335 std::forward<TPs>(args)...);
336}
337
338// TupleSplit
339
341namespace detail {
342
343 template<std::size_t...Is>
344 struct SplitIndexList {
345 template<std::size_t J>
346 AMREX_GPU_HOST_DEVICE
347 static constexpr std::size_t get () noexcept {
348 std::size_t arr[sizeof...(Is)] = {Is...};
349 return arr[J];
350 }
351
352 template<std::size_t J>
353 AMREX_GPU_HOST_DEVICE
354 static constexpr std::size_t get_exclusive_sum () noexcept {
355 std::size_t arr[sizeof...(Is)] = {Is...};
356 std::size_t sum = 0;
357 for (std::size_t k=0; k<J; ++k) {
358 sum += arr[k];
359 }
360 return sum;
361 }
362 };
363
364 template <std::size_t start, typename... Args, std::size_t... Is>
365 AMREX_GPU_HOST_DEVICE
366 constexpr auto
367 GetSubTuple (const GpuTuple<Args...>& tup, std::index_sequence<Is...>) noexcept
368 {
369 return makeTuple(amrex::get<start+Is>(tup)...);
370 }
371
372 template <typename... Args, std::size_t... Is, typename SIL>
373 AMREX_GPU_HOST_DEVICE
374 constexpr auto
375 TupleSplitImp (const GpuTuple<Args...>& tup, std::index_sequence<Is...>, SIL) noexcept
376 {
377 return makeTuple(
378 GetSubTuple<(SIL::template get_exclusive_sum<Is>())>(
379 tup,
380 std::make_index_sequence<SIL::template get<Is>()>()
381 )...
382 );
383 }
384}
386
391template <std::size_t... Is, typename... Args>
392AMREX_GPU_HOST_DEVICE
393constexpr auto
394TupleSplit (const GpuTuple<Args...>& tup) noexcept
395{
396 static_assert((0 + ... + Is) == sizeof...(Args), "Incorrect total size in TupleSplit");
397 return detail::TupleSplitImp(
398 tup,
399 std::make_index_sequence<sizeof...(Is)>(),
400 detail::SplitIndexList<Is...>()
401 );
402}
403
404// Apply
405
407namespace detail {
408
409 template <typename F, typename... Args>
410 AMREX_GPU_HOST_DEVICE
411 auto INVOKE (F&& f, Args&&... args) -> decltype(f(std::forward<Args>(args)...));
412
413 template <typename V, typename F, typename... Args> struct invoke_result {};
414
415 template <typename F, typename... Args>
416 struct invoke_result<decltype(void(INVOKE(std::declval<F>(), std::declval<Args>()...))),
417 F, Args...>
418 {
419 using type = decltype(INVOKE(std::declval<F>(), std::declval<Args>()...));
420 };
421
422 template <typename F, typename...> struct apply_result {};
423
424 template <typename F, typename... Ts>
425 struct apply_result<F, GpuTuple<Ts...> >
426 {
427 using type = typename invoke_result<void, F, Ts...>::type;
428 };
429
430 template <typename F, typename TP, std::size_t... N>
431 AMREX_GPU_HOST_DEVICE
432 constexpr auto
433 apply_impl (F&& f, TP&& t, std::index_sequence<N...> /*is*/)
434 -> typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
435 {
436 return std::forward<F>(f)(amrex::get<N>(std::forward<TP>(t))...);
437 }
438}
440
441template <typename F, typename TP>
442AMREX_GPU_HOST_DEVICE
443constexpr auto
444Apply (F&& f, TP&& t) -> typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
445{
446 return detail::apply_impl(std::forward<F>(f), std::forward<TP>(t),
447 std::make_index_sequence<GpuTupleSize<std::decay_t<TP>>::value>());
448}
449
450// Tie
451
452template <typename... Args>
453AMREX_GPU_HOST_DEVICE
454constexpr GpuTuple<Args&...>
455Tie (Args&... args) noexcept
456{
457 return GpuTuple<Args&...>(args...);
458}
459
460// ForwardAsTuple
461
462template <typename... Ts>
463AMREX_GPU_HOST_DEVICE
464constexpr GpuTuple<Ts&&...>
465ForwardAsTuple (Ts&&... args) noexcept
466{
467 return GpuTuple<Ts&&...>(std::forward<Ts>(args)...);
468}
469
470// MakeZeroTuple
471
476template <typename... Ts>
477AMREX_GPU_HOST_DEVICE
478constexpr GpuTuple<Ts...>
479MakeZeroTuple (GpuTuple<Ts...>) noexcept
480{
481 return GpuTuple<Ts...>(static_cast<Ts>(0)...);
482}
483
485namespace detail {
486 template <typename T, std::size_t... I>
487 AMREX_GPU_HOST_DEVICE constexpr
488 auto tuple_to_array_helper (T const& tup, std::index_sequence<I...>) {
489 return GpuArray<typename GpuTupleElement<0,T>::type, sizeof...(I)>{amrex::get<I>(tup)...};
490 }
491}
493
494template <typename T>
495AMREX_GPU_HOST_DEVICE constexpr
496auto tupleToArray (GpuTuple<T> const& tup)
497{
498 return GpuArray<T,1>{amrex::get<0>(tup)};
499}
500
502template <typename T, typename T2, typename... Ts>
503requires (Same<T,T2,Ts...>::value)
504AMREX_GPU_HOST_DEVICE constexpr
505auto tupleToArray (GpuTuple<T,T2,Ts...> const& tup)
506{
507 return detail::tuple_to_array_helper(tup, std::index_sequence_for<T,T2,Ts...>{});
508}
509
510} // namespace amrex
511
512// Spcialize std::tuple_size for GpuTuple. Used by structured bindings.
513template<typename... Ts>
514struct std::tuple_size<amrex::GpuTuple<Ts...>> {
515 static constexpr std::size_t value = sizeof...(Ts);
516};
517
518// Spcialize std::tuple_element for GpuTuple. Used by structured bindings.
519template<typename T, typename... Ts>
520struct std::tuple_element<std::size_t{0}, amrex::GpuTuple<T, Ts...>> {
521 using type = T;
522};
523
524template<std::size_t s, typename T, typename... Ts>
525struct std::tuple_element<s, amrex::GpuTuple<T, Ts...>> {
526 using type = typename std::tuple_element_t<s-1, amrex::GpuTuple<Ts...>>;
527};
528
529#endif /*AMREX_TUPLE_H_*/
AMReX_Array.H
AMReX_Functional.H
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:98
amrex::GpuTuple::operator=
__host__ __device__ GpuTuple< Ts... > & operator=(GpuTuple< Us... > const &rhs)
Definition AMReX_Tuple.H:237
amrex::GpuTuple::GpuTuple
__host__ __device__ constexpr GpuTuple()=default
amrex::GpuTuple::GpuTuple
constexpr GpuTuple(Ts const &... args)
Definition AMReX_Tuple.H:103
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:479
amrex::Tie
__host__ __device__ constexpr GpuTuple< Args &... > Tie(Args &... args) noexcept
Definition AMReX_Tuple.H:455
amrex::makeTuple
__host__ __device__ constexpr GpuTuple< detail::tuple_decay_t< Ts >... > makeTuple(Ts &&... args)
Definition AMReX_Tuple.H:269
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:394
amrex::ForwardAsTuple
__host__ __device__ constexpr GpuTuple< Ts &&... > ForwardAsTuple(Ts &&... args) noexcept
Definition AMReX_Tuple.H:465
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:496
amrex::TupleCat
__host__ __device__ constexpr auto TupleCat(TP &&a) -> typename detail::tuple_cat_result< detail::tuple_decay_t< TP > >::type
Definition AMReX_Tuple.H:306
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:43
amrex::GpuTupleElement
Definition AMReX_Tuple.H:135
amrex::GpuTupleSize
Definition AMReX_Tuple.H:127
std::tuple_element< s, amrex::GpuTuple< T, Ts... > >::type
typename std::tuple_element_t< s-1, amrex::GpuTuple< Ts... > > type
Definition AMReX_Tuple.H:526
std::tuple_element< std::size_t{0}, amrex::GpuTuple< T, Ts... > >::type
T type
Definition AMReX_Tuple.H:521