4#include <AMReX_Config.H>
17 template <
class... Ts>
18 using Tuple = std::tuple<Ts...>;
26template <std::
size_t I,
typename T>
27struct gpu_tuple_element
29 template <
typename U=T, std::enable_if_t<std::is_default_constructible_v<U>,
int> = 0>
31 constexpr gpu_tuple_element () {}
33 explicit constexpr gpu_tuple_element (T
const& a_value)
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)
39 : m_value(std::forward<U>(a_value))
45template <std::size_t I,
typename... Ts>
struct gpu_tuple_impl;
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>
52 template<
typename U=Head, std::enable_if_t<std::is_default_constructible_v<U>,
int> = 0>
54 constexpr gpu_tuple_impl () {}
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)
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))
68template <std::
size_t I,
typename Head>
69struct gpu_tuple_impl<I, Head>
70 :
public gpu_tuple_element<I, Head>
73 template<
typename U=Head, std::enable_if_t<std::is_default_constructible_v<U>,
int> = 0>
75 constexpr gpu_tuple_impl () {}
77 explicit constexpr gpu_tuple_impl (Head
const& a_head)
78 : gpu_tuple_element<I, Head>(a_head)
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)
83 : gpu_tuple_element<I, Head>(std::forward<U>(a_head))
95template <
typename... Ts>
97 :
public detail::gpu_tuple_impl<0, Ts...>
104 : detail::gpu_tuple_impl<0, Ts...>(args...)
107 template <
typename... Us, std::enable_if_t<
sizeof...(Us) ==
sizeof...(Ts),
int> = 0>
109 : detail::gpu_tuple_impl<0, Ts...>(std::forward<Us>(args)...)
112 template <
typename... Us, std::enable_if_t<
sizeof...(Us) ==
sizeof...(Ts),
int> = 0>
117 template <
typename... Us, std::enable_if_t<
sizeof...(Us) ==
sizeof...(Ts),
int> = 0>
127template <
typename... Ts>
129 :
public std::integral_constant<std::size_t, sizeof...(Ts)> {};
135template <std::size_t I,
typename Head,
typename... Tail>
139template <
typename Head,
typename... Tail>
149template <std::size_t I,
typename... Ts>
153get_impl (detail::gpu_tuple_element
159template <std::size_t I,
typename... Ts>
162typename GpuTupleElement<I, GpuTuple<Ts...> >::type
const&
163get_impl (detail::gpu_tuple_element
164 <I,
typename GpuTupleElement<I, GpuTuple<Ts...> >::type>
const& te)
noexcept
169template <std::size_t I,
typename... Ts>
172typename GpuTupleElement<I, GpuTuple<Ts...> >::type &&
173get_impl (detail::gpu_tuple_element
174 <I,
typename GpuTupleElement<I, GpuTuple<Ts...> >::type> && te)
noexcept
176 return std::move(te).m_value;
182template <std::size_t I,
typename... Ts>
185typename GpuTupleElement<I, GpuTuple<Ts...> >::type&
188 return detail::get_impl<I,Ts...>(tup);
191template <std::size_t I,
typename... Ts>
194typename GpuTupleElement<I, GpuTuple<Ts...> >::type
const&
197 return detail::get_impl<I,Ts...>(tup);
200template <std::size_t I,
typename... Ts>
203typename GpuTupleElement<I, GpuTuple<Ts...> >::type &&
206 return detail::get_impl<I,Ts...>(std::move(tup));
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)
216 (amrex::get<I>(a) = amrex::get<I>(std::forward<TP2>(b)),
217 tuple_copy<I+1,N>(a,std::forward<TP2>(b)));
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)
225 amrex::get<I>(a) = amrex::get<I>(std::forward<TP2>(b));
230template <
typename... Ts>
231template <
typename... Us, std::enable_if_t<
sizeof...(Us) ==
sizeof...(Ts),
int> >
235 detail::tuple_copy<0,
sizeof...(Ts)>(*
this, rhs);
239template <
typename... Ts>
240template <
typename... Us, std::enable_if_t<
sizeof...(Us) ==
sizeof...(Ts),
int> >
244 detail::tuple_copy<0,
sizeof...(Ts)>(*
this, std::move(rhs));
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;
259template <
typename... Ts>
262GpuTuple<detail::tuple_decay_t<Ts>...>
270 template <
typename...>
struct tuple_cat_result {};
272 template <
typename... Ts>
273 struct tuple_cat_result<GpuTuple<Ts...> >
275 using type = GpuTuple<Ts...>;
278 template <
typename... T1s,
typename... T2s,
typename... TPs>
279 struct tuple_cat_result<GpuTuple<T1s...>,GpuTuple<T2s...>,TPs...>
281 using type =
typename tuple_cat_result<GpuTuple<T1s..., T2s...>, TPs...>::type;
284 template <
typename R,
typename TP1,
typename TP2, std::size_t... N1, std::size_t... N2>
286 make_tuple (TP1
const& a, TP2
const& b,
287 std::index_sequence<N1...>
const& , std::index_sequence<N2...>
const& )
289 return R(amrex::get<N1>(a)..., amrex::get<N2>(b)...);
296template <
typename TP>
299TupleCat (TP && a) ->
typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type
301 using ReturnType =
typename detail::tuple_cat_result<detail::tuple_decay_t<TP> >::type;
302 return ReturnType(std::forward<TP>(a));
305template <
typename TP1,
typename TP2>
308TupleCat (TP1 && a, TP2 && b) ->
typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
309 detail::tuple_decay_t<TP2> >::type
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>());
319template <
typename TP1,
typename TP2,
typename... TPs>
323 ->
typename detail::tuple_cat_result<detail::tuple_decay_t<TP1>,
324 detail::tuple_decay_t<TP2>,
325 detail::tuple_decay_t<TPs>...>::type
328 std::forward<TPs>(args)...);
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...};
345 template<std::
size_t J>
347 static constexpr std::size_t get_exclusive_sum () noexcept {
348 std::size_t arr[
sizeof...(Is)] = {Is...};
350 for (std::size_t k=0; k<J; ++k) {
357 template <std::size_t start,
typename... Args, std::size_t... Is>
360 GetSubTuple (
const GpuTuple<Args...>& tup, std::index_sequence<Is...>)
noexcept
362 return makeTuple(amrex::get<start+Is>(tup)...);
365 template <
typename... Args, std::size_t... Is,
typename SIL>
368 TupleSplitImp (
const GpuTuple<Args...>& tup, std::index_sequence<Is...>, SIL)
noexcept
371 GetSubTuple<(SIL::template get_exclusive_sum<Is>())>(
373 std::make_index_sequence<SIL::template get<Is>()>()
384template <std::size_t... Is,
typename... Args>
389 static_assert((0 + ... + Is) ==
sizeof...(Args),
"Incorrect total size in TupleSplit");
390 return detail::TupleSplitImp(
392 std::make_index_sequence<
sizeof...(Is)>(),
393 detail::SplitIndexList<Is...>()
402 template <
typename F,
typename... Args>
404 auto INVOKE (F&& f, Args&&... args) ->
decltype(f(std::forward<Args>(args)...));
406 template <
typename V,
typename F,
typename... Args>
struct invoke_result {};
408 template <
typename F,
typename... Args>
409 struct invoke_result<decltype(void(INVOKE(std::declval<F>(), std::declval<Args>()...))),
412 using type =
decltype(INVOKE(std::declval<F>(), std::declval<Args>()...));
415 template <
typename F,
typename...>
struct apply_result {};
417 template <
typename F,
typename... Ts>
418 struct apply_result<
F, GpuTuple<Ts...> >
420 using type =
typename invoke_result<void,
F, Ts...>::type;
423 template <
typename F,
typename TP, std::size_t... N>
426 apply_impl (F&& f, TP&& t, std::index_sequence<N...> )
427 ->
typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
429 return std::forward<F>(f)(amrex::get<N>(std::forward<TP>(t))...);
434template <
typename F,
typename TP>
437Apply (
F&& f, TP&& t) ->
typename detail::apply_result<F,detail::tuple_decay_t<TP> >::type
439 return detail::apply_impl(std::forward<F>(f), std::forward<TP>(t),
440 std::make_index_sequence<
GpuTupleSize<std::decay_t<TP>>::value>());
445template <
typename... Args>
447constexpr GpuTuple<Args&...>
448Tie (Args&... args)
noexcept
455template <
typename... Ts>
457constexpr GpuTuple<Ts&&...>
460 return GpuTuple<Ts&&...>(std::forward<Ts>(args)...);
469template <
typename... Ts>
471constexpr GpuTuple<Ts...>
474 return GpuTuple<Ts...>(
static_cast<Ts
>(0)...);
479 template <
typename T, std::size_t... I>
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)...};
495template <
typename T,
typename T2,
typename... Ts, std::enable_if_t<Same<T,T2,Ts...>::value,
int> = 0>
499 return detail::tuple_to_array_helper(tup, std::index_sequence_for<T,T2,Ts...>{});
505template<
typename... Ts>
506struct std::tuple_size<
amrex::GpuTuple<Ts...>> {
507 static constexpr std::size_t value =
sizeof...(Ts);
511template<
typename T,
typename... Ts>
516template<std::size_t s,
typename T,
typename... Ts>
517struct std::tuple_element<s,
amrex::GpuTuple<T, Ts...>> {
#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
Head type
Definition AMReX_Tuple.H:141
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
T type
Definition AMReX_Tuple.H:513