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 template <typename R, typename TP1, typename TP2, std::size_t... N1, std::size_t... N2>
285 AMREX_GPU_HOST_DEVICE constexpr R
286 make_tuple (TP1 const& a, TP2 const& b,
287 std::index_sequence<N1...> const& /*n1*/, std::index_sequence<N2...> const& /*n2*/)
288 {
289 return R(amrex::get<N1>(a)..., amrex::get<N2>(b)...);
290 }
291}
293
294// TupleCat
295
296template <typename TP>
298constexpr auto
299TupleCat (TP && a) -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type
300{
301 using ReturnType = typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type;
302 return ReturnType(std::forward<TP>(a));
303}
304
305template <typename TP1, typename TP2>
307constexpr auto
308TupleCat (TP1 && a, TP2 && b) -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
309 detail::tuple_decay_t<TP2> >::type
310{
311 using ReturnType = typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
312 detail::tuple_decay_t<TP2> >::type;
313 return detail::make_tuple<ReturnType>
314 (std::forward<TP1>(a), std::forward<TP2>(b),
315 std::make_index_sequence<GpuTupleSize<std::decay_t<TP1>>::value>(),
316 std::make_index_sequence<GpuTupleSize<std::decay_t<TP2>>::value>());
317}
318
319template <typename TP1, typename TP2, typename... TPs>
321constexpr auto
322TupleCat (TP1&& a, TP2&& b, TPs&&... args)
323 -> typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
324 detail::tuple_decay_t<TP2>,
325 detail::tuple_decay_t<TPs>...>::type
326{
327 return TupleCat(TupleCat(std::forward<TP1>(a),std::forward<TP2>(b)),
328 std::forward<TPs>(args)...);
329}
330
331// TupleSplit
332
334namespace detail {
335
336 template<std::size_t...Is>
337 struct SplitIndexList {
338 template<std::size_t J>
340 static constexpr std::size_t get () noexcept {
341 std::size_t arr[sizeof...(Is)] = {Is...};
342 return arr[J];
343 }
344
345 template<std::size_t J>
347 static constexpr std::size_t get_exclusive_sum () noexcept {
348 std::size_t arr[sizeof...(Is)] = {Is...};
349 std::size_t sum = 0;
350 for (std::size_t k=0; k<J; ++k) {
351 sum += arr[k];
352 }
353 return sum;
354 }
355 };
356
357 template <std::size_t start, typename... Args, std::size_t... Is>
359 constexpr auto
360 GetSubTuple (const GpuTuple<Args...>& tup, std::index_sequence<Is...>) noexcept
361 {
362 return makeTuple(amrex::get<start+Is>(tup)...);
363 }
364
365 template <typename... Args, std::size_t... Is, typename SIL>
367 constexpr auto
368 TupleSplitImp (const GpuTuple<Args...>& tup, std::index_sequence<Is...>, SIL) noexcept
369 {
370 return makeTuple(
371 GetSubTuple<(SIL::template get_exclusive_sum<Is>())>(
372 tup,
373 std::make_index_sequence<SIL::template get<Is>()>()
374 )...
375 );
376 }
377}
379
384template <std::size_t... Is, typename... Args>
386constexpr auto
387TupleSplit (const GpuTuple<Args...>& tup) noexcept
388{
389 static_assert((0 + ... + Is) == sizeof...(Args), "Incorrect total size in TupleSplit");
390 return detail::TupleSplitImp(
391 tup,
392 std::make_index_sequence<sizeof...(Is)>(),
393 detail::SplitIndexList<Is...>()
394 );
395}
396
397// Apply
398
400namespace detail {
401
402 template <typename F, typename... Args>
404 auto INVOKE (F&& f, Args&&... args) -> decltype(f(std::forward<Args>(args)...));
405
406 template <typename V, typename F, typename... Args> struct invoke_result {};
407
408 template <typename F, typename... Args>
409 struct invoke_result<decltype(void(INVOKE(std::declval<F>(), std::declval<Args>()...))),
410 F, Args...>
411 {
412 using type = decltype(INVOKE(std::declval<F>(), std::declval<Args>()...));
413 };
414
415 template <typename F, typename...> struct apply_result {};
416
417 template <typename F, typename... Ts>
418 struct apply_result<F, GpuTuple<Ts...> >
419 {
420 using type = typename invoke_result<void, F, Ts...>::type;
421 };
422
423 template <typename F, typename TP, std::size_t... N>
425 constexpr auto
426 apply_impl (F&& f, TP&& t, std::index_sequence<N...> /*is*/)
427 -> typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
428 {
429 return std::forward<F>(f)(amrex::get<N>(std::forward<TP>(t))...);
430 }
431}
433
434template <typename F, typename TP>
436constexpr auto
437Apply (F&& f, TP&& t) -> typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
438{
439 return detail::apply_impl(std::forward<F>(f), std::forward<TP>(t),
440 std::make_index_sequence<GpuTupleSize<std::decay_t<TP>>::value>());
441}
442
443// Tie
444
445template <typename... Args>
447constexpr GpuTuple<Args&...>
448Tie (Args&... args) noexcept
449{
450 return GpuTuple<Args&...>(args...);
451}
452
453// ForwardAsTuple
454
455template <typename... Ts>
457constexpr GpuTuple<Ts&&...>
458ForwardAsTuple (Ts&&... args) noexcept
459{
460 return GpuTuple<Ts&&...>(std::forward<Ts>(args)...);
461}
462
463// MakeZeroTuple
464
469template <typename... Ts>
471constexpr GpuTuple<Ts...>
473{
474 return GpuTuple<Ts...>(static_cast<Ts>(0)...);
475}
476
478namespace detail {
479 template <typename T, std::size_t... I>
480 AMREX_GPU_HOST_DEVICE constexpr
481 auto tuple_to_array_helper (T const& tup, std::index_sequence<I...>) {
482 return GpuArray<typename GpuTupleElement<0,T>::type, sizeof...(I)>{amrex::get<I>(tup)...};
483 }
484}
486
487template <typename T>
488AMREX_GPU_HOST_DEVICE constexpr
489auto tupleToArray (GpuTuple<T> const& tup)
490{
491 return GpuArray<T,1>{amrex::get<0>(tup)};
492}
493
495template <typename T, typename T2, typename... Ts, std::enable_if_t<Same<T,T2,Ts...>::value, int> = 0>
496AMREX_GPU_HOST_DEVICE constexpr
498{
499 return detail::tuple_to_array_helper(tup, std::index_sequence_for<T,T2,Ts...>{});
500}
501
502} // namespace amrex
503
504// Spcialize std::tuple_size for GpuTuple. Used by structured bindings.
505template<typename... Ts>
506struct std::tuple_size<amrex::GpuTuple<Ts...>> {
507 static constexpr std::size_t value = sizeof...(Ts);
508};
509
510// Spcialize std::tuple_element for GpuTuple. Used by structured bindings.
511template<typename T, typename... Ts>
512struct std::tuple_element<std::size_t{0}, amrex::GpuTuple<T, Ts...>> {
513 using type = T;
514};
515
516template<std::size_t s, typename T, typename... Ts>
517struct std::tuple_element<s, amrex::GpuTuple<T, Ts...>> {
518 using type = typename std::tuple_element<s-1, amrex::GpuTuple<Ts...>>::type;
519};
520
521#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:472
__host__ __device__ constexpr auto Apply(F &&f, TP &&t) -> typename detail::apply_result< F, detail::tuple_decay_t< TP > >::type
Definition AMReX_Tuple.H:437
__host__ __device__ constexpr GpuTuple< Args &... > Tie(Args &... args) noexcept
Definition AMReX_Tuple.H:448
__host__ __device__ constexpr GpuTuple< detail::tuple_decay_t< Ts >... > makeTuple(Ts &&... args)
Definition AMReX_Tuple.H:263
__host__ __device__ constexpr GpuTupleElement< I, GpuTuple< Ts... > >::type & get(GpuTuple< Ts... > &tup) noexcept
Definition AMReX_Tuple.H:186
__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:387
__host__ __device__ constexpr GpuTuple< Ts &&... > ForwardAsTuple(Ts &&... args) noexcept
Definition AMReX_Tuple.H:458
__host__ __device__ constexpr auto tupleToArray(GpuTuple< T > const &tup)
Definition AMReX_Tuple.H:489
__host__ __device__ constexpr auto TupleCat(TP &&a) -> typename detail::tuple_cat_result< detail::tuple_decay_t< TP > >::type
Definition AMReX_Tuple.H:299
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:40
Definition AMReX_Tuple.H:133
Definition AMReX_Tuple.H:125
typename std::tuple_element< s-1, amrex::GpuTuple< Ts... > >::type type
Definition AMReX_Tuple.H:518