array.hpp 10.9 KB
Newer Older
1
2
3
#ifndef CK_ARRAY_HPP
#define CK_ARRAY_HPP

Chao Liu's avatar
Chao Liu committed
4
#include "sequence.hpp"
Chao Liu's avatar
Chao Liu committed
5
#include "functional2.hpp"
6

7
8
namespace ck {

Chao Liu's avatar
Chao Liu committed
9
template <typename TData, index_t NSize>
10
11
struct Array
{
Chao Liu's avatar
Chao Liu committed
12
    using type      = Array<TData, NSize>;
Chao Liu's avatar
Chao Liu committed
13
    using data_type = TData;
14

Chao Liu's avatar
Chao Liu committed
15
    index_t mData[NSize];
16

Chao Liu's avatar
Chao Liu committed
17
    __host__ __device__ explicit constexpr Array() {}
18

Chao Liu's avatar
Chao Liu committed
19
    template <typename X, typename... Xs>
Chao Liu's avatar
Chao Liu committed
20
    __host__ __device__ constexpr Array(X x, Xs... xs)
Chao Liu's avatar
Chao Liu committed
21
        : mData{static_cast<TData>(x), static_cast<TData>(xs)...}
22
    {
Chao Liu's avatar
Chao Liu committed
23
        static_assert(sizeof...(Xs) + 1 == NSize, "wrong! size");
24
25
    }

Chao Liu's avatar
Chao Liu committed
26
27
    __host__ __device__ static constexpr index_t Size() { return NSize; }

Chao Liu's avatar
Chao Liu committed
28
    // TODO: remove
Chao Liu's avatar
Chao Liu committed
29
    __host__ __device__ static constexpr index_t GetSize() { return Size(); }
30

Chao Liu's avatar
Chao Liu committed
31
    template <index_t I>
Chao Liu's avatar
Chao Liu committed
32
    __host__ __device__ constexpr const TData& At(Number<I>) const
Chao Liu's avatar
Chao Liu committed
33
    {
Chao Liu's avatar
Chao Liu committed
34
35
        static_assert(I < NSize, "wrong!");

Chao Liu's avatar
Chao Liu committed
36
37
38
39
        return mData[I];
    }

    template <index_t I>
Chao Liu's avatar
Chao Liu committed
40
    __host__ __device__ constexpr TData& At(Number<I>)
Chao Liu's avatar
Chao Liu committed
41
    {
Chao Liu's avatar
Chao Liu committed
42
43
        static_assert(I < NSize, "wrong!");

Chao Liu's avatar
Chao Liu committed
44
45
46
        return mData[I];
    }

Chao Liu's avatar
Chao Liu committed
47
    __host__ __device__ constexpr const TData& At(index_t i) const { return mData[i]; }
48

Chao Liu's avatar
Chao Liu committed
49
50
51
52
    __host__ __device__ constexpr TData& At(index_t i) { return mData[i]; }

    template <typename I>
    __host__ __device__ constexpr const TData& operator[](I i) const
Chao Liu's avatar
Chao Liu committed
53
    {
Chao Liu's avatar
Chao Liu committed
54
55
        return At(i);
    }
Chao Liu's avatar
Chao Liu committed
56

Chao Liu's avatar
Chao Liu committed
57
58
59
60
    template <typename I>
    __host__ __device__ constexpr TData& operator()(I i)
    {
        return At(i);
Chao Liu's avatar
Chao Liu committed
61
62
    }

Chao Liu's avatar
Chao Liu committed
63
64
65
66
67
68
69
    template <typename T>
    __host__ __device__ constexpr type& operator=(const T& x)
    {
        static_for<0, Size(), 1>{}([&](auto i) { operator()(i) = x[i]; });

        return *this;
    }
Chao Liu's avatar
Chao Liu committed
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84

    struct lambda_PushBack // emulate constexpr lambda
    {
        const Array<TData, NSize>& old_array;
        Array<TData, NSize + 1>& new_array;

