Sequence.hpp 13.4 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
// arithmetic sqeuence
Chao Liu's avatar
Chao Liu committed
132
template <index_t IBegin, index_t NSize, index_t Increment>
133
struct arithmetic_sequence_gen_impl
Chao Liu's avatar
Chao Liu committed
134
135
{
    static constexpr index_t NSizeLeft = NSize / 2;
Chao Liu's avatar
Chao Liu committed
136

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

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

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

template <index_t IBegin, index_t IEnd, index_t Increment>
157
struct arithmetic_sequence_gen
Chao Liu's avatar
Chao Liu committed
158
{
Chao Liu's avatar
Chao Liu committed
159
    using type = typename arithmetic_sequence_gen_impl<IBegin, IEnd - IBegin, Increment>::type;
Chao Liu's avatar
Chao Liu committed
160
161
162
163
164
165
166
167
168
169
170
};

// 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; }
    };

Chao Liu's avatar
Chao Liu committed
171
172
    using type = decltype(
        typename arithmetic_sequence_gen<0, NSize, 1>::type{}.Transform(return_constant{}));
Chao Liu's avatar
Chao Liu committed
173
174
175
};

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

Chao Liu's avatar
Chao Liu committed
179
180
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
181
{
Chao Liu's avatar
Chao Liu committed
182
    using old_scan = typename sequence_reverse_inclusive_scan<Sequence<Is...>, Reduce, Init>::type;
Chao Liu's avatar
Chao Liu committed
183
184
185

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

Chao Liu's avatar
Chao Liu committed
186
    using type = typename sequence_merge<Sequence<new_reduce>, old_scan>::type;
Chao Liu's avatar
Chao Liu committed
187
188
};

Chao Liu's avatar
Chao Liu committed
189
190
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
191
{
Chao Liu's avatar
Chao Liu committed
192
    using type = Sequence<Reduce{}(I, Init)>;
Chao Liu's avatar
Chao Liu committed
193
194
};

Chao Liu's avatar
Chao Liu committed
195
196
template <class Reduce, index_t Init>
struct sequence_reverse_inclusive_scan<Sequence<>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
197
{
Chao Liu's avatar
Chao Liu committed
198
    using type = Sequence<>;
Chao Liu's avatar
Chao Liu committed
199
200
};

Chao Liu's avatar
Chao Liu committed
201
// split sequence
Chao Liu's avatar
Chao Liu committed
202
203
204
205
206
template <class Seq, index_t I>
struct sequence_split
{
    static constexpr index_t NSize = Seq{}.GetSize();

Chao Liu's avatar
Chao Liu committed
207
208
    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
209

Chao Liu's avatar
Chao Liu committed
210
211
    using SeqType0 = decltype(Seq::Extract(range0{}));
    using SeqType1 = decltype(Seq::Extract(range1{}));
Chao Liu's avatar
Chao Liu committed
212
213
};

Chao Liu's avatar
Chao Liu committed
214
// reverse sequence
Chao Liu's avatar
Chao Liu committed
215
216
217
218
219
220
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
221
222
223
    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
224
225
226
227
228
};

template <index_t I>
struct sequence_reverse<Sequence<I>>
{
Chao Liu's avatar
Chao Liu committed
229
    using type = Sequence<I>;
Chao Liu's avatar
Chao Liu committed
230
231
232
233
234
};

template <index_t I0, index_t I1>
struct sequence_reverse<Sequence<I0, I1>>
{
Chao Liu's avatar
Chao Liu committed
235
    using type = Sequence<I1, I0>;
Chao Liu's avatar
Chao Liu committed
236
};
Chao Liu's avatar
Chao Liu committed
237

