Sequence.hip.hpp 15.8 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
6
7
template <class Seq>
struct is_valid_sequence_map;

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

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

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

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

Chao Liu's avatar
Chao Liu committed
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
        // the last dummy element is to prevent compiler complain about empty array, when mSize = 0
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[I];
    }

    template <index_t I>
    __host__ __device__ constexpr index_t operator[](Number<I>) const
    {
        static_assert(I < mSize, "wrong! I too large");

        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[I];
    }

    // make sure I is constepxr
    __host__ __device__ constexpr index_t operator[](index_t I) const
    {
39
        const index_t mData[mSize + 1] = {Is..., 0};
40
41
42
        return mData[I];
    }

43
    template <index_t... IRs>
44
    __host__ __device__ static constexpr auto ReorderGivenNew2Old(Sequence<IRs...> /*new2old*/)
45
    {
Chao Liu's avatar
Chao Liu committed
46
        static_assert(sizeof...(Is) == sizeof...(IRs),
Chao Liu's avatar
Chao Liu committed
47
                      "wrong! reorder map should have the same size as Sequence to be rerodered");
Chao Liu's avatar
Chao Liu committed
48

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

        return Sequence<Type::Get(Number<IRs>{})...>{};
52
53
    }

54
55
56
#if 0 // require sequence_sort, which is not implemented yet
    template <class MapOld2New>
    __host__ __device__ static constexpr auto ReorderGivenOld2New(MapOld2New /*old2new*/)
57
    {
Chao Liu's avatar
Chao Liu committed
58
59
60
61
62
63
        static_assert(sizeof...(Is) == MapOld2New::GetSize(),
                      "wrong! reorder map should have the same size as Sequence to be rerodered");

        static_assert(is_valid_sequence_map<MapOld2New>::value, 
                      "wrong! invalid reorder map");

64
65
66
        constexpr auto map_new2old = typename sequence_map_inverse<MapOld2New>::SeqMapType{};

        return ReorderGivenNew2Old(map_new2old);
67
    }
68
#endif
69

70
    __host__ __device__ static constexpr auto Reverse();
Chao Liu's avatar
Chao Liu committed
71

72
73
74
75
76
    __host__ __device__ static constexpr index_t Front()
    {
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[0];
    }
77

78
79
80
81
82
    __host__ __device__ static constexpr index_t Back()
    {
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[mSize - 1];
    }
83

84
    template <index_t I>
85
    __host__ __device__ static constexpr auto PushFront(Number<I>)
86
87
88
89
    {
        return Sequence<I, Is...>{};
    }

Chao Liu's avatar
Chao Liu committed
90
    template <index_t I>
91
    __host__ __device__ static constexpr auto PushBack(Number<I>)
92
93
94
95
    {
        return Sequence<Is..., I>{};
    }

96
    __host__ __device__ static constexpr auto PopFront();
97

98
    __host__ __device__ static constexpr auto PopBack();
99

Chao Liu's avatar
Chao Liu committed
100
    template <index_t... Xs>
101
    __host__ __device__ static constexpr auto Append(Sequence<Xs...>)
102
    {
Chao Liu's avatar
Chao Liu committed
103
104
        return Sequence<Is..., Xs...>{};
    }
Chao Liu's avatar
Chao Liu committed
105

Chao Liu's avatar
Chao Liu committed
106
    template <index_t... Ns>
107
    __host__ __device__ static constexpr auto Extract(Number<Ns>...)
