Sequence.hip.hpp 16.3 KB
Newer Older
1
#pragma once
Chao Liu's avatar
Chao Liu committed
2
#include "integral_constant.hip.hpp"
3
4
#include "functional.hip.hpp"

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

10
    static constexpr index_t mSize = sizeof...(Is);
11

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

Chao Liu's avatar
Chao Liu committed
14
    template <index_t I>
15
    __host__ __device__ static constexpr index_t Get(Number<I>)
16
    {
17
18
19
20
        static_assert(I < mSize, "wrong! I too large");

        // the last dummy element is to prevent compiler complain about empty Sequence
        const index_t mData[mSize + 1] = {Is..., 0};
21
22
23
        return mData[I];
    }

24
    template <index_t... IRs>
25
    __host__ __device__ static constexpr auto ReorderGivenNew2Old(Sequence<IRs...> /*new2old*/)
26
    {
27
28
29
30
31
#if 0 // require sequence_sort, which is not implemented yet
        static_assert(is_same<sequence_sort<Sequence<IRs...>>::SortedSeqType,
                              arithmetic_sequence_gen<0, mSize, 1>::SeqType>::value,
                      "wrong! invalid new2old map");
#endif
32

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

36
37
38
#if 0 // require sequence_sort, which is not implemented yet
    template <class MapOld2New>
    __host__ __device__ static constexpr auto ReorderGivenOld2New(MapOld2New /*old2new*/)
39
    {
Chao Liu's avatar
Chao Liu committed
40
#if 0
41
42
43
        static_assert(is_same<sequence_sort<MapOld2New>::SortedSeqType,
                              arithmetic_sequence_gen<0, mSize, 1>::SeqType>::value,
                      "wrong! invalid old2new map");
Chao Liu's avatar
Chao Liu committed
44
#endif
45
46
47
        constexpr auto map_new2old = typename sequence_map_inverse<MapOld2New>::SeqMapType{};

        return ReorderGivenNew2Old(map_new2old);
48
    }
49
#endif
50

51
    __host__ __device__ static constexpr auto Reverse();
Chao Liu's avatar
Chao Liu committed
52

53
54
55
56
57
    __host__ __device__ static constexpr index_t Front()
    {
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[0];
    }
58

59
60
61
62
63
    __host__ __device__ static constexpr index_t Back()
    {
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[mSize - 1];
    }
64

65
    template <index_t I>
66
    __host__ __device__ static constexpr auto PushFront(Number<I>)
67
68
69
70
    {
        return Sequence<I, Is...>{};
    }

Chao Liu's avatar
Chao Liu committed
71
    template <index_t I>
72
    __host__ __device__ static constexpr auto PushBack(Number<I>)
73
74
75
76
    {
        return Sequence<Is..., I>{};
    }

77
    __host__ __device__ static constexpr auto PopFront();
78

79
    __host__ __device__ static constexpr auto PopBack();
80

Chao Liu's avatar
Chao Liu committed
81
    template <index_t... Xs>
82
    __host__ __device__ static constexpr auto Append(Sequence<Xs...>)
83
    {
Chao Liu's avatar
Chao Liu committed
84
85
        return Sequence<Is..., Xs...>{};
    }
Chao Liu's avatar
Chao Liu committed
86

Chao Liu's avatar
Chao Liu committed
87
    template <index_t... Ns>
88
    __host__ __device__ static constexpr auto Extract(Number<Ns>...)
Chao Liu's avatar
Chao Liu committed
89
    {
Chao Liu's avatar
Chao Liu committed
90
        return Sequence<Type{}.Get(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
91
    }
Chao Liu's avatar
Chao Liu committed
92

Chao Liu's avatar
Chao Liu committed
93
    template <index_t... Ns>
94
    __host__ __device__ static constexpr auto Extract(Sequence<Ns...>)
Chao Liu's avatar
Chao Liu committed
95
    {
Chao Liu's avatar
Chao Liu committed
96
        return Sequence<Type{}.Get(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
97
    }
98
99
100

    template <index_t I, index_t X>
    __host__ __device__ static constexpr auto Modify(Number<I>, Number<X>);
101
102
};

Chao Liu's avatar
Chao Liu committed
103
// merge sequence
Chao Liu's avatar
Chao Liu committed
104
105
template <class, class>
struct sequence_merge;
Chao Liu's avatar
Chao Liu committed
106

Chao Liu's avatar
Chao Liu committed
107
108
109
template <index_t... Xs, index_t... Ys>
struct sequence_merge<Sequence<Xs...>, Sequence<Ys...>>
{
Chao Liu's avatar
Chao Liu committed
110
    using SeqType = Sequence<Xs..., Ys...>;
Chao Liu's avatar
Chao Liu committed
111
};
Chao Liu's avatar
Chao Liu committed
112

Chao Liu's avatar
Chao Liu committed
113
// arithmetic sqeuence
Chao Liu's avatar
Chao Liu committed
114
template <index_t IBegin, index_t NSize, index_t Increment>
115
struct arithmetic_sequence_gen_impl
Chao Liu's avatar
Chao Liu committed
116
117
{
    static constexpr index_t NSizeLeft = NSize / 2;
Chao Liu's avatar
Chao Liu committed
118

Chao Liu's avatar
Chao Liu committed
119
    using SeqType = typename sequence_merge<
120
121
        typename arithmetic_sequence_gen_impl<IBegin, NSizeLeft, Increment>::SeqType,
        typename arithmetic_sequence_gen_impl<IBegin + NSizeLeft * Increment,
Chao Liu's avatar
Chao Liu committed
122
123
                                              NSize - NSizeLeft,
                                              Increment>::SeqType>::SeqType;
Chao Liu's avatar
Chao Liu committed
124
125
};

Chao Liu's avatar
Chao Liu committed
126
template <index_t IBegin, index_t Increment>
127
struct arithmetic_sequence_gen_impl<IBegin, 1, Increment>
Chao Liu's avatar
Chao Liu committed
128
{
Chao Liu's avatar
Chao Liu committed
129
    using SeqType = Sequence<IBegin>;
Chao Liu's avatar
Chao Liu committed
130
};
Chao Liu's avatar
Chao Liu committed
131

Chao Liu's avatar
Chao Liu committed
132
template <index_t IBegin, index_t Increment>
133
struct arithmetic_sequence_gen_impl<IBegin, 0, Increment>
Chao Liu's avatar
Chao Liu committed
134
{
Chao Liu's avatar
Chao Liu committed
135
136
137
138
    using SeqType = Sequence<>;
};

template <index_t IBegin, index_t IEnd, index_t Increment>
139
struct arithmetic_sequence_gen
Chao Liu's avatar
Chao Liu committed
140
141
{
    using SeqType =
142
        typename arithmetic_sequence_gen_impl<IBegin, IEnd - IBegin, Increment>::SeqType;
Chao Liu's avatar
Chao Liu committed
143
};
Chao Liu's avatar
Chao Liu committed
144

Chao Liu's avatar
Chao Liu committed
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
// transform sequence
template <class, class>
struct sequence_transform;

template <class F, index_t... Is>
struct sequence_transform<F, Sequence<Is...>>
{
    using SeqType = Sequence<F{}(Is)...>;
};

// uniform sequence
template <index_t NSize, index_t I>
struct uniform_sequence_gen
{
    struct return_constant
    {
        __host__ __device__ constexpr index_t operator()(index_t) const { return I; }
    };

    using SeqType = typename sequence_transform<
        return_constant,
        typename arithmetic_sequence_gen<0, NSize, 1>::SeqType>::SeqType;
};

// reverse inclusive scan (with init) sequence
Chao Liu's avatar
Chao Liu committed
170
template <class, class, index_t>
Chao Liu's avatar
Chao Liu committed
171
struct sequence_reverse_inclusive_scan;
Chao Liu's avatar
Chao Liu committed
172

Chao Liu's avatar
Chao Liu committed
173
174
template <index_t I, index_t... Is, class Reduce, index_t Init>
struct sequence_reverse_inclusive_scan<Sequence<I, Is...>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
175
{
Chao Liu's avatar
Chao Liu committed
176
177
    using old_scan =
        typename sequence_reverse_inclusive_scan<Sequence<Is...>, Reduce, Init>::SeqType;
Chao Liu's avatar
Chao Liu committed
178
179
180
181
182
183

    static constexpr index_t new_reduce = Reduce{}(I, old_scan{}.Front());

    using SeqType = typename sequence_merge<Sequence<new_reduce>, old_scan>::SeqType;
};

Chao Liu's avatar
Chao Liu committed
184
185
template <index_t I, class Reduce, index_t Init>
struct sequence_reverse_inclusive_scan<Sequence<I>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
186
{
Chao Liu's avatar
Chao Liu committed
187
    using SeqType = Sequence<Reduce{}(I, Init)>;
Chao Liu's avatar
Chao Liu committed
188
189
};

Chao Liu's avatar
Chao Liu committed
190
191
template <class Reduce, index_t Init>
struct sequence_reverse_inclusive_scan<Sequence<>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
192
193
194
195
{
    using SeqType = Sequence<>;
};

Chao Liu's avatar
Chao Liu committed
196
// extract sequence
Chao Liu's avatar
Chao Liu committed
197
198
199
200
201
202
203
204
205
template <class, class>
struct sequence_extract;

template <class Seq, index_t... Is>
struct sequence_extract<Seq, Sequence<Is...>>
{
    using SeqType = Sequence<Seq{}.Get(Number<Is>{})...>;
};

Chao Liu's avatar
Chao Liu committed
206
// split sequence
Chao Liu's avatar
Chao Liu committed
207
208
209
210
211
template <class Seq, index_t I>
struct sequence_split
{
    static constexpr index_t NSize = Seq{}.GetSize();

212
213
    using range0 = typename arithmetic_sequence_gen<0, I, 1>::SeqType;
    using range1 = typename arithmetic_sequence_gen<I, NSize, 1>::SeqType;
Chao Liu's avatar
Chao Liu committed
214
215
216
217
218

    using SeqType0 = typename sequence_extract<Seq, range0>::SeqType;
    using SeqType1 = typename sequence_extract<Seq, range1>::SeqType;
};

Chao Liu's avatar
Chao Liu committed
219
// reverse sequence
Chao Liu's avatar
Chao Liu committed
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
template <class Seq>
struct sequence_reverse
{
    static constexpr index_t NSize = Seq{}.GetSize();

    using seq_split = sequence_split<Seq, NSize / 2>;
    using SeqType   = typename sequence_merge<
        typename sequence_reverse<typename seq_split::SeqType1>::SeqType,
        typename sequence_reverse<typename seq_split::SeqType0>::SeqType>::SeqType;
};

template <index_t I>
struct sequence_reverse<Sequence<I>>
{
    using SeqType = Sequence<I>;
};

template <index_t I0, index_t I1>
struct sequence_reverse<Sequence<I0, I1>>
{
    using SeqType = Sequence<I1, I0>;
};
Chao Liu's avatar
Chao Liu committed
242

243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
#if 0 // not fully implemented
template <class KeySeq0, class ValSeq0, class KeySeq1, class ValSeq1>
struct sequence_sort_merge_impl;

template <index_t Key0,
          index_t... Keys0,
          index_t Val0,
          index_t... Vals0,
          index_t Key1,
          index_t... Keys1,
          index_t Val0,
          index_t... Vals1>
struct sequence_sort_merge_impl<Sequence<Key0, Keys0...>,
                                Sequence<Val0, Vals0...>,
                                Sequence<Key1, Keys1...>,
                                Sequence<Val1, Vals1...>>
{
};

template <class>
struct sequence_sort;

template <index_t... Is>
struct sequence_sort<Sequence<Is...>>
{
    using OriginalSeqType        = Sequence<Is...>;
    using SortedSeqType          = xxxxx;
    using MapSorted2OriginalType = xxx;
};

template <class Seq, class IsValidSeqMap>
struct sequence_map_inverse_impl;

// impl for valid map, no impl for invalid map
template <index_t... Is>
struct sequence_map_inverse_impl<Sequence<Is...>, true>
{
    using SeqMapType = sequence_sort<Sequence<Is...>>::MapSorted2OriginalType;
};

template <class>
struct sequence_map_inverse;

template <class Is...>
struct sequence_map_inverse<Sequence<Is...>>
{
    // TODO: make sure the map to be inversed is valid: [0, sizeof...(Is))
    static constexpr bool is_valid_sequence_map =
        is_same<typename sequence_sort<Sequence<Is...>>::SortedSeqType,
                typename arithmetic_sequence_gen<0, sizeof...(Is), 1>::SeqType>::value;

    // make compiler fails, if is_valid_map != true
    using SeqMapType =
        typename sequence_map_inverse_impl<Sequence<Is...>, is_valid_map>::SeqMapType;
};
Chao Liu's avatar
Chao Liu committed
298

299
#endif
Chao Liu's avatar
Chao Liu committed
300
301
302
303
304
305
306
307
308
309
310
template <class Seq>
struct is_valid_sequence_map
{
    static constexpr bool value =
#if 0 // sequence_sort is not implemented yet
        is_same<typename arithmetic_sequence_gen<0, Seq::GetSize(), 1>::SeqType,
                typename sequence_sort<Seq>::SortedSeqType>::value;
#else
        true;
#endif
};
311

Chao Liu's avatar
Chao Liu committed
312
template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
313
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
314
315
316
317
318
319
320
{
    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
321
__host__ __device__ constexpr auto operator-(Sequence<Xs...> seq_x, Sequence<Ys...> seq_y)
Chao Liu's avatar
Chao Liu committed
322
323
324
{
    static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");

Chao Liu's avatar
Chao Liu committed
325
#if 0
Chao Liu's avatar
Chao Liu committed
326
327
    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
328
#endif
Chao Liu's avatar
Chao Liu committed
329
330
331
332
333

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

template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
334
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
335
336
337
338
339
340
341
{
    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
342
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
343
344
345
346
347
348
349
{
    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
350
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
351
352
353
354
355
356
357
{
    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
358
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
359
{
Chao Liu's avatar
Chao Liu committed
360
    return Sequence<(Xs + Y)...>{};
Chao Liu's avatar
Chao Liu committed
361
362
363
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
364
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
365
{
366
#if 0 // TODO: turn it on. Doesn't compile
Chao Liu's avatar
Chao Liu committed
367
368
369
370
371
372
373
374
375
    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");
    });
#endif

    return Sequence<(Xs - Y)...>{};
Chao Liu's avatar
Chao Liu committed
376
377
378
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
379
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
380
{
Chao Liu's avatar
Chao Liu committed
381
    return Sequence<(Xs * Y)...>{};
Chao Liu's avatar
Chao Liu committed
382
383
384
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
385
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
386
{
Chao Liu's avatar
Chao Liu committed
387
    return Sequence<(Xs / Y)...>{};
Chao Liu's avatar
Chao Liu committed
388
389
390
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
391
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
392
{
Chao Liu's avatar
Chao Liu committed
393
    return Sequence<(Xs % Y)...>{};
Chao Liu's avatar
Chao Liu committed
394
395
}

Chao Liu's avatar
Chao Liu committed
396
397
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator+(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
398
{
Chao Liu's avatar
Chao Liu committed
399
    return Sequence<(Y + Xs)...>{};
Chao Liu's avatar
Chao Liu committed
400
401
}

Chao Liu's avatar
Chao Liu committed
402
403
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
404
{
Chao Liu's avatar
Chao Liu committed
405
406
    constexpr auto seq_x = Sequence<Xs...>{};

Chao Liu's avatar
Chao Liu committed
407
#if 0
Chao Liu's avatar
Chao Liu committed
408
409
410
411
    static_for<0, sizeof...(Xs), 1>{}([&](auto Iter) {
        constexpr auto I = decltype(Iter){};
        static_assert(seq_x.Get(I) <= Y, "wrong! going to underflow");
    });
Chao Liu's avatar
Chao Liu committed
412
#endif
Chao Liu's avatar
Chao Liu committed
413
414

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

Chao Liu's avatar
Chao Liu committed
417
418
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator*(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
419
{
Chao Liu's avatar
Chao Liu committed
420
    return Sequence<(Y * Xs)...>{};
Chao Liu's avatar
Chao Liu committed
421
422
}

Chao Liu's avatar
Chao Liu committed
423
424
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator/(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
425
{
Chao Liu's avatar
Chao Liu committed
426
    return Sequence<(Y / Xs)...>{};
Chao Liu's avatar
Chao Liu committed
427
428
}

Chao Liu's avatar
Chao Liu committed
429
430
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator%(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
431
{
Chao Liu's avatar
Chao Liu committed
432
    return Sequence<(Y % Xs)...>{};
Chao Liu's avatar
Chao Liu committed
433
434
}

435
436
437
438
439
440
template <index_t I, index_t... Is>
__host__ __device__ constexpr auto sequence_pop_front(Sequence<I, Is...>)
{
    return Sequence<Is...>{};
}

Chao Liu's avatar
Chao Liu committed
441
442
template <class Seq>
__host__ __device__ constexpr auto sequence_pop_back(Seq)
443
{
444
    static_assert(Seq{}.GetSize() > 0, "wrong! cannot pop an empty Sequence!");
Chao Liu's avatar
Chao Liu committed
445
    return sequence_pop_front(Seq{}.Reverse()).Reverse();
446
}
447

Chao Liu's avatar
Chao Liu committed
448
449
450
451
452
453
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
454
template <class F, index_t... Xs, index_t... Ys>
455
__host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>)
456
{
457
    static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize, "Dim not the same");
458
459
460
461

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

462
463
464
465
466
467
468
469
470
471
472
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)...>{};
}

Chao Liu's avatar
Chao Liu committed
473
474
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto reverse_inclusive_scan_sequence(Seq, Reduce, Number<Init>)
475
{
Chao Liu's avatar
Chao Liu committed
476
    return typename sequence_reverse_inclusive_scan<Seq, Reduce, Init>::SeqType{};
477
478
}

Chao Liu's avatar
Chao Liu committed
479
480
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto inclusive_scan_sequence(Seq, Reduce, Number<Init>)
481
{
Chao Liu's avatar
Chao Liu committed
482
    return reverse_inclusive_scan_sequence(Seq{}.Reverse(), Reduce{}, Number<Init>{}).Reverse();
483
}
484
485

template <class Seq>
Chao Liu's avatar
Chao Liu committed
486
struct accumulate_on_sequence_impl
487
488
489
490
491
492
493
494
495
{
    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
496
497
__host__ __device__ constexpr index_t
    accumulate_on_sequence(Seq, Reduce, Number<I> /*initial_value*/)
498
499
{
    constexpr index_t a =
Chao Liu's avatar
Chao Liu committed
500
        static_const_reduce_n<Seq::mSize>{}(accumulate_on_sequence_impl<Seq>{}, Reduce{});
501
502
    return Reduce{}(a, I);
}
Chao Liu's avatar
Chao Liu committed
503

Chao Liu's avatar
Chao Liu committed
504
template <index_t... Is>
505
__host__ __device__ constexpr auto Sequence<Is...>::PopFront()
Chao Liu's avatar
Chao Liu committed
506
{
507
    return sequence_pop_front(Type{});
Chao Liu's avatar
Chao Liu committed
508
}
Chao Liu's avatar
Chao Liu committed
509

510
511
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::PopBack()
Chao Liu's avatar
Chao Liu committed
512
{
513
    return sequence_pop_back(Type{});
Chao Liu's avatar
Chao Liu committed
514
515
}

516
517
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Reverse()
Chao Liu's avatar
Chao Liu committed
518
{
519
520
521
522
523
524
525
526
527
528
529
530
531
532
    return typename sequence_reverse<Sequence<Is...>>::SeqType{};
}

template <index_t... Is>
template <index_t I, index_t X>
__host__ __device__ constexpr auto Sequence<Is...>::Modify(Number<I>, Number<X>)
{
    static_assert(I < GetSize(), "wrong!");

    using seq_split          = sequence_split<Type, I>;
    constexpr auto seq_left  = typename seq_split::SeqType0{};
    constexpr auto seq_right = typename seq_split::SeqType1{}.PopFront();

    return seq_left.PushBack(Number<X>{}).Append(seq_right);
Chao Liu's avatar
Chao Liu committed
533
}
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569

template <index_t... Xs>
__host__ __device__ void print_Sequence(const char* s, Sequence<Xs...>)
{
    constexpr index_t nsize = Sequence<Xs...>::GetSize();

    static_assert(nsize <= 10, "wrong!");

    static_if<nsize == 0>{}([&](auto) { printf("%s size %u, {}\n", s, nsize, Xs...); });

    static_if<nsize == 1>{}([&](auto) { printf("%s size %u, {%u}\n", s, nsize, Xs...); });

    static_if<nsize == 2>{}([&](auto) { printf("%s size %u, {%u %u}\n", s, nsize, Xs...); });

    static_if<nsize == 3>{}([&](auto) { printf("%s size %u, {%u %u %u}\n", s, nsize, Xs...); });

    static_if<nsize == 4>{}([&](auto) { printf("%s size %u, {%u %u %u %u}\n", s, nsize, Xs...); });

    static_if<nsize == 5>{}(
        [&](auto) { printf("%s size %u, {%u %u %u %u %u}\n", s, nsize, Xs...); });

    static_if<nsize == 6>{}(
        [&](auto) { printf("%s size %u, {%u %u %u %u %u %u}\n", s, nsize, Xs...); });

    static_if<nsize == 7>{}(
        [&](auto) { printf("%s size %u, {%u %u %u %u %u %u %u}\n", s, nsize, Xs...); });

    static_if<nsize == 8>{}(
        [&](auto) { printf("%s size %u, {%u %u %u %u %u %u %u %u}\n", s, nsize, Xs...); });

    static_if<nsize == 9>{}(
        [&](auto) { printf("%s size %u, {%u %u %u %u %u %u %u %u %u}\n", s, nsize, Xs...); });

    static_if<nsize == 10>{}(
        [&](auto) { printf("%s size %u, {%u %u %u %u %u %u %u %u %u %u}\n", s, nsize, Xs...); });
}