"git@developer.sourcefind.cn:Fzc7075/nunchaku.git" did not exist on "349cf1427c080c951f8576eed153267ec1cdf609"
tuple.hpp 4.84 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
2
3
4
#ifndef CK_TUPLE_HPP
#define CK_TUPLE_HPP

#include "integral_constant.hpp"
Chao Liu's avatar
Chao Liu committed
5
#include "sequence.hpp"
Chao Liu's avatar
Chao Liu committed
6
#include "type.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
19
20
template <typename Key, typename Data>
struct TupleElement
{
Chao Liu's avatar
Chao Liu committed
21
    __host__ __device__ constexpr TupleElement() = default;
Chao Liu's avatar
Chao Liu committed
22

Chao Liu's avatar
Chao Liu committed
23
24
25
26
27
    template <
        typename T,
        typename std::enable_if<!is_same<remove_reference_t<remove_cv_t<T>>, TupleElement>::value,
                                bool>::type = false>
    __host__ __device__ constexpr TupleElement(T&& v) : mData(std::forward<T>(v))
Chao Liu's avatar
Chao Liu committed
28
29
30
    {
    }

Chao Liu's avatar
Chao Liu committed
31
    Data mData;
Chao Liu's avatar
Chao Liu committed
32
33
};

Chao Liu's avatar
Chao Liu committed
34
35
template <typename Key, typename Data>
__host__ __device__ constexpr const Data& get_tuple_element(const TupleElement<Key, Data>& x)
Chao Liu's avatar
Chao Liu committed
36
{
Chao Liu's avatar
Chao Liu committed
37
    return static_cast<const Data&>(x.mData);
Chao Liu's avatar
Chao Liu committed
38
}
Chao Liu's avatar
Chao Liu committed
39

Chao Liu's avatar
Chao Liu committed
40
41
template <typename Key, typename Data>
__host__ __device__ constexpr Data& get_tuple_element(TupleElement<Key, Data>& x)
Chao Liu's avatar
Chao Liu committed
42
{
Chao Liu's avatar
Chao Liu committed
43
44
    return x.mData;
}
Chao Liu's avatar
Chao Liu committed
45

Chao Liu's avatar
Chao Liu committed
46
// TODO: not sure the use of reference is correct
Chao Liu's avatar
Chao Liu committed
47
48
template <typename Key, typename Data>
__host__ __device__ constexpr Data&& get_tuple_element(TupleElement<Key, Data>&& x)
Chao Liu's avatar
Chao Liu committed
49
{
Chao Liu's avatar
Chao Liu committed
50
51
    return static_cast<Data&&>(x.mData);
}
Chao Liu's avatar
Chao Liu committed
52

Chao Liu's avatar
Chao Liu committed
53
54
55
56
57
template <typename Indices, typename... Xs>
struct TupleImpl;

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

    template <
        typename Y,
        typename std::enable_if<sizeof...(Is) == 1 && sizeof...(Xs) == 1 &&
                                    !is_same<remove_reference_t<remove_cv_t<Y>>, TupleImpl>::value,
                                bool>::type = false>
    __host__ __device__ constexpr TupleImpl(Y&& y)
        : TupleElement<TupleElementKey<Is>, Xs>(std::forward<Y>(y))...
Chao Liu's avatar
Chao Liu committed
68
69
70
    {
    }

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

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

    template <index_t I>
    __host__ __device__ constexpr const auto& GetElementByKey(TupleElementKey<I>) const
    {
        return get_tuple_element<TupleElementKey<I>>(*this);
    }

    template <index_t I>
    __host__ __device__ constexpr auto& GetElementByKey(TupleElementKey<I>)
    {
        return get_tuple_element<TupleElementKey<I>>(*this);
    }
Chao Liu's avatar
Chao Liu committed
92
93
};

Chao Liu's avatar
Chao Liu committed
94
95
96
97
} // 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
98
{
Chao Liu's avatar
Chao Liu committed
99
100
101
    using base =
        detail::TupleImpl<typename arithmetic_sequence_gen<0, sizeof...(Xs), 1>::type, Xs...>;

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

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

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

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

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

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

Chao Liu's avatar
Chao Liu committed
135
136
137
138
139
    template <index_t I>
    __host__ __device__ constexpr const auto& operator[](Number<I> i) const
    {
        return At(i);
    }
140

Chao Liu's avatar
Chao Liu committed
141
142
143
144
145
    template <index_t I>
    __host__ __device__ constexpr auto& operator()(Number<I> i)
    {
        return At(i);
    }
Chao Liu's avatar
Chao Liu committed
146

Chao Liu's avatar
Chao Liu committed
147
148
149
150
    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
151

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

Chao Liu's avatar
Chao Liu committed
154
155
        return *this;
    }
zjing14's avatar
zjing14 committed
156
157

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

Chao Liu's avatar
Chao Liu committed
160
161
template <typename... Xs>
__host__ __device__ constexpr auto make_tuple(Xs&&... xs)
Chao Liu's avatar
Chao Liu committed
162
{
Chao Liu's avatar
Chao Liu committed
163
    return Tuple<remove_cv_t<remove_reference_t<Xs>>...>(std::forward<Xs>(xs)...);
Chao Liu's avatar
Chao Liu committed
164
165
}

Chao Liu's avatar
Chao Liu committed
166
167
} // namespace ck
#endif