Sequence.hip.hpp 13.4 KB
Newer Older
1
2
3
4
#pragma once
#include "constant_integral.hip.hpp"
#include "functional.hip.hpp"

Chao Liu's avatar
Chao Liu committed
5
struct EmptySequence;
Chao Liu's avatar
Chao Liu committed
6

Chao Liu's avatar
Chao Liu committed
7
template <index_t... Is>
8
9
10
11
struct Sequence
{
    using Type = Sequence<Is...>;

12
    static constexpr index_t mSize = sizeof...(Is);
13

14
15
16
    const index_t mData[mSize] = {Is...};

    __host__ __device__ static constexpr index_t GetSize() { return mSize; }
17

Chao Liu's avatar
Chao Liu committed
18
19
    template <index_t I>
    __host__ __device__ constexpr index_t Get(Number<I>) const
20
21
22
23
    {
        return mData[I];
    }

24
25
    __host__ __device__ index_t operator[](index_t i) const { return mData[i]; }

26
27
    template <index_t... IRs>
    __host__ __device__ constexpr auto ReorderGivenNew2Old(Sequence<IRs...> /*new2old*/) const
28
    {
29
        static_assert(mSize == sizeof...(IRs), "mSize not consistent");
30

31
        constexpr auto old = Type{};
32

33
        return Sequence<old.Get(Number<IRs>{})...>{};
34
35
    }

36
37
    template <index_t... IRs>
    __host__ __device__ constexpr auto ReorderGivenOld2New(Sequence<IRs...> /*old2new*/) const
38
    {
39
        // TODO: don't know how to implement this
40
        printf("Sequence::ReorderGivenOld2New not implemented");
41
42
43
        assert(false);
    }

Chao Liu's avatar
Chao Liu committed
44
45
46
47
48
    __host__ __device__ constexpr auto Reverse() const
    {
        // not implemented
    }

49
50
51
52
    __host__ __device__ constexpr index_t Front() const { return mData[0]; }

