"test/git@developer.sourcefind.cn:zhaoyu6/sglang.git" did not exist on "2d54d4bb648dea1bcb13af2ddab30fcf4e1fa37f"
Sequence.hpp 13.6 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
16
    using Type      = Sequence;
    using data_type = index_t;
17

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

Chao Liu's avatar
Chao Liu committed
20
    __host__ __device__ static constexpr auto GetSize() { return Number<mSize>{}; }
21

Chao Liu's avatar
Chao Liu committed
22
    __host__ __device__ static constexpr index_t GetImpl(index_t I)
23
    {
Chao Liu's avatar
Chao Liu committed
24
25
26
27
28
29
        // 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>
Chao Liu's avatar
Chao Liu committed
30
    __host__ __device__ static constexpr auto Get(Number<I>)
Chao Liu's avatar
Chao Liu committed
31
    {
Chao Liu's avatar
Chao Liu committed
32
33
34
        static_assert(I < mSize, "wrong! I too large");

        return Number<GetImpl(Number<I>{})>{};
Chao Liu's avatar
Chao Liu committed
35
36
    }

Chao Liu's avatar
Chao Liu committed
37
38
    template <index_t I>
    __host__ __device__ constexpr auto operator[](Number<I>) const
Chao Liu's avatar
Chao Liu committed
39
    {
Chao Liu's avatar
Chao Liu committed
40
        return Get(Number<I>{});
41
42
    }

Chao Liu's avatar
Chao Liu committed
43
44
45
    // make sure I is constepxr if you want a constexpr return type
    __host__ __device__ constexpr index_t operator[](index_t I) const { return GetImpl(I); }

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

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

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

57
    __host__ __device__ static constexpr auto Reverse();
Chao Liu's avatar
Chao Liu committed
58

Chao Liu's avatar
Chao Liu committed
59
    __host__ __device__ static constexpr auto Front()
60
    {
Chao Liu's avatar
Chao Liu committed
61
62
        static_assert(mSize > 0, "wrong!");
        return Get(Number<0>{});
63
    }
64

Chao Liu's avatar
Chao Liu committed
65
    __host__ __device__ static constexpr auto Back()
66
    {
Chao Liu's avatar
Chao Liu committed
67
68
        static_assert(mSize > 0, "wrong!");
        return Get(Number<mSize - 1>{});
69
    }
70

Chao Liu's avatar
Chao Liu committed
71
72
73
74
75
76
    __host__ __device__ static constexpr auto PopFront();

    __host__ __device__ static constexpr auto PopBack();

    template <index_t... Xs>
    __host__ __device__ static constexpr auto PushFront(Sequence<Xs...>)
77
    {
Chao Liu's avatar
Chao Liu committed
78
        return Sequence<Xs..., Is...>{};
79
80
    }

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

Chao Liu's avatar
Chao Liu committed
87
88
89
90
91
    template <index_t... Xs>
    __host__ __device__ static constexpr auto PushBack(Sequence<Xs...>)
    {
        return Sequence<Is..., Xs...>{};
    }
92

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

Chao Liu's avatar
Chao Liu committed
99
    template <index_t... Ns>
100
    __host__ __device__ static constexpr auto Extract(Number<Ns>...)
Chao Liu's avatar
Chao Liu committed
101
    {
Chao Liu's avatar
Chao Liu committed
102
        return Sequence<Type::Get(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
103
    }
Chao Liu's avatar
Chao Liu committed
104

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

    template <index_t I, index_t X>
    __host__ __device__ static constexpr auto Modify(Number<I>, Number<X>);
Chao Liu's avatar
Chao Liu committed
113
114
115
116
117
118

    template <class F>
    __host__ __device__ static constexpr auto Transform(F f)
    {
        return Sequence<f(Is)...>{};
    }
119
120
};

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

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

Chao Liu's avatar
Chao Liu committed
131
132
133
// generate sequence
template <index_t IBegin, index_t NRemain, class F>
struct sequence_gen_impl
Chao Liu's avatar
Chao Liu committed
134
{
Chao Liu's avatar
Chao Liu committed
135
136
137
    static constexpr index_t NRemainLeft  = NRemain / 2;
    static constexpr index_t NRemainRight = NRemain - NRemainLeft;
    static constexpr index_t IMiddle      = IBegin + NRemainLeft;
Chao Liu's avatar
Chao Liu committed
138

Chao Liu's avatar
Chao Liu committed
139
140
141
    using type =
        typename sequence_merge<typename sequence_gen_impl<IBegin, NRemainLeft, F>::type,
                                typename sequence_gen_impl<IMiddle, NRemainRight, F>::type>::type;
Chao Liu's avatar
Chao Liu committed
142
143
};

Chao Liu's avatar
Chao Liu committed
144
145
template <index_t I, class F>
struct sequence_gen_impl<I, 1, F>
Chao Liu's avatar
Chao Liu committed
146
{
Chao Liu's avatar
Chao Liu committed
147
148
    static constexpr index_t Is = F{}(Number<I>{});
    using type                  = Sequence<Is>;
Chao Liu's avatar
Chao Liu committed
149
};
Chao Liu's avatar
Chao Liu committed
150

Chao Liu's avatar
Chao Liu committed
151
152
template <index_t I, class F>
struct sequence_gen_impl<I, 0, F>
Chao Liu's avatar
Chao Liu committed
153
{
Chao Liu's avatar
Chao Liu committed
154
    using type = Sequence<>;
Chao Liu's avatar
Chao Liu committed
155
156
};

Chao Liu's avatar
Chao Liu committed
157
158
159
160
161
162
163
template <index_t NSize, class F>
struct sequence_gen
{
    using type = typename sequence_gen_impl<0, NSize, F>::type;
};

// arithmetic sequence
Chao Liu's avatar
Chao Liu committed
164
template <index_t IBegin, index_t IEnd, index_t Increment>
165
struct arithmetic_sequence_gen
Chao Liu's avatar
Chao Liu committed
166
{
Chao Liu's avatar
Chao Liu committed
167
168
169
170
171
172
173
174
175
    struct F
    {
        __host__ __device__ constexpr index_t operator()(index_t i) const
        {
            return i * Increment + IBegin;
        }
    };

    using type = typename sequence_gen<(IEnd - IBegin) / Increment, F>::type;
Chao Liu's avatar
Chao Liu committed
176
177
178
179
180
181
};

// uniform sequence
template <index_t NSize, index_t I>
struct uniform_sequence_gen
{
Chao Liu's avatar
Chao Liu committed
182
    struct F
Chao Liu's avatar
Chao Liu committed
183
184
185
186
    {
        __host__ __device__ constexpr index_t operator()(index_t) const { return I; }
    };

Chao Liu's avatar
Chao Liu committed
187
    using type = typename sequence_gen<NSize, F>::type;
Chao Liu's avatar
Chao Liu committed
188
189
190
};

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

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

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

Chao Liu's avatar
Chao Liu committed
201
    using type = typename sequence_merge<Sequence<new_reduce>, old_scan>::type;
Chao Liu's avatar
Chao Liu committed
202
203
};

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

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

