tuple.hpp 5.78 KB
Newer Older
1
#pragma once
Chao Liu's avatar
Chao Liu committed
2
3

#include "integral_constant.hpp"
Chao Liu's avatar
Chao Liu committed
4
#include "sequence.hpp"
Chao Liu's avatar
Chao Liu committed
5
#include "type.hpp"
Chao Liu's avatar
Chao Liu committed
6
#include "enable_if.hpp"
Chao Liu's avatar
Chao Liu committed
7
8
9

namespace ck {

Chao Liu's avatar
Chao Liu committed
10
namespace detail {
Chao Liu's avatar
Chao Liu committed
11

Chao Liu's avatar
Chao Liu committed
12
13
14
template <index_t>
struct TupleElementKey
{
Chao Liu's avatar
Chao Liu committed
15
    __host__ __device__ constexpr TupleElementKey() = default;
Chao Liu's avatar
Chao Liu committed
16
};
Chao Liu's avatar
Chao Liu committed
17

Chao Liu's avatar
Chao Liu committed
18
template <typename Key, typename Data>
19
struct TupleElementKeyData
Chao Liu's avatar
Chao Liu committed
20
{
21
22
23
24
25
#if 0 // workaround compiler complaint about implicitly-deleted default constructor
    __host__ __device__ constexpr TupleElementKeyData() = default;
#else
    __host__ __device__ constexpr TupleElementKeyData() : mData{} {}
#endif
Chao Liu's avatar
Chao Liu committed
26

27
28
29
30
    template <typename T,
              typename enable_if<!is_same<remove_cvref_t<T>, TupleElementKeyData>::value,
                                 bool>::type = false>
    __host__ __device__ constexpr TupleElementKeyData(T&& v) : mData(std::forward<T>(v))
Chao Liu's avatar
Chao Liu committed
31
32
33
    {
    }

Chao Liu's avatar
Chao Liu committed
34
    Data mData;
Chao Liu's avatar
Chao Liu committed
35
36
};

Chao Liu's avatar
Chao Liu committed
37
template <typename Key, typename Data>
38
39
__host__ __device__ constexpr const Data&
get_tuple_element_data(const TupleElementKeyData<Key, Data>& x)
Chao Liu's avatar
Chao Liu committed
40
{
Chao Liu's avatar
Chao Liu committed
41
    return static_cast<const Data&>(x.mData);
Chao Liu's avatar
Chao Liu committed
42
}
Chao Liu's avatar
Chao Liu committed
43

Chao Liu's avatar
Chao Liu committed
44
template <typename Key, typename Data>
45
__host__ __device__ constexpr Data& get_tuple_element_data(TupleElementKeyData<Key, Data>& x)
Chao Liu's avatar
Chao Liu committed
46
{
Chao Liu's avatar
Chao Liu committed
47
48
    return x.mData;
}
Chao Liu's avatar
Chao Liu committed
49

Chao Liu's avatar
Chao Liu committed
50
// TODO: not sure the use of reference is correct
Chao Liu's avatar
Chao Liu committed
51
template <typename Key, typename Data>
52
__host__ __device__ constexpr Data&& get_tuple_element_data(TupleElementKeyData<Key, Data>&& x)
Chao Liu's avatar
Chao Liu committed
53
{
Chao Liu's avatar
Chao Liu committed
54
55
    return static_cast<Data&&>(x.mData);
}
Chao Liu's avatar
Chao Liu committed
56

Chao Liu's avatar
Chao Liu committed
57
58
59
60
template <typename Indices, typename... Xs>
struct TupleImpl;

template <index_t... Is, typename... Xs>
61
struct TupleImpl<Sequence<Is...>, Xs...> : TupleElementKeyData<TupleElementKey<Is>, Xs>...
Chao Liu's avatar
Chao Liu committed
62
{
Chao Liu's avatar
Chao Liu committed
63
64
    __host__ __device__ constexpr TupleImpl() = default;

Chao Liu's avatar
Chao Liu committed
65
66
    template <typename Y,
              typename enable_if<sizeof...(Is) == 1 && sizeof...(Xs) == 1 &&
Chao Liu's avatar
Chao Liu committed
67
                                     !is_same<remove_cvref_t<Y>, TupleImpl>::value,
Chao Liu's avatar
Chao Liu committed
68
                                 bool>::type = false>
Chao Liu's avatar
Chao Liu committed
69
    __host__ __device__ constexpr TupleImpl(Y&& y)
70
        : TupleElementKeyData<TupleElementKey<Is>, Xs>(std::forward<Y>(y))...
Chao Liu's avatar
Chao Liu committed
71
72
73
    {
    }

Chao Liu's avatar
Chao Liu committed
74
    template <typename... Ys, typename enable_if<sizeof...(Ys) >= 2, bool>::type = false>
Chao Liu's avatar
Chao Liu committed
75
    __host__ __device__ constexpr TupleImpl(Ys&&... ys)
76
        : TupleElementKeyData<TupleElementKey<Is>, Xs>(std::forward<Ys>(ys))...
Chao Liu's avatar
Chao Liu committed
77
    {
Chao Liu's avatar
Chao Liu committed
78
79
        static_assert(sizeof...(Is) == sizeof...(Xs) && sizeof...(Is) == sizeof...(Ys),
                      "wrong! inconsistent size");
Chao Liu's avatar
Chao Liu committed
80
81
82
83
84
    }

    __host__ __device__ static constexpr index_t Size() { return sizeof...(Xs); }

    template <index_t I>
85
    __host__ __device__ constexpr const auto& GetElementDataByKey(TupleElementKey<I>) const
Chao Liu's avatar
Chao Liu committed
86
    {
87
        return get_tuple_element_data<TupleElementKey<I>>(*this);
Chao Liu's avatar
Chao Liu committed
88
89
90
    }

    template <index_t I>
91
    __host__ __device__ constexpr auto& GetElementDataByKey(TupleElementKey<I>)
Chao Liu's avatar
Chao Liu committed
92
    {
93
        return get_tuple_element_data<TupleElementKey<I>>(*this);
Chao Liu's avatar
Chao Liu committed
94
    }
Chao Liu's avatar
Chao Liu committed
95
96
};

Chao Liu's avatar
Chao Liu committed
97
98
99
100
} // namespace detail

template <typename... Xs>
struct Tuple : detail::TupleImpl<typename arithmetic_sequence_gen<0, sizeof...(Xs), 1>::type, Xs...>
Chao Liu's avatar
Chao Liu committed
101
{
Chao Liu's avatar
Chao Liu committed
102
103
104
    using base =
        detail::TupleImpl<typename arithmetic_sequence_gen<0, sizeof...(Xs), 1>::type, Xs...>;

Chao Liu's avatar
Chao Liu committed
105
106
107
    __host__ __device__ constexpr Tuple() = default;

    template <typename Y,
Chao Liu's avatar
Chao Liu committed
108
              typename enable_if<sizeof...(Xs) == 1 && !is_same<remove_cvref_t<Y>, Tuple>::value,
Chao Liu's avatar
Chao Liu committed
109
                                 bool>::type = false>
Chao Liu's avatar
Chao Liu committed
110
    __host__ __device__ constexpr Tuple(Y&& y) : base(std::forward<Y>(y))
Chao Liu's avatar
Chao Liu committed
111
112
113
    {
    }

Chao Liu's avatar
Chao Liu committed
114
    template <typename... Ys,
Chao Liu's avatar
Chao Liu committed
115
116
              typename enable_if<sizeof...(Ys) == sizeof...(Xs) && sizeof...(Ys) >= 2, bool>::type =
                  false>
Chao Liu's avatar
Chao Liu committed
117
118
119
120
121
122
    __host__ __device__ constexpr Tuple(Ys&&... ys) : base(std::forward<Ys>(ys)...)
    {
    }

    __host__ __device__ static constexpr index_t Size() { return sizeof...(Xs); }

123
    // read access
Chao Liu's avatar
Chao Liu committed
124
125
126
127
    template <index_t I>
    __host__ __device__ constexpr const auto& At(Number<I>) const
    {
        static_assert(I < base::Size(), "wrong! out of range");
128
        return base::GetElementDataByKey(detail::TupleElementKey<I>{});
Chao Liu's avatar
Chao Liu committed
129
130
    }

131
    // write access
Chao Liu's avatar
Chao Liu committed
132
133
134
135
    template <index_t I>
    __host__ __device__ constexpr auto& At(Number<I>)
    {
        static_assert(I < base::Size(), "wrong! out of range");
136
        return base::GetElementDataByKey(detail::TupleElementKey<I>{});
Chao Liu's avatar
Chao Liu committed
137
    }
Chao Liu's avatar
Chao Liu committed
138

139
    // read access
Chao Liu's avatar
Chao Liu committed
140
141
142
143
144
    template <index_t I>
    __host__ __device__ constexpr const auto& operator[](Number<I> i) const
    {
        return At(i);
    }
145

146
    // write access
Chao Liu's avatar
Chao Liu committed
147
148
149
150
151
    template <index_t I>
    __host__ __device__ constexpr auto& operator()(Number<I> i)
    {
        return At(i);
    }
Chao Liu's avatar
Chao Liu committed
152

Chao Liu's avatar
Chao Liu committed
153
154
155
156
    template <typename T>
    __host__ __device__ constexpr auto operator=(const T& a)
    {
        static_assert(T::Size() == Size(), "wrong! size not the same");
Chao Liu's avatar
Chao Liu committed
157

Chao Liu's avatar
Chao Liu committed
158
        static_for<0, Size(), 1>{}([&](auto i) { operator()(i) = a[i]; });
Chao Liu's avatar
Chao Liu committed
159

Chao Liu's avatar
Chao Liu committed
160
161
        return *this;
    }
zjing14's avatar
zjing14 committed
162
163

    __host__ __device__ static constexpr bool IsStaticBuffer() { return true; }
Chao Liu's avatar
Chao Liu committed
164
};
165

166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
template <>
struct Tuple<>
{
    __host__ __device__ constexpr Tuple() = default;

    __host__ __device__ static constexpr index_t Size() { return 0; }

    template <typename T>
    __host__ __device__ constexpr auto operator=(const T&)
    {
        return *this;
    }

    __host__ __device__ static constexpr bool IsStaticBuffer() { return true; }
};

template <index_t I, typename TTuple>
struct tuple_element
{
    using type = decltype(TTuple{}.At(Number<I>{}));
};

template <index_t I, typename TTuple>
using tuple_element_t = typename tuple_element<I, TTuple>::type;

Chao Liu's avatar
Chao Liu committed
191
192
template <typename... Xs>
__host__ __device__ constexpr auto make_tuple(Xs&&... xs)
Chao Liu's avatar
Chao Liu committed
193
{
Chao Liu's avatar
Chao Liu committed
194
    return Tuple<remove_cvref_t<Xs>...>(std::forward<Xs>(xs)...);
Chao Liu's avatar
Chao Liu committed
195
196
}

197
198
199
200
201
202
203
// https://en.cppreference.com/w/cpp/utility/tuple/tie
template <typename... Args>
constexpr Tuple<Args&...> tie(Args&... args) noexcept
{
    return {args...};
}

Chao Liu's avatar
Chao Liu committed
204
} // namespace ck