    __host__ __device__ constexpr index_t Back() const { return mData[mSize - 1]; }

53
54
55
56
57
58
    template <index_t I>
    __host__ __device__ constexpr auto PushFront(Number<I>) const
    {
        return Sequence<I, Is...>{};
    }

Chao Liu's avatar
Chao Liu committed
59
    template <index_t I>
60
61
62
63
64
    __host__ __device__ constexpr auto PushBack(Number<I>) const
    {
        return Sequence<Is..., I>{};
    }

65
66
    __host__ __device__ constexpr auto PopFront() const;

67
68
    __host__ __device__ constexpr auto PopBack() const;

Chao Liu's avatar
Chao Liu committed
69
    template <index_t... Xs>
Chao Liu's avatar
Chao Liu committed
70
    __host__ __device__ constexpr auto Append(Sequence<Xs...>) const
71
    {
Chao Liu's avatar
Chao Liu committed
72
73
        return Sequence<Is..., Xs...>{};
    }
Chao Liu's avatar
Chao Liu committed
74

Chao Liu's avatar
Chao Liu committed
75
    __host__ __device__ constexpr auto Append(EmptySequence) const;
Chao Liu's avatar
Chao Liu committed
76

Chao Liu's avatar
Chao Liu committed
77
78
79
    template <index_t... Ns>
    __host__ __device__ constexpr auto Extract(Number<Ns>...) const
    {
Chao Liu's avatar
Chao Liu committed
80
        return Sequence<Get(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
81
    }
Chao Liu's avatar
Chao Liu committed
82

Chao Liu's avatar
Chao Liu committed
83
84
85
86
87
88
    template <index_t N>
    struct split_impl
    {
        template <class FirstSeq, class SecondSeq>
        __host__ __device__ constexpr auto operator()(FirstSeq, SecondSeq) const
        {
Chao Liu's avatar
Chao Liu committed
89
90
            constexpr index_t new_first  = FirstSeq{}.PushBack(Number<SecondSeq{}.Front()>{});
            constexpr index_t new_second = SecondSeq{}.PopFront();
Chao Liu's avatar
Chao Liu committed
91
92
93
94
95
96
97

            static_if<(N > 0)>{}([&](auto fwd) {
                return split_impl<N - 1>{}(new_first, fwd(new_second));
            }).else_([&](auto fwd) { return std::make_pair(new_first, fwd(new_second)); });
        }
    };

Chao Liu's avatar
Chao Liu committed
98
    // split one sequence to two sequnces: [0, I) and [I, mSize)
Chao Liu's avatar
Chao Liu committed
99
100
    // return type is std::pair
    template <index_t I>
Chao Liu's avatar
Chao Liu committed
101
    __host__ __device__ constexpr auto Split(Number<I>) const;
Chao Liu's avatar
Chao Liu committed
102
103
104
105
106
107
108
109
110
111
112

    template <index_t I, index_t X>
    __host__ __device__ constexpr auto Modify(Number<I>, Number<X>) const
    {
        constexpr auto first_second = Split(Number<I>{});

        constexpr auto left  = first_second.first;
        constexpr auto right = first_second.second.PopFront();

        return left.PushBack(Number<X>{}).Append(right);
    }
113
114
};

Chao Liu's avatar
Chao Liu committed
115
struct EmptySequence
Chao Liu's avatar
Chao Liu committed
116
{
Chao Liu's avatar
Chao Liu committed
117
    __host__ __device__ static constexpr index_t GetSize() { return 0; }
Chao Liu's avatar
Chao Liu committed
118

Chao Liu's avatar
Chao Liu committed
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
    template <index_t I>
    __host__ __device__ constexpr auto PushFront(Number<I>) const
    {
        return Sequence<I>{};
    }

    template <index_t I>
    __host__ __device__ constexpr auto PushBack(Number<I>) const
    {
        return Sequence<I>{};
    }

    template <class Seq>
    __host__ __device__ constexpr Seq Append(Seq) const
    {
        return Seq{};
    }
};

template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Append(EmptySequence) const
{
    return Type{};
}

// split one sequence to two sequnces: [0, I) and [I, mSize)
// return type is std::pair
template <index_t... Is>
template <index_t I>
__host__ __device__ constexpr auto Sequence<Is...>::Split(Number<I>) const
{
    static_assert(I <= GetSize(), "wrong! split position is too high!");

    static_if<(I == 0)>{}([&](auto fwd) { return std::make_pair(EmptySequence{}, fwd(Type{})); });

    static_if<(I == GetSize())>{}(
        [&](auto fwd) { return std::make_pair(Type{}, fwd(EmptySequence{})); });

    static_if<(I > 0 && I < GetSize())>{}(
        [&](auto fwd) { return split_impl<I>{}(EmptySequence{}, fwd(Type{})); });
Chao Liu's avatar
Chao Liu committed
159
160
}

Chao Liu's avatar
Chao Liu committed
161
162
163
#if 0
template <index_t IBegin, index_t IEnd, index_t Increment>
__host__ __device__ auto make_increasing_sequence(Number<IBegin>, Number<IEnd>, Number<Increment>)
Chao Liu's avatar
Chao Liu committed
164
{
Chao Liu's avatar
Chao Liu committed
165
166
    static_assert(IBegin < IEnd, (IEnd - IBegin) % Increment == 0, "wrong!");

Chao Liu's avatar
Chao Liu committed
167
168
    // not implemented
}
Chao Liu's avatar
Chao Liu committed
169
#endif
Chao Liu's avatar
Chao Liu committed
170
171

template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
172
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
173
174
175
176
177
178
179
{
    static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");

    return Sequence<(Xs + Ys)...>{};
}

template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
180
__host__ __device__ constexpr auto operator-(Sequence<Xs...> seq_x, Sequence<Ys...> seq_y)
Chao Liu's avatar
Chao Liu committed
181
182
183
{
    static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");

Chao Liu's avatar
Chao Liu committed
184
185
    static_for<0, seq_x.GetSize(), 1>{}(
        [&](auto I) { static_assert(seq_x.Get(I) >= seq_y.Get(I), "wrong! going to undeflow"); });
Chao Liu's avatar
Chao Liu committed
186
187
188
189
190

    return Sequence<(Xs - Ys)...>{};
}

template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
191
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
192
193
194
195
196
197
198
{
    static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");

    return Sequence<(Xs * Ys)...>{};
}

template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
199
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
200
201
202
203
204
205
206
{
    static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");

    return Sequence<(Xs / Ys)...>{};
}

template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
207
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
208
209
210
211
212
213
214
{
    static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");

    return Sequence<(Xs % Ys)...>{};
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
215
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
216
{
Chao Liu's avatar
Chao Liu committed
217
    return Sequence<(Xs + Y)...>{};
Chao Liu's avatar
Chao Liu committed
218
219
220
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
221
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
222
{
Chao Liu's avatar
Chao Liu committed
223
224
225
226
227
228
229
230
231
232
    constexpr auto seq_x = Sequence<Xs...>{};

#if 0
    static_for<0, sizeof...(Xs), 1>{}([&](auto Iter) {
        constexpr auto I = decltype(Iter){};
        static_assert(seq_x.Get(I) >= Y, "wrong! going to underflow");
    });
#endif

    return Sequence<(Xs - Y)...>{};
Chao Liu's avatar
Chao Liu committed
233
234
235
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
236
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
237
{
Chao Liu's avatar
Chao Liu committed
238
    return Sequence<(Xs * Y)...>{};
Chao Liu's avatar
Chao Liu committed
239
240
241
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
242
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
243
{
Chao Liu's avatar
Chao Liu committed
244
    return Sequence<(Xs / Y)...>{};
Chao Liu's avatar
Chao Liu committed
245
246
247
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
248
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
249
{
Chao Liu's avatar
Chao Liu committed
250
    return Sequence<(Xs % Y)...>{};
Chao Liu's avatar
Chao Liu committed
251
252
}

Chao Liu's avatar
Chao Liu committed
253
254
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator+(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
255
{
Chao Liu's avatar
Chao Liu committed
256
    return Sequence<(Y + Xs)...>{};
Chao Liu's avatar
Chao Liu committed
257
258
}

Chao Liu's avatar
Chao Liu committed
259
260
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
261
{
Chao Liu's avatar
Chao Liu committed
262
263
264
265
266
267
268
269
    constexpr auto seq_x = Sequence<Xs...>{};

    static_for<0, sizeof...(Xs), 1>{}([&](auto Iter) {
        constexpr auto I = decltype(Iter){};
        static_assert(seq_x.Get(I) <= Y, "wrong! going to underflow");
    });

    return Sequence<(Y - Xs)...>{};
Chao Liu's avatar
Chao Liu committed
270
271
}

Chao Liu's avatar
Chao Liu committed
272
273
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator*(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
274
{
Chao Liu's avatar
Chao Liu committed
275
    return Sequence<(Y * Xs)...>{};
Chao Liu's avatar
Chao Liu committed
276
277
}

Chao Liu's avatar
Chao Liu committed
278
279
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator/(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
280
{
Chao Liu's avatar
Chao Liu committed
281
    return Sequence<(Y / Xs)...>{};
Chao Liu's avatar
Chao Liu committed
282
283
}

Chao Liu's avatar
Chao Liu committed
284
285
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator%(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
286
{
Chao Liu's avatar
Chao Liu committed
287
    return Sequence<(Y % Xs)...>{};
Chao Liu's avatar
Chao Liu committed
288
289
}

290
291
292
293
294
295
296
template <index_t I, index_t... Is>
__host__ __device__ constexpr auto sequence_pop_front(Sequence<I, Is...>)
{
    static_assert(sizeof...(Is) > 0, "empty Sequence!");
    return Sequence<Is...>{};
}

297
298
#if 0
// TODO: for some reason, compiler cannot instantiate this template
Chao Liu's avatar
Chao Liu committed
299
template <index_t... Is, index_t I>
300
301
__host__ __device__ constexpr auto sequence_pop_back(Sequence<Is..., I>)
{
302
    static_assert(sizeof...(Is) > 0, "empty Sequence!");
303
304
    return Sequence<Is...>{};
}
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
#else
// TODO: delete these very ugly mess
template <index_t I0, index_t I1>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<I0, I1>)
{
    return Sequence<I0>{};
}

template <index_t I0, index_t I1, index_t I2>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<I0, I1, I2>)
{
    return Sequence<I0, I1>{};
}

template <index_t I0, index_t I1, index_t I2, index_t I3>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<I0, I1, I2, I3>)
{
    return Sequence<I0, I1, I2>{};
}

template <index_t I0, index_t I1, index_t I2, index_t I3, index_t I4>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<I0, I1, I2, I3, I4>)
{
    return Sequence<I0, I1, I2, I3>{};
}

template <index_t I0, index_t I1, index_t I2, index_t I3, index_t I4, index_t I5>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<I0, I1, I2, I3, I4, I5>)
{
    return Sequence<I0, I1, I2, I3, I4>{};
}

template <index_t I0, index_t I1, index_t I2, index_t I3, index_t I4, index_t I5, index_t I6>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<I0, I1, I2, I3, I4, I5, I6>)
{
    return Sequence<I0, I1, I2, I3, I4, I5>{};
}

template <index_t I0,
          index_t I1,
          index_t I2,
          index_t I3,
          index_t I4,
          index_t I5,
          index_t I6,
          index_t I7>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<I0, I1, I2, I3, I4, I5, I6, I7>)
{
    return Sequence<I0, I1, I2, I3, I4, I5, I6>{};
}

template <index_t I0,
          index_t I1,
          index_t I2,
          index_t I3,
          index_t I4,
          index_t I5,
          index_t I6,
          index_t I7,
          index_t I8>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<I0, I1, I2, I3, I4, I5, I6, I7, I8>)
{
    return Sequence<I0, I1, I2, I3, I4, I5, I6, I7>{};
}

template <index_t I0,
          index_t I1,
          index_t I2,
          index_t I3,
          index_t I4,
          index_t I5,
          index_t I6,
          index_t I7,
          index_t I8,
          index_t I9>
__host__ __device__ constexpr auto
    sequence_pop_back(Sequence<I0, I1, I2, I3, I4, I5, I6, I7, I8, I9>)
{
    return Sequence<I0, I1, I2, I3, I4, I5, I6, I7, I8>{};
}
#endif
386

Chao Liu's avatar
Chao Liu committed
387
388
389
390
391
392
template <class F, index_t... Xs>
__host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>)
{
    return Sequence<f(Xs)...>{};
}

Chao Liu's avatar
Chao Liu committed
393
template <class F, index_t... Xs, index_t... Ys>
394
__host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>)
395
{
396
    static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize, "Dim not the same");