        __host__ __device__ constexpr lambda_PushBack(const Array<TData, NSize>& old_array_,
                                                      Array<TData, NSize + 1>& new_array_)
            : old_array(old_array_), new_array(new_array_)
        {
        }

        template <index_t I>
        __host__ __device__ constexpr void operator()(Number<I>) const
        {
Chao Liu's avatar
Chao Liu committed
85
            new_array(Number<I>{}) = old_array[I];
Chao Liu's avatar
Chao Liu committed
86
87
88
        }
    };

Chao Liu's avatar
Chao Liu committed
89
    __host__ __device__ constexpr auto PushBack(TData x) const
90
91
92
    {
        Array<TData, NSize + 1> new_array;

Chao Liu's avatar
Chao Liu committed
93
        static_for<0, NSize, 1>{}(lambda_PushBack(*this, new_array));
94

Chao Liu's avatar
Chao Liu committed
95
        new_array(Number<NSize>{}) = x;
96
97
98

        return new_array;
    }
99
};
100

Chao Liu's avatar
Chao Liu committed
101
// Arr: Array
Chao Liu's avatar
Chao Liu committed
102
// Picks: Sequence<...>
Chao Liu's avatar
Chao Liu committed
103
template <typename Arr, typename Picks>
Chao Liu's avatar
Chao Liu committed
104
struct ArrayElementPicker
Chao Liu's avatar
Chao Liu committed
105
{
Chao Liu's avatar
Chao Liu committed
106
    using type      = ArrayElementPicker;
Chao Liu's avatar
Chao Liu committed
107
108
    using data_type = typename Arr::data_type;

Chao Liu's avatar
Chao Liu committed
109
110
111
    __host__ __device__ constexpr ArrayElementPicker() = delete;

    __host__ __device__ explicit constexpr ArrayElementPicker(Arr& array) : mArray{array}
Chao Liu's avatar
Chao Liu committed
112
113
114
115
    {
        constexpr index_t imax =
            accumulate_on_sequence(Picks{}, math::maxer<index_t>{}, Number<0>{});

Chao Liu's avatar
Chao Liu committed
116
        static_assert(imax < Arr::Size(), "wrong! exceeding # array element");
Chao Liu's avatar
Chao Liu committed
117
118
    }

Chao Liu's avatar
Chao Liu committed
119
    __host__ __device__ static constexpr auto Size() { return Picks::Size(); }
Chao Liu's avatar
Chao Liu committed
120
121

    template <index_t I>
Chao Liu's avatar
Chao Liu committed
122
    __host__ __device__ constexpr const data_type& At(Number<I>) const
Chao Liu's avatar
Chao Liu committed
123
    {
Chao Liu's avatar
Chao Liu committed
124
125
126
127
        static_assert(I < Size(), "wrong!");

        constexpr auto IP = Picks{}[I];
        return mArray[IP];
Chao Liu's avatar
Chao Liu committed
128
129
    }

Chao Liu's avatar
Chao Liu committed
130
131
    template <index_t I>
    __host__ __device__ constexpr data_type& At(Number<I>)
Chao Liu's avatar
Chao Liu committed
132
    {
Chao Liu's avatar
Chao Liu committed
133
134
135
136
        static_assert(I < Size(), "wrong!");

        constexpr auto IP = Picks{}[I];
        return mArray(IP);
Chao Liu's avatar
Chao Liu committed
137
138
    }

Chao Liu's avatar
Chao Liu committed
139
140
    template <typename I>
    __host__ __device__ constexpr const data_type& operator[](I i) const
Chao Liu's avatar
Chao Liu committed
141
    {
Chao Liu's avatar
Chao Liu committed
142
        return At(i);
Chao Liu's avatar
Chao Liu committed
143
144
    }

Chao Liu's avatar
Chao Liu committed
145
146
    template <typename I>
    __host__ __device__ constexpr data_type& operator()(I i)
Chao Liu's avatar
Chao Liu committed
147
    {
Chao Liu's avatar
Chao Liu committed
148
        return At(i);
Chao Liu's avatar
Chao Liu committed
149
150
    }

Chao Liu's avatar
Chao Liu committed
151
152
153
154
155
156
157
158
159
    template <typename T>
    __host__ __device__ constexpr type& operator=(const T& a)
    {
        static_for<0, Size(), 1>{}([&](auto i) { operator()(i) = a[i]; });

        return *this;
    }

    Arr& mArray;
Chao Liu's avatar
Chao Liu committed
160
161
};

Chao Liu's avatar
Chao Liu committed
162
template <typename Arr, typename Picks>
Chao Liu's avatar
Chao Liu committed
163
164
165
166
167
__host__ __device__ constexpr auto pick_array_element(Arr& a, Picks)
{
    return ArrayElementPicker<Arr, Picks>(a);
}

Chao Liu's avatar
Chao Liu committed
168
169
170
171
172
173
174
175
176
177
template <typename T>
__host__ __device__ constexpr auto to_array(const T& x)
{
    Array<typename T::data_type, T::Size()> y;

    static_for<0, T::Size(), 1>{}([&](auto i) { y.At(i) = x.At(i); });

    return y;
}

Chao Liu's avatar
Chao Liu committed
178
// TODO: remove this
Chao Liu's avatar
Chao Liu committed
179
180
181
182
183
184
template <index_t... Is>
__host__ __device__ constexpr auto sequence2array(Sequence<Is...>)
{
    return Array<index_t, sizeof...(Is)>{Is...};
}

Chao Liu's avatar
Chao Liu committed
185
template <typename TData, index_t NSize>
Chao Liu's avatar
Chao Liu committed
186
187
__host__ __device__ constexpr auto make_zero_array()
{
Chao Liu's avatar
Chao Liu committed
188
    constexpr auto zero_sequence = typename uniform_sequence_gen<NSize, 0>::type{};
Chao Liu's avatar
Chao Liu committed
189
190
    constexpr auto zero_array    = sequence2array(zero_sequence);
    return zero_array;
Chao Liu's avatar
Chao Liu committed
191
192
}

Chao Liu's avatar
Chao Liu committed
193
template <typename TData, index_t NSize, index_t... IRs>
Chao Liu's avatar
Chao Liu committed
194
__host__ __device__ constexpr auto reorder_array_given_new2old(const Array<TData, NSize>& old_array,
Chao Liu's avatar
Chao Liu committed
195
                                                               Sequence<IRs...> /*new2old*/)
196
197
198
{
    static_assert(NSize == sizeof...(IRs), "NSize not consistent");

Chao Liu's avatar
Chao Liu committed
199
    static_assert(is_valid_sequence_map<Sequence<IRs...>>{}, "wrong! invalid reorder map");
200

201
    return Array<TData, NSize>{old_array[IRs]...};
202
203
}

Chao Liu's avatar
Chao Liu committed
204
template <typename TData, index_t NSize, typename MapOld2New>
Chao Liu's avatar
Chao Liu committed
205
struct lambda_reorder_array_given_old2new
Chao Liu's avatar
Chao Liu committed
206
{
Chao Liu's avatar
Chao Liu committed
207
208
    const Array<TData, NSize>& old_array;
    Array<TData, NSize>& new_array;
Chao Liu's avatar
Chao Liu committed
209

Chao Liu's avatar
Chao Liu committed
210
211
212
    __host__ __device__ constexpr lambda_reorder_array_given_old2new(
        const Array<TData, NSize>& old_array_, Array<TData, NSize>& new_array_)
        : old_array(old_array_), new_array(new_array_)
Chao Liu's avatar
Chao Liu committed
213
214
215
216
217
218
    {
    }

    template <index_t IOldDim>
    __host__ __device__ constexpr void operator()(Number<IOldDim>) const
    {
Chao Liu's avatar
Chao Liu committed
219
        TData old_data = old_array[IOldDim];
Chao Liu's avatar
Chao Liu committed
220

Chao Liu's avatar
Chao Liu committed
221
        constexpr index_t INewDim = MapOld2New::At(Number<IOldDim>{});
Chao Liu's avatar
Chao Liu committed
222

Chao Liu's avatar
Chao Liu committed
223
        new_array(Number<INewDim>{}) = old_data;
Chao Liu's avatar
Chao Liu committed
224
225
226
    }
};

Chao Liu's avatar
Chao Liu committed
227
template <typename TData, index_t NSize, index_t... IRs>
Chao Liu's avatar
Chao Liu committed
228
__host__ __device__ constexpr auto reorder_array_given_old2new(const Array<TData, NSize>& old_array,
Chao Liu's avatar
Chao Liu committed
229
                                                               Sequence<IRs...> /*old2new*/)
Chao Liu's avatar
Chao Liu committed
230
231
232
233
234
{
    Array<TData, NSize> new_array;

    static_assert(NSize == sizeof...(IRs), "NSize not consistent");

Chao Liu's avatar
Chao Liu committed
235
236
    static_assert(is_valid_sequence_map<Sequence<IRs...>>::value, "wrong! invalid reorder map");

Chao Liu's avatar
Chao Liu committed
237
    static_for<0, NSize, 1>{}(
Chao Liu's avatar
Chao Liu committed
238
        lambda_reorder_array_given_old2new<TData, NSize, Sequence<IRs...>>(old_array, new_array));
Chao Liu's avatar
Chao Liu committed
239
240
241

    return new_array;
}
Chao Liu's avatar
Chao Liu committed
242

Chao Liu's avatar
Chao Liu committed
243
template <typename TData, index_t NSize, typename ExtractSeq>
Chao Liu's avatar
Chao Liu committed
244
__host__ __device__ constexpr auto extract_array(const Array<TData, NSize>& old_array, ExtractSeq)
245
246
247
248
249
250
251
{
    Array<TData, ExtractSeq::GetSize()> new_array;

    constexpr index_t new_size = ExtractSeq::GetSize();

    static_assert(new_size <= NSize, "wrong! too many extract");

Chao Liu's avatar
Chao Liu committed
252
    static_for<0, new_size, 1>{}([&](auto I) { new_array(I) = old_array[ExtractSeq::At(I)]; });
253
254
255
256

    return new_array;
}

Chao Liu's avatar
Chao Liu committed
257
258
// emulate constepxr lambda for array
template <typename F, typename X, typename Y, typename Z>
Chao Liu's avatar
Chao Liu committed
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
struct lambda_array_math
{
    const F& f;
    const X& x;
    const Y& y;
    Z& z;

