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
    using Type = Sequence;
16

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

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

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

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

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

Chao Liu's avatar
Chao Liu committed
42
43
44
    // 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); }

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

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

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

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

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

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

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

    __host__ __device__ static constexpr auto PopBack();

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

// 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
170
171
    using type = decltype(
        typename arithmetic_sequence_gen<0, NSize, 1>::type{}.Transform(return_constant{}));
Chao Liu's avatar
Chao Liu committed
172
173
174
};

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

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

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

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

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

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

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

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

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

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

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

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

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

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

Chao Liu's avatar
Chao Liu committed
248
template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
249
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
250
251
252
253
254
255
256
{
    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
257
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
258
259
260
261
262
263
264
{
    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
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 Y>
Chao Liu's avatar
Chao Liu committed
289
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
290
{
Chao Liu's avatar
Chao Liu committed
291
    return Sequence<(Xs + Y)...>{};
Chao Liu's avatar
Chao Liu committed
292
293
294
}

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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
428
    return seq_left.PushBack(Number<X>{}).PushBack(seq_right);
Chao Liu's avatar
Chao Liu committed
429
}
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

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

} // namespace ck
#endif