397
398
399
400

    return Sequence<f(Xs, Ys)...>{};
}

401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
template <class F, index_t... Xs, index_t... Ys, index_t... Zs>
__host__ __device__ constexpr auto
transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>, Sequence<Zs...>)
{
    static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize &&
                      Sequence<Xs...>::mSize == Sequence<Zs...>::mSize,
                  "Dim not the same");

    return Sequence<f(Xs, Ys, Zs)...>{};
}

template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::PopFront() const
{
    return sequence_pop_front(Type{});
416
417
}

Chao Liu's avatar
Chao Liu committed
418
template <index_t... Is>
419
420
421
422
__host__ __device__ constexpr auto Sequence<Is...>::PopBack() const
{
    return sequence_pop_back(Type{});
}
423
424

template <class Seq>
Chao Liu's avatar
Chao Liu committed
425
struct accumulate_on_sequence_impl
426
427
428
429
430
431
432
433
434
{
    template <class IDim>
    __host__ __device__ constexpr index_t operator()(IDim) const
    {
        return Seq{}.Get(IDim{});
    }
};

template <class Seq, class Reduce, index_t I>
Chao Liu's avatar
Chao Liu committed
435
436
__host__ __device__ constexpr index_t
    accumulate_on_sequence(Seq, Reduce, Number<I> /*initial_value*/)