Chao Liu's avatar
Chao Liu committed
238
239
240
template <class Seq>
struct is_valid_sequence_map
{
Chao Liu's avatar
Chao Liu committed
241
    static constexpr integral_constant<bool, true> value = integral_constant<bool, true>{};
Chao Liu's avatar
Chao Liu committed
242
243
244

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

Chao Liu's avatar
Chao Liu committed
249
template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
250
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
251
252
253
254
255
256
257
{
    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
258
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
259
260
261
262
263
264
265
{
    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
266
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
267
268
269
270
271
272
273
{
    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
274
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
275
276
277
278
279
280
281
{
    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
282
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
283
284
285
286
287
288
289
{
    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
290
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
291
{
Chao Liu's avatar
Chao Liu committed
292
    return Sequence<(Xs + Y)...>{};
Chao Liu's avatar
Chao Liu committed
293
294
295
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
296
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
297
{
Chao Liu's avatar
Chao Liu committed
298
    return Sequence<(Xs - Y)...>{};
Chao Liu's avatar
Chao Liu committed
299
300
301
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
302
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
303
{
Chao Liu's avatar
Chao Liu committed
304
    return Sequence<(Xs * Y)...>{};
Chao Liu's avatar
Chao Liu committed
305
306
307
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
308
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
309
{
Chao Liu's avatar
Chao Liu committed
310
    return Sequence<(Xs / Y)...>{};
Chao Liu's avatar
Chao Liu committed
311
312
313
}

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

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

Chao Liu's avatar
Chao Liu committed
325
326
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
327
{
Chao Liu's avatar
Chao Liu committed
328
329
330
    constexpr auto seq_x = Sequence<Xs...>{};

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

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

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

Chao Liu's avatar
Chao Liu committed
345
346
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator%(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
347
{
Chao Liu's avatar
Chao Liu committed
348
    return Sequence<(Y % Xs)...>{};
Chao Liu's avatar
Chao Liu committed
349
350
}

351
352
353
354
355
356
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
357
358
template <class Seq>
__host__ __device__ constexpr auto sequence_pop_back(Seq)
359
{
360
    static_assert(Seq{}.GetSize() > 0, "wrong! cannot pop an empty Sequence!");
Chao Liu's avatar
Chao Liu committed
361
    return sequence_pop_front(Seq{}.Reverse()).Reverse();
362
}
363

Chao Liu's avatar
Chao Liu committed
364
365
366
367
368
369
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
370
template <class F, index_t... Xs, index_t... Ys>
371
__host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>)
372
{
373
    static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize, "Dim not the same");
374
375
376
377

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

378
379
380
381
382
383
384
385
386
387
388
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
389
390
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto reverse_inclusive_scan_sequence(Seq, Reduce, Number<Init>)
391
{
Chao Liu's avatar
Chao Liu committed
392
    return typename sequence_reverse_inclusive_scan<Seq, Reduce, Init>::type{};
393
394
}

Chao Liu's avatar
Chao Liu committed
395
396
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto inclusive_scan_sequence(Seq, Reduce, Number<Init>)
397
{
Chao Liu's avatar
Chao Liu committed
398
    return reverse_inclusive_scan_sequence(Seq{}.Reverse(), Reduce{}, Number<Init>{}).Reverse();
399
}
400

Chao Liu's avatar
Chao Liu committed
401
template <index_t... Is>
402
__host__ __device__ constexpr auto Sequence<Is...>::PopFront()
Chao Liu's avatar
Chao Liu committed
403
{
404
    return sequence_pop_front(Type{});
Chao Liu's avatar
Chao Liu committed
405
}
Chao Liu's avatar
Chao Liu committed
406

407
408
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::PopBack()
Chao Liu's avatar
Chao Liu committed
409
{
410
    return sequence_pop_back(Type{});
Chao Liu's avatar
Chao Liu committed
411
412
}

413
414
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Reverse()
Chao Liu's avatar
Chao Liu committed
415
{
Chao Liu's avatar
Chao Liu committed
416
    return typename sequence_reverse<Sequence<Is...>>::type{};
417
418
419
420
421
422
423
424
425
426
427
428
}

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
429
    return seq_left.PushBack(Number<X>{}).PushBack(seq_right);
Chao Liu's avatar
Chao Liu committed
430
}
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466

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

} // namespace ck
#endif