"test/vscode:/vscode.git/clone" did not exist on "fd7fcf939f214223ea4abe92b7c482cd5a17c579"
Sequence.hpp 13.3 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
        // 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
32
    __host__ __device__ constexpr auto operator[](Number<I>) const
Chao Liu's avatar
Chao Liu committed
33
    {
Chao Liu's avatar
Chao Liu committed
34
        return Number<Get(Number<I>{})>{};
Chao Liu's avatar
Chao Liu committed
35
36
37
38
39
    }

    // make sure I is constepxr
    __host__ __device__ constexpr index_t operator[](index_t I) const
    {
40
        const index_t mData[mSize + 1] = {Is..., 0};
41
42
43
        return mData[I];
    }

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

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

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

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

57
58
59
60
61
    __host__ __device__ static constexpr index_t Front()
    {
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[0];
    }
62

63
64
65
66
67
    __host__ __device__ static constexpr index_t Back()
    {
        const index_t mData[mSize + 1] = {Is..., 0};
        return mData[mSize - 1];
    }
68

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

    __host__ __device__ static constexpr auto PopBack();

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

} // namespace ck
#endif