437
438
{
    constexpr index_t a =
Chao Liu's avatar
Chao Liu committed
439
        static_const_reduce_n<Seq::mSize>{}(accumulate_on_sequence_impl<Seq>{}, Reduce{});
440
441
    return Reduce{}(a, I);
}
Chao Liu's avatar
Chao Liu committed
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477

template <index_t NRemain>
struct scan_sequence_impl
{
    template <class ScanedSeq, class RemainSeq, class Reduce>
    __host__ __device__ constexpr auto operator()(ScanedSeq, RemainSeq, Reduce) const
    {
        static_assert(RemainSeq{}.GetSize() == NRemain,
                      "wrong! RemainSeq and NRemain not consistent!");

        constexpr index_t a       = Reduce{}(ScanedSeq{}.Back(), RemainSeq{}.Front());
        constexpr auto scaned_seq = ScanedSeq{}.PushBack(Number<a>{});

        static_if<(NRemain > 1)>{}([&](auto fwd) {
            return scan_sequence_impl<NRemain - 1>{}(
                scaned_seq, RemainSeq{}.PopFront(), fwd(Reduce{}));
        }).else_([&](auto fwd) { return fwd(scaned_seq); });
    }
};

template <class Seq, class Reduce>
__host__ __device__ constexpr auto scan_sequence(Seq, Reduce)
{
    constexpr auto scaned_seq = Sequence<Seq{}.front()>{};
    constexpr auto remain_seq = Seq{}.PopFront();

    constexpr index_t remain_size = Seq::GetSize() - 1;

    return scan_sequence_impl<remain_size>{}(scaned_seq, remain_seq, Reduce{});
}

template <class Seq, class Reduce>
__host__ __device__ constexpr auto reverse_scan_sequence(Seq, Reduce)
{
    return scan_seqeunce(Seq{}.Reverse(), Reduce{}).Reverse();
}