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
    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
    __host__ __device__ static constexpr auto Reverse();
Chao Liu's avatar
Chao Liu committed
59

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

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

72
    template <index_t I>
73
    __host__ __device__ static constexpr auto PushFront(Number<I>)
74
75
76
77
    {
        return Sequence<I, Is...>{};
    }

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

84
    __host__ __device__ static constexpr auto PopFront();
85

86
    __host__ __device__ static constexpr auto PopBack();
87

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

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

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

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

Chao Liu's avatar
Chao Liu committed
110
// merge sequence
Chao Liu's avatar
Chao Liu committed
111
112
template <class, class>
struct sequence_merge;
Chao Liu's avatar
Chao Liu committed
113

Chao Liu's avatar
Chao Liu committed
114
115
116
template <index_t... Xs, index_t... Ys>
struct sequence_merge<Sequence<Xs...>, Sequence<Ys...>>
{
Chao Liu's avatar
Chao Liu committed
117
    using SeqType = Sequence<Xs..., Ys...>;
Chao Liu's avatar
Chao Liu committed
118
};
Chao Liu's avatar
Chao Liu committed
119

Chao Liu's avatar
Chao Liu committed
120
// arithmetic sqeuence
Chao Liu's avatar
Chao Liu committed
121
template <index_t IBegin, index_t NSize, index_t Increment>
122
struct arithmetic_sequence_gen_impl
Chao Liu's avatar
Chao Liu committed
123
124
{
    static constexpr index_t NSizeLeft = NSize / 2;
Chao Liu's avatar
Chao Liu committed
125

Chao Liu's avatar
Chao Liu committed
126
    using SeqType = typename sequence_merge<
127
128
        typename arithmetic_sequence_gen_impl<IBegin, NSizeLeft, Increment>::SeqType,
        typename arithmetic_sequence_gen_impl<IBegin + NSizeLeft * Increment,
Chao Liu's avatar
Chao Liu committed
129
130
                                              NSize - NSizeLeft,
                                              Increment>::SeqType>::SeqType;
Chao Liu's avatar
Chao Liu committed
131
132
};

Chao Liu's avatar
Chao Liu committed
133
template <index_t IBegin, index_t Increment>
134
struct arithmetic_sequence_gen_impl<IBegin, 1, Increment>
Chao Liu's avatar
Chao Liu committed
135
{
Chao Liu's avatar
Chao Liu committed
136
    using SeqType = Sequence<IBegin>;
Chao Liu's avatar
Chao Liu committed
137
};
Chao Liu's avatar
Chao Liu committed
138

Chao Liu's avatar
Chao Liu committed
139
template <index_t IBegin, index_t Increment>
140
struct arithmetic_sequence_gen_impl<IBegin, 0, Increment>
Chao Liu's avatar
Chao Liu committed
141
{
Chao Liu's avatar
Chao Liu committed
142
143
144
145
    using SeqType = Sequence<>;
};

template <index_t IBegin, index_t IEnd, index_t Increment>
146
struct arithmetic_sequence_gen
Chao Liu's avatar
Chao Liu committed
147
148
{
    using SeqType =
149
        typename arithmetic_sequence_gen_impl<IBegin, IEnd - IBegin, Increment>::SeqType;
Chao Liu's avatar
Chao Liu committed
150
};
Chao Liu's avatar
Chao Liu committed
151

Chao Liu's avatar
Chao Liu committed
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
// 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
177
template <class, class, index_t>
Chao Liu's avatar
Chao Liu committed
178
struct sequence_reverse_inclusive_scan;
Chao Liu's avatar
Chao Liu committed
179

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

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

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

Chao Liu's avatar
Chao Liu committed
203
// extract sequence
Chao Liu's avatar
Chao Liu committed
204
205
206
207
208
209
210
211
212
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
213
// split sequence
Chao Liu's avatar
Chao Liu committed
214
215
216
217
218
template <class Seq, index_t I>
struct sequence_split
{
    static constexpr index_t NSize = Seq{}.GetSize();

219
220
    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
221
222
223
224
225

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

Chao Liu's avatar
Chao Liu committed
226
// reverse sequence
Chao Liu's avatar
Chao Liu committed
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
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
249

Chao Liu's avatar
Chao Liu committed
250
251
252
template <class Seq>
struct is_valid_sequence_map
{
Chao Liu's avatar
Chao Liu committed
253
254
255
256
257
258
    static constexpr bool value = true;

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

Chao Liu's avatar
Chao Liu committed
261
template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
262
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
263
264
265
266
267
268
269
{
    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
270
__host__ __device__ constexpr auto operator-(Sequence<Xs...> seq_x, Sequence<Ys...> seq_y)
Chao Liu's avatar
Chao Liu committed
271
272
273
274
275
276
277
{
    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
278
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
279
280
281
282
283
284
285
{
    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
286
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
287
288
289
290
291
292
293
{
    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
294
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
295
296
297
298
299
300
301
{
    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
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
319
}

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

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
326
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
327
{
Chao Liu's avatar
Chao Liu committed
328
    return Sequence<(Xs % Y)...>{};
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
341
342
    constexpr auto seq_x = Sequence<Xs...>{};

    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
}

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

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

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

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

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

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

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

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

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

425
426
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Reverse()
Chao Liu's avatar
Chao Liu committed
427
{
428
429
430
431
432
433
434
435
436
437
438
439
440
441
    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
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
467
468
469
470
471
472
473
474
475
476
477
478

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...); });
}
479
480
481

} // namespace ck
#endif