Chao Liu's avatar
Chao Liu committed
108
    {
Chao Liu's avatar
Chao Liu committed
109
        return Sequence<Type::Get(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
110
    }
Chao Liu's avatar
Chao Liu committed
111

Chao Liu's avatar
Chao Liu committed
112
    template <index_t... Ns>
113
    __host__ __device__ static constexpr auto Extract(Sequence<Ns...>)
Chao Liu's avatar
Chao Liu committed
114
    {
Chao Liu's avatar
Chao Liu committed
115
        return Sequence<Type::Get(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
116
    }
117
118
119

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

Chao Liu's avatar
Chao Liu committed
122
// merge sequence
Chao Liu's avatar
Chao Liu committed
123
124
template <class, class>
struct sequence_merge;
Chao Liu's avatar
Chao Liu committed
125

Chao Liu's avatar
Chao Liu committed
126
127
128
template <index_t... Xs, index_t... Ys>
struct sequence_merge<Sequence<Xs...>, Sequence<Ys...>>
{
Chao Liu's avatar
Chao Liu committed
129
    using SeqType = Sequence<Xs..., Ys...>;
Chao Liu's avatar
Chao Liu committed
130
};
Chao Liu's avatar
Chao Liu committed
131

Chao Liu's avatar
Chao Liu committed
132
// arithmetic sqeuence
Chao Liu's avatar
Chao Liu committed
133
template <index_t IBegin, index_t NSize, index_t Increment>
134
struct arithmetic_sequence_gen_impl
Chao Liu's avatar
Chao Liu committed
135
136
{
    static constexpr index_t NSizeLeft = NSize / 2;
Chao Liu's avatar
Chao Liu committed
137

Chao Liu's avatar
Chao Liu committed
138
    using SeqType = typename sequence_merge<
139
140
        typename arithmetic_sequence_gen_impl<IBegin, NSizeLeft, Increment>::SeqType,
        typename arithmetic_sequence_gen_impl<IBegin + NSizeLeft * Increment,
Chao Liu's avatar
Chao Liu committed
141
142
                                              NSize - NSizeLeft,
                                              Increment>::SeqType>::SeqType;
Chao Liu's avatar
Chao Liu committed
143
144
};

Chao Liu's avatar
Chao Liu committed
145
template <index_t IBegin, index_t Increment>
146
struct arithmetic_sequence_gen_impl<IBegin, 1, Increment>
Chao Liu's avatar
Chao Liu committed
147
{
Chao Liu's avatar
Chao Liu committed
148
    using SeqType = Sequence<IBegin>;
Chao Liu's avatar
Chao Liu committed
149
};
Chao Liu's avatar
Chao Liu committed
150

Chao Liu's avatar
Chao Liu committed
151
template <index_t IBegin, index_t Increment>
152
struct arithmetic_sequence_gen_impl<IBegin, 0, Increment>
Chao Liu's avatar
Chao Liu committed
153
{
Chao Liu's avatar
Chao Liu committed
154
155
156
157
    using SeqType = Sequence<>;
};

template <index_t IBegin, index_t IEnd, index_t Increment>
158
struct arithmetic_sequence_gen
Chao Liu's avatar
Chao Liu committed
159
160
{
    using SeqType =
161
        typename arithmetic_sequence_gen_impl<IBegin, IEnd - IBegin, Increment>::SeqType;
Chao Liu's avatar
Chao Liu committed
162
};
Chao Liu's avatar
Chao Liu committed
163

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
// 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
189
template <class, class, index_t>
Chao Liu's avatar
Chao Liu committed
190
struct sequence_reverse_inclusive_scan;
Chao Liu's avatar
Chao Liu committed
191

Chao Liu's avatar
Chao Liu committed
192
193
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
194
{
Chao Liu's avatar
Chao Liu committed
195
196
    using old_scan =
        typename sequence_reverse_inclusive_scan<Sequence<Is...>, Reduce, Init>::SeqType;
Chao Liu's avatar
Chao Liu committed
197
198
199
200
201
202

    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
203
204
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
205
{
Chao Liu's avatar
Chao Liu committed
206
    using SeqType = Sequence<Reduce{}(I, Init)>;
Chao Liu's avatar
Chao Liu committed
207
208
};

Chao Liu's avatar
Chao Liu committed
209
210
template <class Reduce, index_t Init>
struct sequence_reverse_inclusive_scan<Sequence<>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
211
212
213
214
{
    using SeqType = Sequence<>;
};

Chao Liu's avatar
Chao Liu committed
215
// extract sequence
Chao Liu's avatar
Chao Liu committed
216
217
218
219
220
221
222
223
224
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
225
// split sequence
Chao Liu's avatar
Chao Liu committed
226
227
228
229
230
template <class Seq, index_t I>
struct sequence_split
{
    static constexpr index_t NSize = Seq{}.GetSize();

231
232
    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
233
234
235
236
237

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

Chao Liu's avatar
Chao Liu committed
238
// reverse sequence
Chao Liu's avatar
Chao Liu committed
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
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
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
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
#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
317

318
#endif
Chao Liu's avatar
Chao Liu committed
319

Chao Liu's avatar
Chao Liu committed
320
321
322
323
324
325
326
327
328
329
330
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
};
331

Chao Liu's avatar
Chao Liu committed
332
template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
333
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
334
335
336
337
338
339
340
{
    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
341
__host__ __device__ constexpr auto operator-(Sequence<Xs...> seq_x, Sequence<Ys...> seq_y)
Chao Liu's avatar
Chao Liu committed
342
343
344
345
346
347
348
{
    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
349
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
350
351
352
353
354
355
356
{
    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
357
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
358
359
360
361
362
363
364
{
    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
365
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
366
367
368
369
370
371
372
{
    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
373
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
374
{
Chao Liu's avatar
Chao Liu committed
375
    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
396
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
397
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
398
{
Chao Liu's avatar
Chao Liu committed
399
    return Sequence<(Xs % Y)...>{};
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
    return Sequence<(Y + Xs)...>{};
Chao Liu's avatar
Chao Liu committed
406
407
}

Chao Liu's avatar
Chao Liu committed
408
409
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
410
{
Chao Liu's avatar
Chao Liu committed
411
412
413
    constexpr auto seq_x = Sequence<Xs...>{};

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

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

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

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

434
435
436
437
438
439
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
440
441
template <class Seq>
__host__ __device__ constexpr auto sequence_pop_back(Seq)
442
{
443
    static_assert(Seq{}.GetSize() > 0, "wrong! cannot pop an empty Sequence!");
Chao Liu's avatar
Chao Liu committed
444
    return sequence_pop_front(Seq{}.Reverse()).Reverse();
445
}
446

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

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

461
462
463
464
465
466
467
468
469
470
471
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
472
473
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto reverse_inclusive_scan_sequence(Seq, Reduce, Number<Init>)
474
{
Chao Liu's avatar
Chao Liu committed
475
    return typename sequence_reverse_inclusive_scan<Seq, Reduce, Init>::SeqType{};
476
477
}

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

Chao Liu's avatar
Chao Liu committed
484
template <index_t... Is>
485
__host__ __device__ constexpr auto Sequence<Is...>::PopFront()
Chao Liu's avatar
Chao Liu committed
486
{
487
    return sequence_pop_front(Type{});
Chao Liu's avatar
Chao Liu committed
488
}
Chao Liu's avatar
Chao Liu committed
489

490
491
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::PopBack()
Chao Liu's avatar
Chao Liu committed
492
{
493
    return sequence_pop_back(Type{});
Chao Liu's avatar
Chao Liu committed
494
495
}

496
497
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Reverse()
Chao Liu's avatar
Chao Liu committed
498
{
499
500
501
502
503
504
505
506
507
508
509
510
511
512
    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
513
}
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549

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...); });
}