Chao Liu's avatar
Chao Liu committed
216
// split sequence
Chao Liu's avatar
Chao Liu committed
217
218
219
220
221
template <class Seq, index_t I>
struct sequence_split
{
    static constexpr index_t NSize = Seq{}.GetSize();

Chao Liu's avatar
Chao Liu committed
222
223
    using range0 = typename arithmetic_sequence_gen<0, I, 1>::type;
    using range1 = typename arithmetic_sequence_gen<I, NSize, 1>::type;
Chao Liu's avatar
Chao Liu committed
224

Chao Liu's avatar
Chao Liu committed
225
226
    using SeqType0 = decltype(Seq::Extract(range0{}));
    using SeqType1 = decltype(Seq::Extract(range1{}));
Chao Liu's avatar
Chao Liu committed
227
228
};

Chao Liu's avatar
Chao Liu committed
229
// reverse sequence
Chao Liu's avatar
Chao Liu committed
230
231
232
233
234
235
template <class Seq>
struct sequence_reverse
{
    static constexpr index_t NSize = Seq{}.GetSize();

    using seq_split = sequence_split<Seq, NSize / 2>;
Chao Liu's avatar
Chao Liu committed
236
237
238
    using type      = typename sequence_merge<
        typename sequence_reverse<typename seq_split::SeqType1>::type,
        typename sequence_reverse<typename seq_split::SeqType0>::type>::type;
Chao Liu's avatar
Chao Liu committed
239
240
241
242
243
};

template <index_t I>
struct sequence_reverse<Sequence<I>>
{
Chao Liu's avatar
Chao Liu committed
244
    using type = Sequence<I>;
Chao Liu's avatar
Chao Liu committed
245
246
247
248
249
};

template <index_t I0, index_t I1>
struct sequence_reverse<Sequence<I0, I1>>
{
Chao Liu's avatar
Chao Liu committed
250
    using type = Sequence<I1, I0>;
Chao Liu's avatar
Chao Liu committed
251
};
Chao Liu's avatar
Chao Liu committed
252

Chao Liu's avatar
Chao Liu committed
253
254
255
template <class Seq>
struct is_valid_sequence_map
{
Chao Liu's avatar
Chao Liu committed
256
    static constexpr integral_constant<bool, true> value = integral_constant<bool, true>{};
Chao Liu's avatar
Chao Liu committed
257
258
259

    // TODO: add proper check for is_valid, something like:
    // static constexpr bool value =
Chao Liu's avatar
Chao Liu committed
260
    //     is_same<typename arithmetic_sequence_gen<0, Seq::GetSize(), 1>::type,
Chao Liu's avatar
Chao Liu committed
261
    //             typename sequence_sort<Seq>::SortedSeqType>{};
Chao Liu's avatar
Chao Liu committed
262
};
263

