sequence.hpp 30.4 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
// SPDX-License-Identifier: MIT
Illia Silin's avatar
Illia Silin committed
2
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
Chao Liu's avatar
Chao Liu committed
3

4
#pragma once
5

6
7
8
9
#include "ck/utility/integral_constant.hpp"
#include "ck/utility/type.hpp"
#include "ck/utility/functional.hpp"
#include "ck/utility/math.hpp"
10

11
12
namespace ck {

Chao Liu's avatar
Chao Liu committed
13
14
15
template <index_t, index_t, index_t>
struct static_for;

16
17
18
template <index_t...>
struct Sequence;

Chao Liu's avatar
Chao Liu committed
19
template <typename Seq, index_t I>
20
21
struct sequence_split;

Chao Liu's avatar
Chao Liu committed
22
template <typename>
23
struct sequence_reverse;
Chao Liu's avatar
Chao Liu committed
24

Chao Liu's avatar
Chao Liu committed
25
template <typename>
Chao Liu's avatar
Chao Liu committed
26
27
struct sequence_map_inverse;

Chao Liu's avatar
Chao Liu committed
28
template <typename>
29
30
31
32
33
struct is_valid_sequence_map;

template <index_t I, index_t... Is>
__host__ __device__ constexpr auto sequence_pop_front(Sequence<I, Is...>);

Chao Liu's avatar
Chao Liu committed
34
template <typename Seq>
35
36
__host__ __device__ constexpr auto sequence_pop_back(Seq);

Chao Liu's avatar
Chao Liu committed
37
template <index_t... Is>
38
39
struct Sequence
{
Chao Liu's avatar
Chao Liu committed
40
41
    using Type      = Sequence;
    using data_type = index_t;
42

43
    static constexpr index_t mSize = sizeof...(Is);
44

Chao Liu's avatar
Chao Liu committed
45
    __host__ __device__ static constexpr index_t Size() { return mSize; }
46

Chao Liu's avatar
Chao Liu committed
47
    // TODO: deprecate
Chao Liu's avatar
Chao Liu committed
48
49
50
    __host__ __device__ static constexpr auto GetSize() { return Size(); }

