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
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
        static_assert(I < mSize, "wrong! I too large");

Chao Liu's avatar
Chao Liu committed
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
        // 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
    {
36
        const index_t mData[mSize + 1] = {Is..., 0};
37
38
39
        return mData[I];
    }

40
    template <index_t... IRs>
41
    __host__ __device__ static constexpr auto ReorderGivenNew2Old(Sequence<IRs...> /*new2old*/)
42
    {
43
44
45
46
47
#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
48

Chao Liu's avatar
Chao Liu committed
49
50
51
        static_assert(sizeof...(Is) == sizeof...(IRs),
                      "wrong! new2old map should have the same size as Sequence to be rerodered");

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

55
56
57
#if 0 // require sequence_sort, which is not implemented yet
    template <class MapOld2New>
    __host__ __device__ static constexpr auto ReorderGivenOld2New(MapOld2New /*old2new*/)
58
    {
Chao Liu's avatar
Chao Liu committed
59
#if 0
60
61
62
        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
63
#endif
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
320
321
322
323
324
325
326
327
328
329
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
};
330

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

495
496
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Reverse()
Chao Liu's avatar
Chao Liu committed
497
{
498
499
500
501
502
503
504
505
506
507
508
509
510
511
    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
512
}
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

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