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>
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 template <typename U=T, std::enable_if_t<std::is_default_constructible_v<U>,int> = 0>
31 constexpr gpu_tuple_element () {} // NOLINT
32
33 explicit constexpr gpu_tuple_element (T const& a_value)
34 : m_value(a_value)
35 {}
36
37 template <typename U, std::enable_if_t<std::is_convertible_v<U&&,T>,int> = 0>
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 template<typename U=Head, std::enable_if_t<std::is_default_constructible_v<U>,int> = 0>
54 constexpr gpu_tuple_impl () {} // NOLINT
55
56 constexpr gpu_tuple_impl (Head const& a_head, Tail const&... a_tail)
57 : gpu_tuple_impl<I+1, Tail...>(a_tail...),
58 gpu_tuple_element<I, Head>(a_head)
59 {}
60
61 template <typename UH, typename... UT, std::enable_if_t<std::is_convertible_v<UH&&,Head>,int> = 0>
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 template<typename U=Head, std::enable_if_t<std::is_default_constructible_v<U>,int> = 0>
75 constexpr gpu_tuple_impl () {} // NOLINT
76
77 explicit constexpr gpu_tuple_impl (Head const& a_head)
78 : gpu_tuple_element<I, Head>(a_head)
79 {}
80
81 template <typename U, std::enable_if_t<std::is_convertible_v<U&&,Head>,int> = 0>
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>
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, std::enable_if_t<sizeof...(Us) == sizeof...(Ts),int> = 0>
108 constexpr GpuTuple (Us&&... args)
109 : detail::gpu_tuple_impl<0, Ts...>(std::forward<Us>(args)...)
110 {}
111
112 template <typename... Us, std::enable_if_t<sizeof...(Us) == sizeof...(Ts),int> = 0>
114 inline GpuTuple<Ts...>&
116
117 template <typename... Us, std::enable_if_t<sizeof...(Us) == sizeof...(Ts),int> = 0>
119 inline GpuTuple<Ts...>&
121};
122
123// GpuTupleSize
124
125template <typename T> struct GpuTupleSize;
126
127template <typename... Ts>
128struct GpuTupleSize<GpuTuple<Ts...> >
129 : public std::integral_constant<std::size_t, sizeof...(Ts)> {};
130
131// GpuTupleElement
132
133template <std::size_t I, typename T> struct GpuTupleElement;
134
135template <std::size_t I, typename Head, typename... Tail>
136struct GpuTupleElement<I, GpuTuple<Head, Tail...> >
137 : GpuTupleElement<I-1, GpuTuple<Tail...> > {};
138
139template <typename Head, typename... Tail>
140struct GpuTupleElement<0, GpuTuple<Head, Tail...> > {
141 using type = Head;
142};
143
144// get
145
147namespace detail {
148
149template <std::size_t I, typename... Ts>
151constexpr
152typename GpuTupleElement<I, GpuTuple<Ts...> >::type&
153get_impl (detail::gpu_tuple_element
154 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type>& te) noexcept
155{
156 return te.m_value;
157}
158
159template <std::size_t I, typename... Ts>
161constexpr
162typename GpuTupleElement<I, GpuTuple<Ts...> >::type const&
163get_impl (detail::gpu_tuple_element
164 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type> const& te) noexcept
165{
166 return te.m_value;
167}
168
169template <std::size_t I, typename... Ts>
171constexpr
172typename GpuTupleElement<I, GpuTuple<Ts...> >::type &&
173get_impl (detail::gpu_tuple_element
174 <I, typename GpuTupleElement<I, GpuTuple<Ts...> >::type> && te) noexcept
175{
176 return std::move(te).m_value;
177}
178
179} // detail
181
182template <std::size_t I, typename... Ts>
184constexpr
185typename GpuTupleElement<I, GpuTuple<Ts...> >::type&
186get (GpuTuple<Ts...>& tup) noexcept
187{
188 return detail::get_impl<I,Ts...>(tup);
189}
190
191template <std::size_t I, typename... Ts>
193constexpr
194typename GpuTupleElement<I, GpuTuple<Ts...> >::type const&
195get (GpuTuple<Ts...> const& tup) noexcept
196{
197 return detail::get_impl<I,Ts...>(tup);
198}
199
200template <std::size_t I, typename... Ts>
202constexpr
203typename GpuTupleElement<I, GpuTuple<Ts...> >::type &&
204get (GpuTuple<Ts...> && tup) noexcept
205{
206 return detail::get_impl<I,Ts...>(std::move(tup));
207}
208
210namespace detail {
211 template <std::size_t I, std::size_t N, typename TP1, typename TP2>
213 std::enable_if_t<(I<N-1),void>
214 tuple_copy (TP1 & a, TP2 && b)
215 {
216 (amrex::get<I>(a) = amrex::get<I>(std::forward<TP2>(b)),
217 tuple_copy<I+1,N>(a,std::forward<TP2>(b)));
218 }
219
220 template <std::size_t I, std::size_t N, typename TP1, typename TP2>
222 std::enable_if_t<I==N-1,void>
223 tuple_copy (TP1 & a, TP2 && b)
224 {
225 amrex::get<I>(a) = amrex::get<I>(std::forward<TP2>(b));
226 }
227}
229
230template <typename... Ts>
231template <typename... Us, std::enable_if_t<sizeof...(Us) == sizeof...(Ts),int> >
232AMREX_GPU_HOST_DEVICE inline GpuTuple<Ts...>&
234{
235 detail::tuple_copy<0,sizeof...(Ts)>(*this, rhs);
236 return *this;
237}
238
239template <typename... Ts>
240template <typename... Us, std::enable_if_t<sizeof...(Us) == sizeof...(Ts),int> >
241AMREX_GPU_HOST_DEVICE inline GpuTuple<Ts...>&
243{
244 detail::tuple_copy<0,sizeof...(Ts)>(*this, std::move(rhs));
245 return *this;
246}
247
248// makeTuple
249
251namespace detail {
252 template <typename T> struct unwrap { using type = T; };
253 template <typename T> struct unwrap<std::reference_wrapper<T> > { using type = T&; };
254 template <typename T>
255 using tuple_decay_t = typename unwrap<std::decay_t<T>>::type;
256}
258
259template <typename... Ts>
261constexpr
262GpuTuple<detail::tuple_decay_t<Ts>...>
263makeTuple (Ts &&... args)
264{
265 return GpuTuple<detail::tuple_decay_t<Ts>...>(std::forward<Ts>(args)...);
266}
267
269namespace detail {
270 template <typename...> struct tuple_cat_result {};
271
272 template <typename... Ts>
273 struct tuple_cat_result<GpuTuple<Ts...> >
274 {
275 using type = GpuTuple<Ts...>;
276 };
277
278 template <typename... T1s, typename... T2s, typename... TPs>
279 struct tuple_cat_result<GpuTuple<T1s...>,GpuTuple<T2s...>,TPs...>
280 {
281 using type = typename tuple_cat_result<GpuTuple<T1s..., T2s...>, TPs...>::type;
282 };
283
284 // NOLINTNEXTLINE(misc-confusable-identifiers)
285 template <typename R, typename TP1, typename TP2, std::size_t... N1, std::size_t... N2>
286 AMREX_GPU_HOST_DEVICE constexpr R
287 make_tuple (TP1 const& a, TP2 const& b,
288 std::index_sequence<N1...> const& /*n1*/, std::index_sequence<N2...> const& /*n2*/)
289 {
290 return R(amrex::get<N1>(a)..., amrex::get<N2>(b)...);
291 }
292}
294
295// TupleCat
296
297template <typename TP>
299constexpr auto
300TupleCat (TP && a) -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type
301{
302 using ReturnType = typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type;
303 return ReturnType(std::forward<TP>(a));
304}
305
306template <typename TP1, typename TP2>
308constexpr auto
309TupleCat (TP1 && a, TP2 && b) -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
310 detail::tuple_decay_t<TP2> >::type
311{
312 using ReturnType = typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
313 detail::tuple_decay_t<TP2> >::type;
314 return detail::make_tuple<ReturnType>
315 (std::forward<TP1>(a), std::forward<TP2>(b),
316 std::make_index_sequence<GpuTupleSize<std::decay_t<TP1>>::value>(),
317 std::make_index_sequence<GpuTupleSize<std::decay_t<TP2>>::value>());
318}
319
320template <typename TP1, typename TP2, typename... TPs>
322constexpr auto
323TupleCat (TP1&& a, TP2&& b, TPs&&... args)
324 -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
325 detail::tuple_decay_t<TP2>,
326 detail::tuple_decay_t<TPs>...>::type
327{
328 return TupleCat(TupleCat(std::forward<TP1>(a),std::forward<TP2>(b)),
329 std::forward<TPs>(args)...);
330}
331
332// TupleSplit
333
335namespace detail {
336
337 template<std::size_t...Is>
338 struct SplitIndexList {
339 template<std::size_t J>
341 static constexpr std::size_t get () noexcept {
342 std::size_t arr[sizeof...(Is)] = {Is...};
343 return arr[J];
344 }
345
346 template<std::size_t J>
348 static constexpr std::size_t get_exclusive_sum () noexcept {
349 std::size_t arr[sizeof...(Is)] = {Is...};
350 std::size_t sum = 0;
351 for (std::size_t k=0; k<J; ++k) {
352 sum += arr[k];
353 }
354 return sum;
355 }
356 };
357
358 template <std::size_t start, typename... Args, std::size_t... Is>
360 constexpr auto
361 GetSubTuple (const GpuTuple<Args...>& tup, std::index_sequence<Is...>) noexcept
362 {
363 return makeTuple(amrex::get<start+Is>(tup)...);
364 }
365
366 template <typename... Args, std::size_t... Is, typename SIL>
368 constexpr auto
369 TupleSplitImp (const GpuTuple<Args...>& tup, std::index_sequence<Is...>, SIL) noexcept
370 {
371 return makeTuple(
372 GetSubTuple<(SIL::template get_exclusive_sum<Is>())>(
373 tup,
374 std::make_index_sequence<SIL::template get<Is>()>()
375 )...
376 );
377 }
378}
380
385template <std::size_t... Is, typename... Args>
387constexpr auto
388TupleSplit (const GpuTuple<Args...>& tup) noexcept
389{
390 static_assert((0 + ... + Is) == sizeof...(Args), "Incorrect total size in TupleSplit");
391 return detail::TupleSplitImp(
392 tup,
393 std::make_index_sequence<sizeof...(Is)>(),
394 detail::SplitIndexList<Is...>()
395 );
396}
397
398// Apply
399
401namespace detail {
402
403 template <typename F, typename... Args>
405 auto INVOKE (F&& f, Args&&... args) -> decltype(f(std::forward<Args>(args)...));
406
407 template <typename V, typename F, typename... Args> struct invoke_result {};
408
409 template <typename F, typename... Args>
410 struct invoke_result<decltype(void(INVOKE(std::declval<F>(), std::declval<Args>()...))),
411 F, Args...>
412 {
413 using type = decltype(INVOKE(std::declval<F>(), std::declval<Args>()...));
414 };
415
416 template <typename F, typename...> struct apply_result {};
417
418 template <typename F, typename... Ts>
419 struct apply_result<F, GpuTuple<Ts...> >
420 {
421 using type = typename invoke_result<void, F, Ts...>::type;
422 };
423
424 template <typename F, typename TP, std::size_t... N>
426 constexpr auto
427 apply_impl (F&& f, TP&& t, std::index_sequence<N...> /*is*/)
428 -> typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
429 {
430 return std::forward<F>(f)(amrex::get<N>(std::forward<TP>(t))...);
431 }
432}
434
435template <typename F, typename TP>
437constexpr auto
438Apply (F&& f, TP&& t) -> typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
439{
440 return detail::apply_impl(std::forward<F>(f), std::forward<TP>(t),
441 std::make_index_sequence<GpuTupleSize<std::decay_t<TP>>::value>());
442}
443
444// Tie
445
446template <typename... Args>
448constexpr GpuTuple<Args&...>
449Tie (Args&... args) noexcept
450{
451 return GpuTuple<Args&...>(args...);
452}
453
454// ForwardAsTuple
455
456template <typename... Ts>
458constexpr GpuTuple<Ts&&...>
459ForwardAsTuple (Ts&&... args) noexcept
460{
461 return GpuTuple<Ts&&...>(std::forward<Ts>(args)...);
462}
463
464// MakeZeroTuple
465
470template <typename... Ts>
472constexpr GpuTuple<Ts...>
474{
475 return GpuTuple<Ts...>(static_cast<Ts>(0)...);
476}
477
479namespace detail {
480 template <typename T, std::size_t... I>
481 AMREX_GPU_HOST_DEVICE constexpr
482 auto tuple_to_array_helper (T const& tup, std::index_sequence<I...>) {
483 return GpuArray<typename GpuTupleElement<0,T>::type, sizeof...(I)>{amrex::get<I>(tup)...};
484 }
485}
487
488template <typename T>
489AMREX_GPU_HOST_DEVICE constexpr
490auto tupleToArray (GpuTuple<T> const& tup)
491{
492 return GpuArray<T,1>{amrex::get<0>(tup)};
493}
494
496template <typename T, typename T2, typename... Ts, std::enable_if_t<Same<T,T2,Ts...>::value, int> = 0>
497AMREX_GPU_HOST_DEVICE constexpr
499{
500 return detail::tuple_to_array_helper(tup, std::index_sequence_for<T,T2,Ts...>{});
501}
502
503} // namespace amrex
504
505// Spcialize std::tuple_size for GpuTuple. Used by structured bindings.
506template<typename... Ts>
507struct std::tuple_size<amrex::GpuTuple<Ts...>> {
508 static constexpr std::size_t value = sizeof...(Ts);
509};
510
511// Spcialize std::tuple_element for GpuTuple. Used by structured bindings.
512template<typename T, typename... Ts>
513struct std::tuple_element<std::size_t{0}, amrex::GpuTuple<T, Ts...>> {
514 using type = T;
515};
516
517template<std::size_t s, typename T, typename... Ts>
518struct std::tuple_element<s, amrex::GpuTuple<T, Ts...>> {
519 using type = typename std::tuple_element_t<s-1, amrex::GpuTuple<Ts...>>;
520};
521
522#endif /*AMREX_TUPLE_H_*/
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
GPU-compatible tuple.
Definition AMReX_Tuple.H:98
__host__ __device__ GpuTuple< Ts... > & operator=(GpuTuple< Us... > const &rhs)
Definition AMReX_Tuple.H:233
constexpr GpuTuple(Us &&... args)
Definition AMReX_Tuple.H:108
__host__ __device__ constexpr GpuTuple()=default
constexpr GpuTuple(Ts const &... args)
Definition AMReX_Tuple.H:103
Definition AMReX_Amr.cpp:49
std::tuple< Ts... > Tuple
Definition AMReX_Tuple.H:18
__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:473
__host__ __device__ constexpr GpuTuple< Args &... > Tie(Args &... args) noexcept
Definition AMReX_Tuple.H:449
__host__ __device__ constexpr GpuTuple< detail::tuple_decay_t< Ts >... > makeTuple(Ts &&... args)
Definition AMReX_Tuple.H:263
__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:388
__host__ __device__ constexpr GpuTuple< Ts &&... > ForwardAsTuple(Ts &&... args) noexcept
Definition AMReX_Tuple.H:459
__host__ __device__ constexpr auto Apply(F &&f, IntVectND< dim > const &iv)
Definition AMReX_IntVect.H:1246
__host__ __device__ constexpr auto tupleToArray(GpuTuple< T > const &tup)
Definition AMReX_Tuple.H:490
__host__ __device__ constexpr auto TupleCat(TP &&a) -> typename detail::tuple_cat_result< detail::tuple_decay_t< TP > >::type
Definition AMReX_Tuple.H:300
__host__ __device__ constexpr int get(IntVectND< dim > const &iv) noexcept
Get I'th element of IntVectND<dim>
Definition AMReX_IntVect.H:1230
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:43
Definition AMReX_Tuple.H:133
Definition AMReX_Tuple.H:125
typename std::tuple_element_t< s-1, amrex::GpuTuple< Ts... > > type
Definition AMReX_Tuple.H:519