/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/tuple.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/tuple.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/tuple.hpp Source File
tuple.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
8 #include "ck/utility/type.hpp"
10 
11 namespace ck {
12 
13 namespace detail {
14 
15 template <index_t>
17 {
18  __host__ __device__ constexpr TupleElementKey() = default;
19 };
20 
21 template <typename Key, typename Data>
23 {
24  using DataType = Data;
25 
26 #if 0 // workaround compiler complaint about implicitly-deleted default constructor
27  __host__ __device__ constexpr TupleElementKeyData() = default;
28 #else
29  __host__ __device__ constexpr TupleElementKeyData() : mData{} {}
30 #endif
31 
32  template <typename T,
34  bool>::type = false>
35  __host__ __device__ constexpr TupleElementKeyData(T&& v) : mData(ck::forward<T>(v))
36  {
37  }
38 
40 };
41 
42 // for read access of tuple element
43 template <typename Key, typename Data>
44 __host__ __device__ constexpr const Data&
46 {
47  return static_cast<const Data&>(x.mData);
48 }
49 
50 // for write access of tuple element
51 template <typename Key, typename Data>
52 __host__ __device__ constexpr Data&
54 {
55  return x.mData;
56 }
57 
58 // TODO: not sure the use of reference is correct
59 template <typename Key, typename Data>
60 __host__ __device__ constexpr Data&&
62 {
63  return static_cast<Data&&>(x.mData);
64 }
65 
66 // for infering type of tuple element
67 template <typename Key, typename Data>
68 __host__ __device__ constexpr Data get_tuple_element_data(const TupleElementKeyData<Key, Data>& x)
69 {
70  return ck::forward(x.mData);
71 }
72 
73 template <typename Indices, typename... Xs>
74 struct TupleImpl;
75 
76 template <index_t... Is, typename... Xs>
77 struct TupleImpl<Sequence<Is...>, Xs...> : TupleElementKeyData<TupleElementKey<Is>, Xs>...
78 {
79  __host__ __device__ constexpr TupleImpl() = default;
80 
81  template <typename Y,
82  typename enable_if<sizeof...(Is) == 1 && sizeof...(Xs) == 1 &&
84  bool>::type = false>
85  __host__ __device__ constexpr TupleImpl(Y&& y)
86  : TupleElementKeyData<TupleElementKey<Is>, Xs>(ck::forward<Y>(y))...
87  {
88  }
89 
90  template <typename... Ys, typename enable_if<sizeof...(Ys) >= 2, bool>::type = false>
91  __host__ __device__ constexpr TupleImpl(Ys&&... ys)
92  : TupleElementKeyData<TupleElementKey<Is>, Xs>(ck::forward<Ys>(ys))...
93  {
94  static_assert(sizeof...(Is) == sizeof...(Xs) && sizeof...(Is) == sizeof...(Ys),
95  "wrong! inconsistent size");
96  }
97 
98  __host__ __device__ static constexpr index_t Size() { return sizeof...(Xs); }
99 
100  template <index_t I>
101  __host__ __device__ constexpr const auto& GetElementDataByKey(TupleElementKey<I>) const
102  {
103  return get_tuple_element_data_reference<TupleElementKey<I>>(*this);
104  }
105 
106  template <index_t I>
107  __host__ __device__ constexpr auto& GetElementDataByKey(TupleElementKey<I>)
108  {
109  return get_tuple_element_data_reference<TupleElementKey<I>>(*this);
110  }
111 };
112 
113 } // namespace detail
114 
115 template <typename... Xs>
116 struct Tuple : detail::TupleImpl<typename arithmetic_sequence_gen<0, sizeof...(Xs), 1>::type, Xs...>
117 {
118  using base =
119  detail::TupleImpl<typename arithmetic_sequence_gen<0, sizeof...(Xs), 1>::type, Xs...>;
120 
121  __host__ __device__ constexpr Tuple() = default;
122 
123  template <typename Y,
124  typename enable_if<sizeof...(Xs) == 1 && !is_same<remove_cvref_t<Y>, Tuple>::value,
125  bool>::type = false>
126  __host__ __device__ constexpr Tuple(Y&& y) : base(ck::forward<Y>(y))
127  {
128  }
129 
130  template <typename... Ys,
131  typename enable_if<sizeof...(Ys) == sizeof...(Xs) && sizeof...(Ys) >= 2, bool>::type =
132  false>
133  __host__ __device__ constexpr Tuple(Ys&&... ys) : base(ck::forward<Ys>(ys)...)
134  {
135  }
136 
137  __host__ __device__ static constexpr index_t Size() { return sizeof...(Xs); }
138 
139  // read access
140  template <index_t I>
141  __host__ __device__ constexpr const auto& At(Number<I>) const
142  {
143  static_assert(I < base::Size(), "wrong! out of range");
144  return base::GetElementDataByKey(detail::TupleElementKey<I>{});
145  }
146 
147  // write access
148  template <index_t I>
149  __host__ __device__ constexpr auto& At(Number<I>)
150  {
151  static_assert(I < base::Size(), "wrong! out of range");
152  return base::GetElementDataByKey(detail::TupleElementKey<I>{});
153  }
154 
155  // read access
156  template <index_t I>
157  __host__ __device__ constexpr const auto& operator[](Number<I> i) const
158  {
159  return At(i);
160  }
161 
162  // write access
163  template <index_t I>
164  __host__ __device__ constexpr auto& operator()(Number<I> i)
165  {
166  return At(i);
167  }
168 
169  template <typename T>
170  __host__ __device__ constexpr auto operator=(const T& a)
171  {
172  static_assert(T::Size() == Size(), "wrong! size not the same");
173 
174  static_for<0, Size(), 1>{}([&](auto i) { operator()(i) = a[i]; });
175 
176  return *this;
177  }
178 
179  __host__ __device__ static constexpr bool IsStaticBuffer() { return true; }
180 
181  __host__ __device__ static constexpr bool IsTuple() { return true; }
182 };
183 
184 template <>
185 struct Tuple<>
186 {
187  __host__ __device__ constexpr Tuple() = default;
188 
189  __host__ __device__ static constexpr index_t Size() { return 0; }
190 
191  template <typename T>
192  __host__ __device__ constexpr auto operator=(const T&)
193  {
194  return *this;
195  }
196 
197  __host__ __device__ static constexpr bool IsStaticBuffer() { return true; }
198 };
199 
200 template <index_t I, typename TTuple>
202 {
203  // type should keep the cv/ref qualifier of original tuple element
205 };
206 
207 template <index_t I, typename TTuple>
209 
210 template <typename... Xs>
211 __host__ __device__ constexpr auto make_tuple(Xs&&... xs)
212 {
213  return Tuple<remove_cvref_t<Xs>...>(ck::forward<Xs>(xs)...);
214 }
215 
216 // https://en.cppreference.com/w/cpp/utility/tuple/tie
217 template <typename... Args>
218 constexpr Tuple<Args&...> tie(Args&... args) noexcept
219 {
220  return {args...};
221 }
222 
223 } // namespace ck
__host__ constexpr __device__ Data get_tuple_element_data(const TupleElementKeyData< Key, Data > &x)
Definition: tuple.hpp:68
__host__ constexpr __device__ const Data & get_tuple_element_data_reference(const TupleElementKeyData< Key, Data > &x)
Definition: tuple.hpp:45
Definition: ck.hpp:267
typename tuple_element< I, TTuple >::type tuple_element_t
Definition: tuple.hpp:208
constexpr Tuple< Args &... > tie(Args &... args) noexcept
Definition: tuple.hpp:218
std::enable_if< B, T > enable_if
Definition: enable_if.hpp:24
__host__ constexpr __device__ auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:211
remove_cv_t< remove_reference_t< T > > remove_cvref_t
Definition: type.hpp:297
int32_t index_t
Definition: ck.hpp:298
const GenericPointer< typename T::ValueType > T2 value
Definition: pointer.h:1350
const GenericPointer< typename T::ValueType > T2 T::AllocatorType & a
Definition: pointer.h:1249
Definition: sequence.hpp:43
__host__ static constexpr __device__ index_t Size()
Definition: tuple.hpp:189
__host__ constexpr __device__ auto operator=(const T &)
Definition: tuple.hpp:192
__host__ static constexpr __device__ bool IsStaticBuffer()
Definition: tuple.hpp:197
__host__ constexpr __device__ Tuple()=default
Definition: tuple.hpp:117
__host__ constexpr __device__ Tuple(Y &&y)
Definition: tuple.hpp:126
detail::TupleImpl< typename arithmetic_sequence_gen< 0, sizeof...(Xs), 1 >::type, Xs... > base
Definition: tuple.hpp:119
__host__ constexpr __device__ Tuple()=default
Definition: sequence.hpp:256
Definition: tuple.hpp:23
__host__ constexpr __device__ TupleElementKeyData(T &&v)
Definition: tuple.hpp:35
Data DataType
Definition: tuple.hpp:24
DataType mData
Definition: tuple.hpp:39
__host__ constexpr __device__ TupleElementKeyData()
Definition: tuple.hpp:29
Definition: tuple.hpp:17
__host__ constexpr __device__ TupleElementKey()=default
__host__ constexpr __device__ auto & GetElementDataByKey(TupleElementKey< I >)
Definition: tuple.hpp:107
__host__ constexpr __device__ TupleImpl()=default
__host__ constexpr __device__ TupleImpl(Y &&y)
Definition: tuple.hpp:85
__host__ static constexpr __device__ index_t Size()
Definition: tuple.hpp:98
__host__ constexpr __device__ const auto & GetElementDataByKey(TupleElementKey< I >) const
Definition: tuple.hpp:101
Definition: tuple.hpp:74
Definition: type.hpp:177
Definition: tuple.hpp:202
decltype(detail::get_tuple_element_data< detail::TupleElementKey< I > >(TTuple{})) type
Definition: tuple.hpp:204