    __host__ __device__ static constexpr index_t At(index_t I)
51
    {
Chao Liu's avatar
Chao Liu committed
52
53
54
55
56
57
        // 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
58
    __host__ __device__ static constexpr auto At(Number<I>)
Chao Liu's avatar
Chao Liu committed
59
    {
Chao Liu's avatar
Chao Liu committed
60
61
        static_assert(I < mSize, "wrong! I too large");

Chao Liu's avatar
Chao Liu committed
62
        return Number<At(I)>{};
Chao Liu's avatar
Chao Liu committed
63
64
    }

Chao Liu's avatar
Chao Liu committed
65
66
67
68
69
70
71
72
    template <index_t I>
    __host__ __device__ static constexpr auto At()
    {
        static_assert(I < mSize, "wrong! I too large");

        return Number<At(I)>{};
    }

Chao Liu's avatar
Chao Liu committed
73
    template <index_t I>
Chao Liu's avatar
Chao Liu committed
74
    __host__ __device__ static constexpr auto Get(Number<I>)
Chao Liu's avatar
Chao Liu committed
75
    {
Chao Liu's avatar
Chao Liu committed
76
        return At(I);
77
78
    }

Chao Liu's avatar
Chao Liu committed
79
80
81
82
83
    template <typename I>
    __host__ __device__ constexpr auto operator[](I i) const
    {
        return At(i);
    }
Chao Liu's avatar
Chao Liu committed
84

85
    template <index_t... IRs>
86
    __host__ __device__ static constexpr auto ReorderGivenNew2Old(Sequence<IRs...> /*new2old*/)
87
    {
Chao Liu's avatar
Chao Liu committed
88
        static_assert(sizeof...(Is) == sizeof...(IRs),
Chao Liu's avatar
Chao Liu committed
89
                      "wrong! reorder map should have the same size as Sequence to be rerodered");
Chao Liu's avatar
Chao Liu committed
90

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

Chao Liu's avatar
Chao Liu committed
93
        return Sequence<Type::At(Number<IRs>{})...>{};
94
95
    }

Chao Liu's avatar
Chao Liu committed
96
    // MapOld2New is Sequence<...>
Chao Liu's avatar
Chao Liu committed
97
    template <typename MapOld2New>
Chao Liu's avatar
Chao Liu committed
98
99
    __host__ __device__ static constexpr auto ReorderGivenOld2New(MapOld2New)
    {
Chao Liu's avatar
Chao Liu committed
100
        static_assert(MapOld2New::Size() == Size(),
Chao Liu's avatar
Chao Liu committed
101
102
103
104
105
106
107
                      "wrong! reorder map should have the same size as Sequence to be rerodered");

        static_assert(is_valid_sequence_map<MapOld2New>::value, "wrong! invalid reorder map");

        return ReorderGivenNew2Old(typename sequence_map_inverse<MapOld2New>::type{});
    }

108
109
110
111
    __host__ __device__ static constexpr auto Reverse()
    {
        return typename sequence_reverse<Type>::type{};
    }
Chao Liu's avatar
Chao Liu committed
112

Chao Liu's avatar
Chao Liu committed
113
    __host__ __device__ static constexpr auto Front()
114
    {
Chao Liu's avatar
Chao Liu committed
115
        static_assert(mSize > 0, "wrong!");
Chao Liu's avatar
Chao Liu committed
116
        return At(Number<0>{});
117
    }
118

Chao Liu's avatar
Chao Liu committed
119
    __host__ __device__ static constexpr auto Back()
120
    {
Chao Liu's avatar
Chao Liu committed
121
        static_assert(mSize > 0, "wrong!");
Chao Liu's avatar
Chao Liu committed
122
        return At(Number<mSize - 1>{});
123
    }
124

125
    __host__ __device__ static constexpr auto PopFront() { return sequence_pop_front(Type{}); }
Chao Liu's avatar
Chao Liu committed
126

127
    __host__ __device__ static constexpr auto PopBack() { return sequence_pop_back(Type{}); }
Chao Liu's avatar
Chao Liu committed
128
129
130

    template <index_t... Xs>
    __host__ __device__ static constexpr auto PushFront(Sequence<Xs...>)
131
    {
Chao Liu's avatar
Chao Liu committed
132
        return Sequence<Xs..., Is...>{};
133
134
    }

Chao Liu's avatar
Chao Liu committed
135
136
    template <index_t... Xs>
    __host__ __device__ static constexpr auto PushFront(Number<Xs>...)
137
    {
Chao Liu's avatar
Chao Liu committed
138
        return Sequence<Xs..., Is...>{};
139
140
    }

Chao Liu's avatar
Chao Liu committed
141
142
143
144
145
    template <index_t... Xs>
    __host__ __device__ static constexpr auto PushBack(Sequence<Xs...>)
    {
        return Sequence<Is..., Xs...>{};
    }
146

Chao Liu's avatar
Chao Liu committed
147
    template <index_t... Xs>
Chao Liu's avatar
Chao Liu committed
148
    __host__ __device__ static constexpr auto PushBack(Number<Xs>...)
149
    {
Chao Liu's avatar
Chao Liu committed
150
151
        return Sequence<Is..., Xs...>{};
    }
Chao Liu's avatar
Chao Liu committed
152

Chao Liu's avatar
Chao Liu committed
153
    template <index_t... Ns>
154
    __host__ __device__ static constexpr auto Extract(Number<Ns>...)
Chao Liu's avatar
Chao Liu committed
155
    {
Chao Liu's avatar
Chao Liu committed
156
        return Sequence<Type::At(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
157
    }
Chao Liu's avatar
Chao Liu committed
158

Chao Liu's avatar
Chao Liu committed
159
    template <index_t... Ns>
160
    __host__ __device__ static constexpr auto Extract(Sequence<Ns...>)
Chao Liu's avatar
Chao Liu committed
161
    {
Chao Liu's avatar
Chao Liu committed
162
        return Sequence<Type::At(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
163
    }
164
165

    template <index_t I, index_t X>
166
167
    __host__ __device__ static constexpr auto Modify(Number<I>, Number<X>)
    {
Chao Liu's avatar
Chao Liu committed
168
        static_assert(I < Size(), "wrong!");
169
170

        using seq_split          = sequence_split<Type, I>;
Chao Liu's avatar
Chao Liu committed
171
172
        constexpr auto seq_left  = typename seq_split::left_type{};
        constexpr auto seq_right = typename seq_split::right_type{}.PopFront();
173
174
175

        return seq_left.PushBack(Number<X>{}).PushBack(seq_right);
    }
Chao Liu's avatar
Chao Liu committed
176

Chao Liu's avatar
Chao Liu committed
177
    template <typename F>
Chao Liu's avatar
Chao Liu committed
178
179
180
181
    __host__ __device__ static constexpr auto Transform(F f)
    {
        return Sequence<f(Is)...>{};
    }
Chao Liu's avatar
Chao Liu committed
182

Chao Liu's avatar
Chao Liu committed
183
184
    __host__ __device__ static constexpr bool IsStatic() { return true; };

Chao Liu's avatar
Chao Liu committed
185
186
    __host__ __device__ static void Print()
    {
Chao Liu's avatar
Chao Liu committed
187
188
189
190
191
192
193
194
195
196
197
198
199
        printf("Sequence{size: %d, data: [", Size());

        for(index_t i = 0; i < Size(); i++)
        {
            print(At(i));

            if(i < Size() - 1)
            {
                printf(", ");
            }
        }

        printf("]}");
Chao Liu's avatar
Chao Liu committed
200
    }
201
202
};

Chao Liu's avatar
Chao Liu committed
203
// merge sequence
Chao Liu's avatar
Chao Liu committed
204
205
206
207
208
template <typename Seq, typename... Seqs>
struct sequence_merge
{
    using type = typename sequence_merge<Seq, typename sequence_merge<Seqs...>::type>::type;
};
Chao Liu's avatar
Chao Liu committed
209

Chao Liu's avatar
Chao Liu committed
210
211
212
template <index_t... Xs, index_t... Ys>
struct sequence_merge<Sequence<Xs...>, Sequence<Ys...>>
{
Chao Liu's avatar
Chao Liu committed
213
    using type = Sequence<Xs..., Ys...>;
Chao Liu's avatar
Chao Liu committed
214
};
Chao Liu's avatar
Chao Liu committed
215

Chao Liu's avatar
Chao Liu committed
216
217
218
219
220
221
template <typename Seq>
struct sequence_merge<Seq>
{
    using type = Seq;
};

Chao Liu's avatar
Chao Liu committed
222
// generate sequence
Chao Liu's avatar
Chao Liu committed
223
224
template <index_t NSize, typename F>
struct sequence_gen
Chao Liu's avatar
Chao Liu committed
225
{
Chao Liu's avatar
Chao Liu committed
226
227
228
229
230
231
    template <index_t IBegin, index_t NRemain, typename G>
    struct sequence_gen_impl
    {
        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
232

Chao Liu's avatar
Chao Liu committed
233
234
235
236
        using type = typename sequence_merge<
            typename sequence_gen_impl<IBegin, NRemainLeft, G>::type,
            typename sequence_gen_impl<IMiddle, NRemainRight, G>::type>::type;
    };
Chao Liu's avatar
Chao Liu committed
237

Chao Liu's avatar
Chao Liu committed
238
239
240
241
242
243
    template <index_t I, typename G>
    struct sequence_gen_impl<I, 1, G>
    {
        static constexpr index_t Is = G{}(Number<I>{});
        using type                  = Sequence<Is>;
    };
Chao Liu's avatar
Chao Liu committed
244

Chao Liu's avatar
Chao Liu committed
245
246
247
248
249
    template <index_t I, typename G>
    struct sequence_gen_impl<I, 0, G>
    {
        using type = Sequence<>;
    };
Chao Liu's avatar
Chao Liu committed
250

Chao Liu's avatar
Chao Liu committed
251
252
253
254
    using type = typename sequence_gen_impl<0, NSize, F>::type;
};

// arithmetic sequence
Chao Liu's avatar
Chao Liu committed
255
template <index_t IBegin, index_t IEnd, index_t Increment>
256
struct arithmetic_sequence_gen
Chao Liu's avatar
Chao Liu committed
257
{
Chao Liu's avatar
Chao Liu committed
258
259
260
261
262
263
264
265
    struct F
    {
        __host__ __device__ constexpr index_t operator()(index_t i) const
        {
            return i * Increment + IBegin;
        }
    };

266
267
268
269
270
271
272
    using type0 = typename sequence_gen<(IEnd - IBegin) / Increment, F>::type;
    using type1 = Sequence<>;

    static constexpr bool kHasContent =
        (Increment > 0 && IBegin < IEnd) || (Increment < 0 && IBegin > IEnd);

    using type = typename conditional<kHasContent, type0, type1>::type;
Chao Liu's avatar
Chao Liu committed
273
274
275
276
277
278
};

// uniform sequence
template <index_t NSize, index_t I>
struct uniform_sequence_gen
{
Chao Liu's avatar
Chao Liu committed
279
    struct F
Chao Liu's avatar
Chao Liu committed
280
281
282
283
    {
        __host__ __device__ constexpr index_t operator()(index_t) const { return I; }
    };

Chao Liu's avatar
Chao Liu committed
284
    using type = typename sequence_gen<NSize, F>::type;
Chao Liu's avatar
Chao Liu committed
285
286
287
};

// reverse inclusive scan (with init) sequence
Chao Liu's avatar
Chao Liu committed
288
template <typename, typename, index_t>
Chao Liu's avatar
Chao Liu committed
289
struct sequence_reverse_inclusive_scan;
Chao Liu's avatar
Chao Liu committed
290

Chao Liu's avatar
Chao Liu committed
291
template <index_t I, index_t... Is, typename Reduce, index_t Init>
Chao Liu's avatar
Chao Liu committed
292
struct sequence_reverse_inclusive_scan<Sequence<I, Is...>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
293
{
Chao Liu's avatar
Chao Liu committed
294
    using old_scan = typename sequence_reverse_inclusive_scan<Sequence<Is...>, Reduce, Init>::type;
Chao Liu's avatar
Chao Liu committed
295
296
297

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

Chao Liu's avatar
Chao Liu committed
298
    using type = typename sequence_merge<Sequence<new_reduce>, old_scan>::type;
Chao Liu's avatar
Chao Liu committed
299
300
};

Chao Liu's avatar
Chao Liu committed
301
template <index_t I, typename Reduce, index_t Init>
Chao Liu's avatar
Chao Liu committed
302
struct sequence_reverse_inclusive_scan<Sequence<I>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
303
{
Chao Liu's avatar
Chao Liu committed
304
    using type = Sequence<Reduce{}(I, Init)>;
Chao Liu's avatar
Chao Liu committed
305
306
};

Chao Liu's avatar
Chao Liu committed
307
template <typename Reduce, index_t Init>
Chao Liu's avatar
Chao Liu committed
308
struct sequence_reverse_inclusive_scan<Sequence<>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
309
{
Chao Liu's avatar
Chao Liu committed
310
    using type = Sequence<>;
Chao Liu's avatar
Chao Liu committed
311
312
};

Chao Liu's avatar
Chao Liu committed
313
// split sequence
Chao Liu's avatar
Chao Liu committed
314
template <typename Seq, index_t I>
Chao Liu's avatar
Chao Liu committed
315
316
struct sequence_split
{
Chao Liu's avatar
Chao Liu committed
317
    static constexpr index_t NSize = Seq{}.Size();
Chao Liu's avatar
Chao Liu committed
318

Chao Liu's avatar
Chao Liu committed
319
320
    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
321

Chao Liu's avatar
Chao Liu committed
322
323
    using left_type  = decltype(Seq::Extract(range0{}));
    using right_type = decltype(Seq::Extract(range1{}));
Chao Liu's avatar
Chao Liu committed
324
325
};

Chao Liu's avatar
Chao Liu committed
326
// reverse sequence
Chao Liu's avatar
Chao Liu committed
327
template <typename Seq>
Chao Liu's avatar
Chao Liu committed
328
329
struct sequence_reverse
{
Chao Liu's avatar
Chao Liu committed
330
    static constexpr index_t NSize = Seq{}.Size();
Chao Liu's avatar
Chao Liu committed
331
332

    using seq_split = sequence_split<Seq, NSize / 2>;
Chao Liu's avatar
Chao Liu committed
333
    using type      = typename sequence_merge<
Chao Liu's avatar
Chao Liu committed
334
335
        typename sequence_reverse<typename seq_split::right_type>::type,
        typename sequence_reverse<typename seq_split::left_type>::type>::type;
Chao Liu's avatar
Chao Liu committed
336
337
338
339
340
};

template <index_t I>
struct sequence_reverse<Sequence<I>>
{
Chao Liu's avatar
Chao Liu committed
341
    using type = Sequence<I>;
Chao Liu's avatar
Chao Liu committed
342
343
344
345
346
};

template <index_t I0, index_t I1>
struct sequence_reverse<Sequence<I0, I1>>
{
Chao Liu's avatar
Chao Liu committed
347
    using type = Sequence<I1, I0>;
Chao Liu's avatar
Chao Liu committed
348
};
Chao Liu's avatar
Chao Liu committed
349

Chao Liu's avatar
Chao Liu committed
350
#if 1
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
template <typename Reduce, typename Seq, typename... Seqs>
struct sequence_reduce
{
    using type = typename sequence_reduce<Reduce,
                                          Seq,
                                          typename sequence_reduce<Reduce, Seqs...>::type>::type;
};

template <typename Reduce, index_t... Xs, index_t... Ys>
struct sequence_reduce<Reduce, Sequence<Xs...>, Sequence<Ys...>>
{
    using type = Sequence<Reduce{}(Xs, Ys)...>;
};

template <typename Reduce, typename Seq>
struct sequence_reduce<Reduce, Seq>
{
    using type = Seq;
};
#endif

Chao Liu's avatar
Chao Liu committed
372
373
template <typename Values, typename Ids, typename Compare>
struct sequence_sort_impl
Chao Liu's avatar
Chao Liu committed
374
{
Chao Liu's avatar
Chao Liu committed
375
376
377
378
379
380
381
    template <typename LeftValues,
              typename LeftIds,
              typename RightValues,
              typename RightIds,
              typename MergedValues,
              typename MergedIds,
              typename Comp>
Chao Liu's avatar
Chao Liu committed
382
383
    struct sorted_sequence_merge_impl
    {
Chao Liu's avatar
Chao Liu committed
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
        static constexpr bool choose_left = LeftValues::Front() < RightValues::Front();

        static constexpr index_t chosen_value =
            choose_left ? LeftValues::Front() : RightValues::Front();
        static constexpr index_t chosen_id = choose_left ? LeftIds::Front() : RightIds::Front();

        using new_merged_values = decltype(MergedValues::PushBack(Number<chosen_value>{}));
        using new_merged_ids    = decltype(MergedIds::PushBack(Number<chosen_id>{}));

        using new_left_values =
            typename conditional<choose_left, decltype(LeftValues::PopFront()), LeftValues>::type;
        using new_left_ids =
            typename conditional<choose_left, decltype(LeftIds::PopFront()), LeftIds>::type;

        using new_right_values =
            typename conditional<choose_left, RightValues, decltype(RightValues::PopFront())>::type;
        using new_right_ids =
            typename conditional<choose_left, RightIds, decltype(RightIds::PopFront())>::type;

        using merge = sorted_sequence_merge_impl<new_left_values,
                                                 new_left_ids,
                                                 new_right_values,
                                                 new_right_ids,
                                                 new_merged_values,
                                                 new_merged_ids,
                                                 Comp>;
        // this is output
        using merged_values = typename merge::merged_values;
        using merged_ids    = typename merge::merged_ids;
Chao Liu's avatar
Chao Liu committed
413
414
    };

Chao Liu's avatar
Chao Liu committed
415
416
417
418
419
420
421
422
423
424
425
426
    template <typename LeftValues,
              typename LeftIds,
              typename MergedValues,
              typename MergedIds,
              typename Comp>
    struct sorted_sequence_merge_impl<LeftValues,
                                      LeftIds,
                                      Sequence<>,
                                      Sequence<>,
                                      MergedValues,
                                      MergedIds,
                                      Comp>
Chao Liu's avatar
Chao Liu committed
427
    {
Chao Liu's avatar
Chao Liu committed
428
429
        using merged_values = typename sequence_merge<MergedValues, LeftValues>::type;
        using merged_ids    = typename sequence_merge<MergedIds, LeftIds>::type;
Chao Liu's avatar
Chao Liu committed
430
431
    };

Chao Liu's avatar
Chao Liu committed
432
433
434
435
436
437
438
439
440
441
442
443
    template <typename RightValues,
              typename RightIds,
              typename MergedValues,
              typename MergedIds,
              typename Comp>
    struct sorted_sequence_merge_impl<Sequence<>,
                                      Sequence<>,
                                      RightValues,
                                      RightIds,
                                      MergedValues,
                                      MergedIds,
                                      Comp>
Chao Liu's avatar
Chao Liu committed
444
    {
Chao Liu's avatar
Chao Liu committed
445
446
        using merged_values = typename sequence_merge<MergedValues, RightValues>::type;
        using merged_ids    = typename sequence_merge<MergedIds, RightIds>::type;
Chao Liu's avatar
Chao Liu committed
447
448
    };

Chao Liu's avatar
Chao Liu committed
449
450
451
452
453
    template <typename LeftValues,
              typename LeftIds,
              typename RightValues,
              typename RightIds,
              typename Comp>
Chao Liu's avatar
Chao Liu committed
454
455
    struct sorted_sequence_merge
    {
Chao Liu's avatar
Chao Liu committed
456
457
458
459
460
461
462
463
464
465
        using merge = sorted_sequence_merge_impl<LeftValues,
                                                 LeftIds,
                                                 RightValues,
                                                 RightIds,
                                                 Sequence<>,
                                                 Sequence<>,
                                                 Comp>;

        using merged_values = typename merge::merged_values;
        using merged_ids    = typename merge::merged_ids;
Chao Liu's avatar
Chao Liu committed
466
467
    };

Chao Liu's avatar
Chao Liu committed
468
469
470
471
    static constexpr index_t nsize = Values::Size();

    using split_unsorted_values = sequence_split<Values, nsize / 2>;
    using split_unsorted_ids    = sequence_split<Ids, nsize / 2>;
Chao Liu's avatar
Chao Liu committed
472

Chao Liu's avatar
Chao Liu committed
473
474
475
476
477
    using left_unsorted_values = typename split_unsorted_values::left_type;
    using left_unsorted_ids    = typename split_unsorted_ids::left_type;
    using left_sort          = sequence_sort_impl<left_unsorted_values, left_unsorted_ids, Compare>;
    using left_sorted_values = typename left_sort::sorted_values;
    using left_sorted_ids    = typename left_sort::sorted_ids;
Chao Liu's avatar
Chao Liu committed
478

Chao Liu's avatar
Chao Liu committed
479
480
481
482
483
484
485
486
487
488
489
490
491
492
    using right_unsorted_values = typename split_unsorted_values::right_type;
    using right_unsorted_ids    = typename split_unsorted_ids::right_type;
    using right_sort = sequence_sort_impl<right_unsorted_values, right_unsorted_ids, Compare>;
    using right_sorted_values = typename right_sort::sorted_values;
    using right_sorted_ids    = typename right_sort::sorted_ids;

    using merged_sorted = sorted_sequence_merge<left_sorted_values,
                                                left_sorted_ids,
                                                right_sorted_values,
                                                right_sorted_ids,
                                                Compare>;

    using sorted_values = typename merged_sorted::merged_values;
    using sorted_ids    = typename merged_sorted::merged_ids;
Chao Liu's avatar
Chao Liu committed
493
494
};

Chao Liu's avatar
Chao Liu committed
495
496
template <index_t ValueX, index_t ValueY, index_t IdX, index_t IdY, typename Compare>
struct sequence_sort_impl<Sequence<ValueX, ValueY>, Sequence<IdX, IdY>, Compare>
Chao Liu's avatar
Chao Liu committed
497
{
Chao Liu's avatar
Chao Liu committed
498
499
500
501
502
503
    static constexpr bool choose_x = Compare{}(ValueX, ValueY);

    using sorted_values =
        typename conditional<choose_x, Sequence<ValueX, ValueY>, Sequence<ValueY, ValueX>>::type;
    using sorted_ids = typename conditional<choose_x, Sequence<IdX, IdY>, Sequence<IdY, IdX>>::type;
};
Chao Liu's avatar
Chao Liu committed
504

Chao Liu's avatar
Chao Liu committed
505
506
507
508
509
template <index_t Value, index_t Id, typename Compare>
struct sequence_sort_impl<Sequence<Value>, Sequence<Id>, Compare>
{
    using sorted_values = Sequence<Value>;
    using sorted_ids    = Sequence<Id>;
Chao Liu's avatar
Chao Liu committed
510
511
};

512
513
514
515
516
517
518
template <typename Compare>
struct sequence_sort_impl<Sequence<>, Sequence<>, Compare>
{
    using sorted_values = Sequence<>;
    using sorted_ids    = Sequence<>;
};

Chao Liu's avatar
Chao Liu committed
519
520
template <typename Values, typename Compare>
struct sequence_sort
Chao Liu's avatar
Chao Liu committed
521
{
Chao Liu's avatar
Chao Liu committed
522
523
524
525
526
527
    using unsorted_ids = typename arithmetic_sequence_gen<0, Values::Size(), 1>::type;
    using sort         = sequence_sort_impl<Values, unsorted_ids, Compare>;

    // this is output
    using type                = typename sort::sorted_values;
    using sorted2unsorted_map = typename sort::sorted_ids;
Chao Liu's avatar
Chao Liu committed
528
529
};

Chao Liu's avatar
Chao Liu committed
530
template <typename Values, typename Less, typename Equal>
Chao Liu's avatar
Chao Liu committed
531
532
struct sequence_unique_sort
{
Chao Liu's avatar
Chao Liu committed
533
534
535
536
537
    template <typename RemainValues,
              typename RemainIds,
              typename UniquifiedValues,
              typename UniquifiedIds,
              typename Eq>
Chao Liu's avatar
Chao Liu committed
538
539
    struct sorted_sequence_uniquify_impl
    {
Chao Liu's avatar
Chao Liu committed
540
541
542
543
544
545
546
547
548
549
550
551
        static constexpr index_t current_value = RemainValues::Front();
        static constexpr index_t current_id    = RemainIds::Front();

        static constexpr bool is_unique_value = (current_value != UniquifiedValues::Back());

        using new_remain_values = decltype(RemainValues::PopFront());
        using new_remain_ids    = decltype(RemainIds::PopFront());

        using new_uniquified_values =
            typename conditional<is_unique_value,
                                 decltype(UniquifiedValues::PushBack(Number<current_value>{})),
                                 UniquifiedValues>::type;
Chao Liu's avatar
Chao Liu committed
552

Chao Liu's avatar
Chao Liu committed
553
554
555
556
557
558
559
560
561
562
563
564
565
566
        using new_uniquified_ids =
            typename conditional<is_unique_value,
                                 decltype(UniquifiedIds::PushBack(Number<current_id>{})),
                                 UniquifiedIds>::type;

        using uniquify = sorted_sequence_uniquify_impl<new_remain_values,
                                                       new_remain_ids,
                                                       new_uniquified_values,
                                                       new_uniquified_ids,
                                                       Eq>;

        // this is output
        using uniquified_values = typename uniquify::uniquified_values;
        using uniquified_ids    = typename uniquify::uniquified_ids;
Chao Liu's avatar
Chao Liu committed
567
568
    };

Chao Liu's avatar
Chao Liu committed
569
570
571
572
573
574
    template <typename UniquifiedValues, typename UniquifiedIds, typename Eq>
    struct sorted_sequence_uniquify_impl<Sequence<>,
                                         Sequence<>,
                                         UniquifiedValues,
                                         UniquifiedIds,
                                         Eq>
Chao Liu's avatar
Chao Liu committed
575
    {
Chao Liu's avatar
Chao Liu committed
576
577
        using uniquified_values = UniquifiedValues;
        using uniquified_ids    = UniquifiedIds;
Chao Liu's avatar
Chao Liu committed
578
579
    };

Chao Liu's avatar
Chao Liu committed
580
    template <typename SortedValues, typename SortedIds, typename Eq>
Chao Liu's avatar
Chao Liu committed
581
582
    struct sorted_sequence_uniquify
    {
Chao Liu's avatar
Chao Liu committed
583
584
585
586
587
588
589
590
        using uniquify = sorted_sequence_uniquify_impl<decltype(SortedValues::PopFront()),
                                                       decltype(SortedIds::PopFront()),
                                                       Sequence<SortedValues::Front()>,
                                                       Sequence<SortedIds::Front()>,
                                                       Eq>;

        using uniquified_values = typename uniquify::uniquified_values;
        using uniquified_ids    = typename uniquify::uniquified_ids;
Chao Liu's avatar
Chao Liu committed
591
592
    };

Chao Liu's avatar
Chao Liu committed
593
594
595
    using sort          = sequence_sort<Values, Less>;
    using sorted_values = typename sort::type;
    using sorted_ids    = typename sort::sorted2unsorted_map;
Chao Liu's avatar
Chao Liu committed
596

Chao Liu's avatar
Chao Liu committed
597
598
599
600
601
    using uniquify = sorted_sequence_uniquify<sorted_values, sorted_ids, Equal>;

    // this is output
    using type                = typename uniquify::uniquified_values;
    using sorted2unsorted_map = typename uniquify::uniquified_ids;
Chao Liu's avatar
Chao Liu committed
602
603
};

Chao Liu's avatar
Chao Liu committed
604
template <typename SeqMap>
Chao Liu's avatar
Chao Liu committed
605
606
struct is_valid_sequence_map : is_same<typename arithmetic_sequence_gen<0, SeqMap::Size(), 1>::type,
                                       typename sequence_sort<SeqMap, math::less<index_t>>::type>
Chao Liu's avatar
Chao Liu committed
607
608
{
};
609

Chao Liu's avatar
Chao Liu committed
610
611
template <typename SeqMap>
struct sequence_map_inverse
Chao Liu's avatar
Chao Liu committed
612
{
Chao Liu's avatar
Chao Liu committed
613
614
615
616
617
    template <typename X2Y, typename WorkingY2X, index_t XBegin, index_t XRemain>
    struct sequence_map_inverse_impl
    {
        static constexpr auto new_y2x =
            WorkingY2X::Modify(X2Y::At(Number<XBegin>{}), Number<XBegin>{});
Chao Liu's avatar
Chao Liu committed
618

Chao Liu's avatar
Chao Liu committed
619
620
621
622
        using type =
            typename sequence_map_inverse_impl<X2Y, decltype(new_y2x), XBegin + 1, XRemain - 1>::
                type;
    };
Chao Liu's avatar
Chao Liu committed
623

Chao Liu's avatar
Chao Liu committed
624
625
626
627
628
    template <typename X2Y, typename WorkingY2X, index_t XBegin>
    struct sequence_map_inverse_impl<X2Y, WorkingY2X, XBegin, 0>
    {
        using type = WorkingY2X;
    };
Chao Liu's avatar
Chao Liu committed
629
630

    using type =
Chao Liu's avatar
Chao Liu committed
631
632
        typename sequence_map_inverse_impl<SeqMap,
                                           typename uniform_sequence_gen<SeqMap::Size(), 0>::type,
Chao Liu's avatar
Chao Liu committed
633
                                           0,
Chao Liu's avatar
Chao Liu committed
634
                                           SeqMap::Size()>::type;
Chao Liu's avatar
Chao Liu committed
635
636
};

637
638
639
640
641
642
template <index_t... Xs, index_t... Ys>
__host__ __device__ constexpr bool operator==(Sequence<Xs...>, Sequence<Ys...>)
{
    return ((Xs == Ys) && ...);
}

Chao Liu's avatar
Chao Liu committed
643
644
645
646
647
648
template <index_t... Xs, index_t... Ys>
__host__ __device__ constexpr bool operator!=(Sequence<Xs...> x, Sequence<Ys...> y)
{
    return !(x == y);
}

Chao Liu's avatar
Chao Liu committed
649
template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
650
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
651
652
653
654
655
656
657
{
    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
658
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
659
660
661
662
663
664
665
{
    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
666
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
667
668
669
670
671
672
673
{
    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
674
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
675
676
677
678
679
680
681
{
    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
682
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
683
684
685
686
687
688
689
{
    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
690
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
691
{
Chao Liu's avatar
Chao Liu committed
692
    return Sequence<(Xs + Y)...>{};
Chao Liu's avatar
Chao Liu committed
693
694
695
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
696
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
697
{
Chao Liu's avatar
Chao Liu committed
698
    return Sequence<(Xs - Y)...>{};
Chao Liu's avatar
Chao Liu committed
699
700
701
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
702
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
703
{
Chao Liu's avatar
Chao Liu committed
704
    return Sequence<(Xs * Y)...>{};
Chao Liu's avatar
Chao Liu committed
705
706
707
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
708
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
709
{
Chao Liu's avatar
Chao Liu committed
710
    return Sequence<(Xs / Y)...>{};
Chao Liu's avatar
Chao Liu committed
711
712
713
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
714
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
715
{
Chao Liu's avatar
Chao Liu committed
716
    return Sequence<(Xs % Y)...>{};
Chao Liu's avatar
Chao Liu committed
717
718
}

Chao Liu's avatar
Chao Liu committed
719
720
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator+(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
721
{
Chao Liu's avatar
Chao Liu committed
722
    return Sequence<(Y + Xs)...>{};
Chao Liu's avatar
Chao Liu committed
723
724
}

Chao Liu's avatar
Chao Liu committed
725
726
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
727
{
Chao Liu's avatar
Chao Liu committed
728
    return Sequence<(Y - Xs)...>{};
Chao Liu's avatar
Chao Liu committed
729
730
}

Chao Liu's avatar
Chao Liu committed
731
732
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator*(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
733
{
Chao Liu's avatar
Chao Liu committed
734
    return Sequence<(Y * Xs)...>{};
Chao Liu's avatar
Chao Liu committed
735
736
}

Chao Liu's avatar
Chao Liu committed
737
738
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator/(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
739
{
Chao Liu's avatar
Chao Liu committed
740
    return Sequence<(Y / Xs)...>{};
Chao Liu's avatar
Chao Liu committed
741
742
}

Chao Liu's avatar
Chao Liu committed
743
744
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator%(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
745
{
Chao Liu's avatar
Chao Liu committed
746
    return Sequence<(Y % Xs)...>{};
Chao Liu's avatar
Chao Liu committed
747
748
}

749
750
751
752
753
754
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
755
template <typename Seq>
Chao Liu's avatar
Chao Liu committed
756
__host__ __device__ constexpr auto sequence_pop_back(Seq)
757
{
Chao Liu's avatar
Chao Liu committed
758
    static_assert(Seq::Size() > 0, "wrong! cannot pop an empty Sequence!");
759
    return sequence_pop_front(Seq::Reverse()).Reverse();
760
}
761

Chao Liu's avatar
Chao Liu committed
762
763
764
765
766
767
template <typename... Seqs>
__host__ __device__ constexpr auto merge_sequences(Seqs...)
{
    return typename sequence_merge<Seqs...>::type{};
}

768
769
770
771
772
773
template <typename 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
774
template <typename F, index_t... Xs, index_t... Ys>
775
__host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>)
776
{
777
    static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize, "Dim not the same");
778
779
780
781

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

Chao Liu's avatar
Chao Liu committed
782
template <typename F, index_t... Xs, index_t... Ys, index_t... Zs>
783
784
785
786
787
788
789
790
791
792
__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
793
template <typename Seq, typename Reduce, index_t Init>
Chao Liu's avatar
Chao Liu committed
794
__host__ __device__ constexpr auto reverse_inclusive_scan_sequence(Seq, Reduce, Number<Init>)
795
{
Chao Liu's avatar
Chao Liu committed
796
    return typename sequence_reverse_inclusive_scan<Seq, Reduce, Init>::type{};
797
798
}

Chao Liu's avatar
Chao Liu committed
799
800
801
802
803
804
805
template <typename Seq, typename Reduce, index_t Init>
__host__ __device__ constexpr auto reverse_exclusive_scan_sequence(Seq, Reduce, Number<Init>)
{
    return reverse_inclusive_scan_sequence(Seq::PopFront(), Reduce{}, Number<Init>{})
        .PushBack(Number<Init>{});
}

Chao Liu's avatar
Chao Liu committed
806
template <typename Seq, typename Reduce, index_t Init>
Chao Liu's avatar
Chao Liu committed
807
__host__ __device__ constexpr auto inclusive_scan_sequence(Seq, Reduce, Number<Init>)
808
{
Chao Liu's avatar
Chao Liu committed
809
    return reverse_inclusive_scan_sequence(Seq{}.Reverse(), Reduce{}, Number<Init>{}).Reverse();
810
}
811

carlushuang's avatar
carlushuang committed
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
835
836
837
838
839
840
841
842
843
844
845
846
847
848
849
850
851
// e.g. Seq<2, 3, 4> --> Seq<0, 2, 5>, Init=0, Reduce=Add
//      ResultSeq  TargetSeq  Reduce
template <typename, typename, typename>
struct sequence_exclusive_scan;

template <index_t... Xs, index_t Y, index_t... Ys, typename Reduce>
struct sequence_exclusive_scan<Sequence<Xs...>, Sequence<Y, Ys...>, Reduce>
{
    using old_scan = typename sequence_merge<Sequence<Xs...>,
                                             Sequence<Reduce{}(Y, Sequence<Xs...>{}.Back())>>::type;
    using type     = typename sequence_exclusive_scan<old_scan, Sequence<Ys...>, Reduce>::type;
};

template <index_t... Xs, index_t Y, typename Reduce>
struct sequence_exclusive_scan<Sequence<Xs...>, Sequence<Y>, Reduce>
{
    using type = Sequence<Xs...>;
};

template <index_t... Xs, typename Reduce>
struct sequence_exclusive_scan<Sequence<Xs...>, Sequence<>, Reduce>
{
    using type = Sequence<Xs...>;
};

template <typename Seq, typename Reduce, index_t Init>
constexpr auto exclusive_scan_sequence(Seq, Reduce, Number<Init>)
{
    // TODO: c++20 and later can pass in Reduce with a lambda expression
    return typename sequence_exclusive_scan<Sequence<Init>, Seq, Reduce>::type{};
}

template <typename Seq>
constexpr auto prefix_sum_sequence(Seq)
{
    return typename sequence_exclusive_scan<Sequence<0>,
                                            typename sequence_merge<Seq, Sequence<0>>::type,
                                            math::plus<index_t>>::type{};
}

Chao Liu's avatar
Chao Liu committed
852
template <typename Seq, index_t... Is>
853
__host__ __device__ constexpr auto pick_sequence_elements_by_ids(Seq, Sequence<Is...> /* ids */)
Chao Liu's avatar
Chao Liu committed
854
855
856
857
{
    return Sequence<Seq::At(Number<Is>{})...>{};
}

Chao Liu's avatar
Chao Liu committed
858
859
860
861
862
863
864
865
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
#if 1
namespace detail {
template <typename WorkSeq, typename RemainSeq, typename RemainMask>
struct pick_sequence_elements_by_mask_impl
{
    using new_work_seq = typename conditional<RemainMask::Front(),
                                              decltype(WorkSeq::PushBack(RemainSeq::Front())),
                                              WorkSeq>::type;

    using type =
        typename pick_sequence_elements_by_mask_impl<new_work_seq,
                                                     decltype(RemainSeq::PopFront()),
                                                     decltype(RemainMask::PopFront())>::type;
};

template <typename WorkSeq>
struct pick_sequence_elements_by_mask_impl<WorkSeq, Sequence<>, Sequence<>>
{
    using type = WorkSeq;
};

} // namespace detail

881
882
883
template <typename Seq, typename Mask>
__host__ __device__ constexpr auto pick_sequence_elements_by_mask(Seq, Mask)
{
Chao Liu's avatar
Chao Liu committed
884
885
886
    static_assert(Seq::Size() == Mask::Size(), "wrong!");

    return typename detail::pick_sequence_elements_by_mask_impl<Sequence<>, Seq, Mask>::type{};
887
888
}

Chao Liu's avatar
Chao Liu committed
889
890
891
namespace detail {
template <typename WorkSeq, typename RemainValues, typename RemainIds>
struct modify_sequence_elements_by_ids_impl
Chao Liu's avatar
Chao Liu committed
892
{
Chao Liu's avatar
Chao Liu committed
893
    using new_work_seq = decltype(WorkSeq::Modify(RemainIds::Front(), RemainValues::Front()));
Chao Liu's avatar
Chao Liu committed
894

Chao Liu's avatar
Chao Liu committed
895
896
897
898
899
    using type =
        typename modify_sequence_elements_by_ids_impl<new_work_seq,
                                                      decltype(RemainValues::PopFront()),
                                                      decltype(RemainIds::PopFront())>::type;
};
Chao Liu's avatar
Chao Liu committed
900

Chao Liu's avatar
Chao Liu committed
901
902
903
904
template <typename WorkSeq>
struct modify_sequence_elements_by_ids_impl<WorkSeq, Sequence<>, Sequence<>>
{
    using type = WorkSeq;
Chao Liu's avatar
Chao Liu committed
905
};
Chao Liu's avatar
Chao Liu committed
906
907
908
909
910
911
912
913
914
915
} // namespace detail

template <typename Seq, typename Values, typename Ids>
__host__ __device__ constexpr auto modify_sequence_elements_by_ids(Seq, Values, Ids)
{
    static_assert(Values::Size() == Ids::Size() && Seq::Size() >= Values::Size(), "wrong!");

    return typename detail::modify_sequence_elements_by_ids_impl<Seq, Values, Ids>::type{};
}
#endif
Chao Liu's avatar
Chao Liu committed
916

Chao Liu's avatar
Chao Liu committed
917
template <typename Seq, typename Reduce, index_t Init>
Chao Liu's avatar
Chao Liu committed
918
__host__ __device__ constexpr index_t
Chao Liu's avatar
Chao Liu committed
919
reduce_on_sequence(Seq, Reduce f, Number<Init> /*initial_value*/)
Chao Liu's avatar
Chao Liu committed
920
921
922
{
    index_t result = Init;

Chao Liu's avatar
Chao Liu committed
923
924
925
926
    for(index_t i = 0; i < Seq::Size(); ++i)
    {
        result = f(result, Seq::At(i));
    }
Chao Liu's avatar
Chao Liu committed
927
928
929
930

    return result;
}

Chao Liu's avatar
Chao Liu committed
931
932
// TODO: a generic any_of for any container
template <typename Seq, typename F>
Chao Liu's avatar
Chao Liu committed
933
__host__ __device__ constexpr bool sequence_any_of(Seq, F f)
Chao Liu's avatar
Chao Liu committed
934
935
936
937
938
939
940
941
942
943
944
945
946
{
    bool flag = false;

    for(index_t i = 0; i < Seq::Size(); ++i)
    {
        flag = flag || f(Seq::At(i));
    }

    return flag;
}

// TODO: a generic all_of for any container
template <typename Seq, typename F>
Chao Liu's avatar
Chao Liu committed
947
__host__ __device__ constexpr bool sequence_all_of(Seq, F f)
Chao Liu's avatar
Chao Liu committed
948
949
950
951
952
953
954
955
956
957
958
{
    bool flag = true;

    for(index_t i = 0; i < Seq::Size(); ++i)
    {
        flag = flag && f(Seq::At(i));
    }

    return flag;
}

Chao Liu's avatar
Chao Liu committed
959
960
template <typename... Seqs>
using sequence_merge_t = typename sequence_merge<Seqs...>::type;
961
962
963
964

template <index_t NSize, index_t I>
using uniform_sequence_gen_t = typename uniform_sequence_gen<NSize, I>::type;

965
} // namespace ck