Sequence.hpp 15.8 KB
Newer Older
1
2
3
#ifndef CK_SEQUENCE_HPP
#define CK_SEQUENCE_HPP

Chao Liu's avatar
Chao Liu committed
4
5
#include "integral_constant.hpp"
#include "functional.hpp"
6

7
8
namespace ck {

Chao Liu's avatar
Chao Liu committed
9
10
11
template <class Seq>
struct is_valid_sequence_map;

Chao Liu's avatar
Chao Liu committed
12
template <index_t... Is>
13
14
struct Sequence
{
Chao Liu's avatar
Chao Liu committed
15
    using Type = Sequence;
16

17
    static constexpr index_t mSize = sizeof...(Is);
18

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

Chao Liu's avatar
Chao Liu committed
21
    template <index_t I>
22
    __host__ __device__ static constexpr index_t Get(Number<I>)
23
    {
24
25
        static_assert(I < mSize, "wrong! I too large");

Chao Liu's avatar
Chao Liu committed
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
        // 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
    {
43
        const index_t mData[mSize + 1] = {Is..., 0};
44
45
46
        return mData[I];
    }

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

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

        return Sequence<Type::Get(Number<IRs>{})...>{};
56
57
    }

58
59
60
#if 0 // require sequence_sort, which is not implemented yet
    template <class MapOld2New>
    __host__ __device__ static constexpr auto ReorderGivenOld2New(MapOld2New /*old2new*/)
61
    {
Chao Liu's avatar
Chao Liu committed
62
63
64
65
66
67
        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");

68
69
70
        constexpr auto map_new2old = typename sequence_map_inverse<MapOld2New>::SeqMapType{};

        return ReorderGivenNew2Old(map_new2old);
71
    }
72
#endif
73

74
    __host__ __device__ static constexpr auto Reverse();
Chao Liu's avatar
Chao Liu committed
75

76
77
78
79
80
    __host__ __device__ static constexpr index_t Front()
    {
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[0];
    }
81

82
83
84
85
86
    __host__ __device__ static constexpr index_t Back()
    {
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[mSize - 1];
    }
87

88
    template <index_t I>
89
    __host__ __device__ static constexpr auto PushFront(Number<I>)
90
91
92
93
    {
        return Sequence<I, Is...>{};
    }

Chao Liu's avatar
Chao Liu committed
94
    template <index_t I>
95
    __host__ __device__ static constexpr auto PushBack(Number<I>)
96
97
98
99
    {
        return Sequence<Is..., I>{};
    }

100
    __host__ __device__ static constexpr auto PopFront();
101

102
    __host__ __device__ static constexpr auto PopBack();
103

Chao Liu's avatar
Chao Liu committed
104
    template <index_t... Xs>
105
    __host__ __device__ static constexpr auto Append(Sequence<Xs...>)
106
    {
Chao Liu's avatar
Chao Liu committed
107
108
        return Sequence<Is..., Xs...>{};
    }
Chao Liu's avatar
Chao Liu committed
109

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

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

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

Chao Liu's avatar
Chao Liu committed
126
// merge sequence
Chao Liu's avatar
Chao Liu committed
127
128
template <class, class>
struct sequence_merge;
Chao Liu's avatar
Chao Liu committed
129

Chao Liu's avatar
Chao Liu committed
130
131
132
template <index_t... Xs, index_t... Ys>
struct sequence_merge<Sequence<Xs...>, Sequence<Ys...>>
{
Chao Liu's avatar
Chao Liu committed
133
    using SeqType = Sequence<Xs..., Ys...>;
Chao Liu's avatar
Chao Liu committed
134
};
Chao Liu's avatar
Chao Liu committed
135

Chao Liu's avatar
Chao Liu committed
136
// arithmetic sqeuence
Chao Liu's avatar
Chao Liu committed
137
template <index_t IBegin, index_t NSize, index_t Increment>
138
struct arithmetic_sequence_gen_impl
Chao Liu's avatar
Chao Liu committed
139
140
{
    static constexpr index_t NSizeLeft = NSize / 2;
Chao Liu's avatar
Chao Liu committed
141

Chao Liu's avatar
Chao Liu committed
142
    using SeqType = typename sequence_merge<
143
144
        typename arithmetic_sequence_gen_impl<IBegin, NSizeLeft, Increment>::SeqType,
        typename arithmetic_sequence_gen_impl<IBegin + NSizeLeft * Increment,
Chao Liu's avatar
Chao Liu committed
145
146
                                              NSize - NSizeLeft,
                                              Increment>::SeqType>::SeqType;
Chao Liu's avatar
Chao Liu committed
147
148
};

Chao Liu's avatar
Chao Liu committed
149
template <index_t IBegin, index_t Increment>
150
struct arithmetic_sequence_gen_impl<IBegin, 1, Increment>
Chao Liu's avatar
Chao Liu committed
151
{
Chao Liu's avatar
Chao Liu committed
152
    using SeqType = Sequence<IBegin>;
Chao Liu's avatar
Chao Liu committed
153
};
Chao Liu's avatar
Chao Liu committed
154

Chao Liu's avatar
Chao Liu committed
155
template <index_t IBegin, index_t Increment>
156
struct arithmetic_sequence_gen_impl<IBegin, 0, Increment>
Chao Liu's avatar
Chao Liu committed
157
{
Chao Liu's avatar
Chao Liu committed
158
159
160
161
    using SeqType = Sequence<>;
};

template <index_t IBegin, index_t IEnd, index_t Increment>
162
struct arithmetic_sequence_gen
Chao Liu's avatar
Chao Liu committed
163
164
{
    using SeqType =
165
        typename arithmetic_sequence_gen_impl<IBegin, IEnd - IBegin, Increment>::SeqType;
Chao Liu's avatar
Chao Liu committed
166
};
Chao Liu's avatar
Chao Liu committed
167

Chao Liu's avatar
Chao Liu committed
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
// 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
193
template <class, class, index_t>
Chao Liu's avatar
Chao Liu committed
194
struct sequence_reverse_inclusive_scan;
Chao Liu's avatar
Chao Liu committed
195

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

    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
207
208
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
209
{
Chao Liu's avatar
Chao Liu committed
210
    using SeqType = Sequence<Reduce{}(I, Init)>;
Chao Liu's avatar
Chao Liu committed
211
212
};

Chao Liu's avatar
Chao Liu committed
213
214
template <class Reduce, index_t Init>
struct sequence_reverse_inclusive_scan<Sequence<>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
215
216
217
218
{
    using SeqType = Sequence<>;
};

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

235
236
    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
237
238
239
240
241

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

Chao Liu's avatar
Chao Liu committed
242
// reverse sequence
Chao Liu's avatar
Chao Liu committed
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
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
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
317
318
319
320
#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
321

322
#endif
Chao Liu's avatar
Chao Liu committed
323

Chao Liu's avatar
Chao Liu committed
324
325
326
327
328
329
330
331
332
333
334
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
};
335

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

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

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

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

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

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

Chao Liu's avatar
Chao Liu committed
412
413
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
414
{
Chao Liu's avatar
Chao Liu committed
415
416
417
    constexpr auto seq_x = Sequence<Xs...>{};

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

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

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

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

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

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

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

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

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

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

494
495
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::PopBack()
Chao Liu's avatar
Chao Liu committed
496
{
497
    return sequence_pop_back(Type{});
Chao Liu's avatar
Chao Liu committed
498
499
}

500
501
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Reverse()
Chao Liu's avatar
Chao Liu committed
502
{
503
504
505
506
507
508
509
510
511
512
513
514
515
516
    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
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
550
551
552
553

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

} // namespace ck
#endif