Chao Liu's avatar
Chao Liu committed
264
template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
265
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
266
267
268
269
270
271
272
{
    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
273
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
274
275
276
277
278
279
280
{
    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
281
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
282
283
284
285
286
287
288
{
    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
289
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
290
291
292
293
294
295
296
{
    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
297
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
298
299
300
301
302
303
304
{
    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
305
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
306
{
Chao Liu's avatar
Chao Liu committed
307
    return Sequence<(Xs + Y)...>{};
Chao Liu's avatar
Chao Liu committed
308
309
310
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
311
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
312
{
Chao Liu's avatar
Chao Liu committed
313
    return Sequence<(Xs - Y)...>{};
Chao Liu's avatar
Chao Liu committed
314
315
316
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
317
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
318
{
Chao Liu's avatar
Chao Liu committed
319
    return Sequence<(Xs * Y)...>{};
Chao Liu's avatar
Chao Liu committed
320
321
322
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
323
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
324
{
Chao Liu's avatar
Chao Liu committed
325
    return Sequence<(Xs / Y)...>{};
Chao Liu's avatar
Chao Liu committed
326
327
328
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
329
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
330
{
Chao Liu's avatar
Chao Liu committed
331
    return Sequence<(Xs % Y)...>{};
Chao Liu's avatar
Chao Liu committed
332
333
}

Chao Liu's avatar
Chao Liu committed
334
335
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator+(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
336
{
Chao Liu's avatar
Chao Liu committed
337
    return Sequence<(Y + Xs)...>{};
Chao Liu's avatar
Chao Liu committed
338
339
}

Chao Liu's avatar
Chao Liu committed
340
341
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
342
{
Chao Liu's avatar
Chao Liu committed
343
344
345
    constexpr auto seq_x = Sequence<Xs...>{};

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

Chao Liu's avatar
Chao Liu committed
348
349
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator*(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
350
{
Chao Liu's avatar
Chao Liu committed
351
    return Sequence<(Y * Xs)...>{};
Chao Liu's avatar
Chao Liu committed
352
353
}

Chao Liu's avatar
Chao Liu committed
354
355
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator/(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
356
{
Chao Liu's avatar
Chao Liu committed
357
    return Sequence<(Y / Xs)...>{};
Chao Liu's avatar
Chao Liu committed
358
359
}

Chao Liu's avatar
Chao Liu committed
360
361
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator%(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
362
{
Chao Liu's avatar
Chao Liu committed
363
    return Sequence<(Y % Xs)...>{};
Chao Liu's avatar
Chao Liu committed
364
365
}

366
367
368
369
370
371
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
372
373
template <class Seq>
__host__ __device__ constexpr auto sequence_pop_back(Seq)
374
{
375
    static_assert(Seq{}.GetSize() > 0, "wrong! cannot pop an empty Sequence!");
Chao Liu's avatar
Chao Liu committed
376
    return sequence_pop_front(Seq{}.Reverse()).Reverse();
377
}
378

Chao Liu's avatar
Chao Liu committed
379
380
381
382
383
384
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
385
template <class F, index_t... Xs, index_t... Ys>
386
__host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>)
387
{
388
    static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize, "Dim not the same");
389
390
391
392

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

393
394
395
396
397
398
399
400
401
402
403
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
404
405
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto reverse_inclusive_scan_sequence(Seq, Reduce, Number<Init>)
406
{
Chao Liu's avatar
Chao Liu committed
407
    return typename sequence_reverse_inclusive_scan<Seq, Reduce, Init>::type{};
408
409
}

Chao Liu's avatar
Chao Liu committed
410
411
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto inclusive_scan_sequence(Seq, Reduce, Number<Init>)
412
{
Chao Liu's avatar
Chao Liu committed
413
    return reverse_inclusive_scan_sequence(Seq{}.Reverse(), Reduce{}, Number<Init>{}).Reverse();
414
}
415

Chao Liu's avatar
Chao Liu committed
416
template <index_t... Is>
417
__host__ __device__ constexpr auto Sequence<Is...>::PopFront()
Chao Liu's avatar
Chao Liu committed
418
{
419
    return sequence_pop_front(Type{});
Chao Liu's avatar
Chao Liu committed
420
}
Chao Liu's avatar
Chao Liu committed
421

422
423
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::PopBack()
Chao Liu's avatar
Chao Liu committed
424
{
425
    return sequence_pop_back(Type{});
Chao Liu's avatar
Chao Liu committed
426
427
}

428
429
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Reverse()
Chao Liu's avatar
Chao Liu committed
430
{
Chao Liu's avatar
Chao Liu committed
431
    return typename sequence_reverse<Sequence<Is...>>::type{};
432
433
434
435
436
437
438
439
440
441
442
443
}

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();

Chao Liu's avatar
Chao Liu committed
444
    return seq_left.PushBack(Number<X>{}).PushBack(seq_right);
Chao Liu's avatar
Chao Liu committed
445
}
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481

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

} // namespace ck
#endif