    __host__ __device__ constexpr lambda_array_math(const F& f_, const X& x_, const Y& y_, Z& z_)
        : f(f_), x(x_), y(y_), z(z_)
    {
    }

    template <index_t IDim_>
    __host__ __device__ constexpr void operator()(Number<IDim_>) const
    {
        constexpr auto IDim = Number<IDim_>{};
Chao Liu's avatar
Chao Liu committed
275
        z(IDim)             = f(x[IDim], y[IDim]);
Chao Liu's avatar
Chao Liu committed
276
277
278
    }
};

279
// Array = Array + Array
Chao Liu's avatar
Chao Liu committed
280
template <typename TData, index_t NSize>
Chao Liu's avatar
Chao Liu committed
281
__host__ __device__ constexpr auto operator+(Array<TData, NSize> a, Array<TData, NSize> b)
Chao Liu's avatar
Chao Liu committed
282
283
284
{
    Array<TData, NSize> result;

285
    auto f = math::plus<index_t>{};
Chao Liu's avatar
Chao Liu committed
286

Chao Liu's avatar
Chao Liu committed
287
288
289
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
Chao Liu's avatar
Chao Liu committed
290
291
292

    return result;
}
Chao Liu's avatar
Chao Liu committed
293

294
// Array = Array - Array
Chao Liu's avatar
Chao Liu committed
295
template <typename TData, index_t NSize>
296
297
298
299
__host__ __device__ constexpr auto operator-(Array<TData, NSize> a, Array<TData, NSize> b)
{
    Array<TData, NSize> result;

300
    auto f = math::minus<index_t>{};
301

Chao Liu's avatar
Chao Liu committed
302
303
304
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
305
306
307
308

    return result;
}

Chao Liu's avatar
Chao Liu committed
309
// Array += Array
Chao Liu's avatar
Chao Liu committed
310
template <typename TData, index_t NSize>
Chao Liu's avatar
Chao Liu committed
311
312
313
314
315
316
317
__host__ __device__ constexpr auto operator+=(Array<TData, NSize>& a, Array<TData, NSize> b)
{
    a = a + b;
    return a;
}

// Array -= Array
Chao Liu's avatar
Chao Liu committed
318
template <typename TData, index_t NSize>
Chao Liu's avatar
Chao Liu committed
319
320
321
322
323
__host__ __device__ constexpr auto operator-=(Array<TData, NSize>& a, Array<TData, NSize> b)
{
    a = a - b;
    return a;
}
324
// Array = Array + Sequence
Chao Liu's avatar
Chao Liu committed
325
template <typename TData, index_t NSize, index_t... Is>
326
327
328
329
330
331
__host__ __device__ constexpr auto operator+(Array<TData, NSize> a, Sequence<Is...> b)
{
    static_assert(sizeof...(Is) == NSize, "wrong! size not the same");

    Array<TData, NSize> result;

332
    auto f = math::plus<index_t>{};
333

Chao Liu's avatar
Chao Liu committed
334
335
336
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
337
338
339
340
341

    return result;
}

// Array = Array - Sequence
Chao Liu's avatar
Chao Liu committed
342
template <typename TData, index_t NSize, index_t... Is>
343
344
345
346
347
348
__host__ __device__ constexpr auto operator-(Array<TData, NSize> a, Sequence<Is...> b)
{
    static_assert(sizeof...(Is) == NSize, "wrong! size not the same");

    Array<TData, NSize> result;

349
    auto f = math::minus<index_t>{};
350

Chao Liu's avatar
Chao Liu committed
351
352
353
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
354
355
356
357

    return result;
}

Chao Liu's avatar
Chao Liu committed
358
// Array = Array * Sequence
Chao Liu's avatar
Chao Liu committed
359
template <typename TData, index_t NSize, index_t... Is>
Chao Liu's avatar
Chao Liu committed
360
361
362
363
364
365
__host__ __device__ constexpr auto operator*(Array<TData, NSize> a, Sequence<Is...> b)
{
    static_assert(sizeof...(Is) == NSize, "wrong! size not the same");

    Array<TData, NSize> result;

366
    auto f = math::multiplies<index_t>{};
Chao Liu's avatar
Chao Liu committed
367

Chao Liu's avatar
Chao Liu committed
368
369
370
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
Chao Liu's avatar
Chao Liu committed
371
372
373

    return result;
}
374

375
// Array = Sequence - Array
Chao Liu's avatar
Chao Liu committed
376
template <typename TData, index_t NSize, index_t... Is>
377
__host__ __device__ constexpr auto operator-(Sequence<Is...> a, Array<TData, NSize> b)
378
{
379
380
381
    static_assert(sizeof...(Is) == NSize, "wrong! size not the same");

    Array<TData, NSize> result;
382

383
    auto f = math::minus<index_t>{};
384

Chao Liu's avatar
Chao Liu committed
385
386
387
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
388
389
390
391

    return result;
}

Chao Liu's avatar
Chao Liu committed
392
template <typename TData, index_t NSize, typename Reduce>
393
394
395
396
397
398
399
__host__ __device__ constexpr TData
accumulate_on_array(const Array<TData, NSize>& a, Reduce f, TData init)
{
    TData result = init;

    static_assert(NSize > 0, "wrong");

Chao Liu's avatar
Chao Liu committed
400
    static_for<0, NSize, 1>{}([&](auto I) { result = f(result, a[I]); });
401
402
403

    return result;
}
404

405
406
} // namespace ck
#endif