"vscode:/vscode.git/clone" did not exist on "21c918162e81211d2b7ec3a555accef1235957eb"
Sequence.hpp 15 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
template <class>
Chao Liu's avatar
Chao Liu committed
10
11
struct is_valid_sequence_map;

Chao Liu's avatar
Chao Liu committed
12
13
14
template <class>
struct sequence_map_inverse;

Chao Liu's avatar
Chao Liu committed
15
template <index_t... Is>
16
17
struct Sequence
{
Chao Liu's avatar
Chao Liu committed
18
19
    using Type      = Sequence;
    using data_type = index_t;
20

21
    static constexpr index_t mSize = sizeof...(Is);
22

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

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

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

Chao Liu's avatar
Chao Liu committed
40
41
    __host__ __device__ static constexpr auto Get(index_t I) { return GetImpl(I); }

Chao Liu's avatar
Chao Liu committed
42
43
    template <index_t I>
    __host__ __device__ constexpr auto operator[](Number<I>) const
Chao Liu's avatar
Chao Liu committed
44
    {
Chao Liu's avatar
Chao Liu committed
45
        return Get(Number<I>{});
46
47
    }

Chao Liu's avatar
Chao Liu committed
48
49
50
    // 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); }

51
    template <index_t... IRs>
52
    __host__ __device__ static constexpr auto ReorderGivenNew2Old(Sequence<IRs...> /*new2old*/)
53
    {
Chao Liu's avatar
Chao Liu committed
54
        static_assert(sizeof...(Is) == sizeof...(IRs),
Chao Liu's avatar
Chao Liu committed
55
                      "wrong! reorder map should have the same size as Sequence to be rerodered");
Chao Liu's avatar
Chao Liu committed
56

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

        return Sequence<Type::Get(Number<IRs>{})...>{};
60
61
    }

Chao Liu's avatar
Chao Liu committed
62
63
64
65
66
67
68
69
70
71
72
73
    // MapOld2New is Sequence<...>
    template <class MapOld2New>
    __host__ __device__ static constexpr auto ReorderGivenOld2New(MapOld2New)
    {
        static_assert(MapOld2New::GetSize() == GetSize(),
                      "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{});
    }

74
    __host__ __device__ static constexpr auto Reverse();
Chao Liu's avatar
Chao Liu committed
75

Chao Liu's avatar
Chao Liu committed
76
    __host__ __device__ static constexpr auto Front()
77
    {
Chao Liu's avatar
Chao Liu committed
78
79
        static_assert(mSize > 0, "wrong!");
        return Get(Number<0>{});
80
    }
81

Chao Liu's avatar
Chao Liu committed
82
    __host__ __device__ static constexpr auto Back()
83
    {
Chao Liu's avatar
Chao Liu committed
84
85
        static_assert(mSize > 0, "wrong!");
        return Get(Number<mSize - 1>{});
86
    }
87

Chao Liu's avatar
Chao Liu committed
88
89
90
91
92
93
    __host__ __device__ static constexpr auto PopFront();

    __host__ __device__ static constexpr auto PopBack();

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

Chao Liu's avatar
Chao Liu committed
98
99
    template <index_t... Xs>
    __host__ __device__ static constexpr auto PushFront(Number<Xs>...)
100
    {
Chao Liu's avatar
Chao Liu committed
101
        return Sequence<Xs..., Is...>{};
102
103
    }

Chao Liu's avatar
Chao Liu committed
104
105
106
107
108
    template <index_t... Xs>
    __host__ __device__ static constexpr auto PushBack(Sequence<Xs...>)
    {
        return Sequence<Is..., Xs...>{};
    }
109

Chao Liu's avatar
Chao Liu committed
110
    template <index_t... Xs>
Chao Liu's avatar
Chao Liu committed
111
    __host__ __device__ static constexpr auto PushBack(Number<Xs>...)
112
    {
Chao Liu's avatar
Chao Liu committed
113
114
        return Sequence<Is..., Xs...>{};
    }
Chao Liu's avatar
Chao Liu committed
115

Chao Liu's avatar
Chao Liu committed
116
    template <index_t... Ns>
117
    __host__ __device__ static constexpr auto Extract(Number<Ns>...)
Chao Liu's avatar
Chao Liu committed
118
    {
Chao Liu's avatar
Chao Liu committed
119
        return Sequence<Type::Get(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
120
    }
Chao Liu's avatar
Chao Liu committed
121

Chao Liu's avatar
Chao Liu committed
122
    template <index_t... Ns>
123
    __host__ __device__ static constexpr auto Extract(Sequence<Ns...>)
Chao Liu's avatar
Chao Liu committed
124
    {
Chao Liu's avatar
Chao Liu committed
125
        return Sequence<Type::Get(Number<Ns>{})...>{};
Chao Liu's avatar
Chao Liu committed
126
    }
127
128
129

    template <index_t I, index_t X>
    __host__ __device__ static constexpr auto Modify(Number<I>, Number<X>);
Chao Liu's avatar
Chao Liu committed
130
131
132
133
134
135

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

Chao Liu's avatar
Chao Liu committed
138
// merge sequence
Chao Liu's avatar
Chao Liu committed
139
140
template <class, class>
struct sequence_merge;
Chao Liu's avatar
Chao Liu committed
141

Chao Liu's avatar
Chao Liu committed
142
143
144
template <index_t... Xs, index_t... Ys>
struct sequence_merge<Sequence<Xs...>, Sequence<Ys...>>
{
Chao Liu's avatar
Chao Liu committed
145
    using type = Sequence<Xs..., Ys...>;
Chao Liu's avatar
Chao Liu committed
146
};
Chao Liu's avatar
Chao Liu committed
147

Chao Liu's avatar
Chao Liu committed
148
149
150
// generate sequence
template <index_t IBegin, index_t NRemain, class F>
struct sequence_gen_impl
Chao Liu's avatar
Chao Liu committed
151
{
Chao Liu's avatar
Chao Liu committed
152
153
154
    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
155

Chao Liu's avatar
Chao Liu committed
156
157
158
    using type =
        typename sequence_merge<typename sequence_gen_impl<IBegin, NRemainLeft, F>::type,
                                typename sequence_gen_impl<IMiddle, NRemainRight, F>::type>::type;
Chao Liu's avatar
Chao Liu committed
159
160
};

Chao Liu's avatar
Chao Liu committed
161
162
template <index_t I, class F>
struct sequence_gen_impl<I, 1, F>
Chao Liu's avatar
Chao Liu committed
163
{
Chao Liu's avatar
Chao Liu committed
164
165
    static constexpr index_t Is = F{}(Number<I>{});
    using type                  = Sequence<Is>;
Chao Liu's avatar
Chao Liu committed
166
};
Chao Liu's avatar
Chao Liu committed
167

Chao Liu's avatar
Chao Liu committed
168
169
template <index_t I, class F>
struct sequence_gen_impl<I, 0, F>
Chao Liu's avatar
Chao Liu committed
170
{
Chao Liu's avatar
Chao Liu committed
171
    using type = Sequence<>;
Chao Liu's avatar
Chao Liu committed
172
173
};

Chao Liu's avatar
Chao Liu committed
174
175
176
177
178
179
180
template <index_t NSize, class F>
struct sequence_gen
{
    using type = typename sequence_gen_impl<0, NSize, F>::type;
};

// arithmetic sequence
Chao Liu's avatar
Chao Liu committed
181
template <index_t IBegin, index_t IEnd, index_t Increment>
182
struct arithmetic_sequence_gen
Chao Liu's avatar
Chao Liu committed
183
{
Chao Liu's avatar
Chao Liu committed
184
185
186
187
188
189
190
191
192
    struct F
    {
        __host__ __device__ constexpr index_t operator()(index_t i) const
        {
            return i * Increment + IBegin;
        }
    };

    using type = typename sequence_gen<(IEnd - IBegin) / Increment, F>::type;
Chao Liu's avatar
Chao Liu committed
193
194
195
196
197
198
};

// uniform sequence
template <index_t NSize, index_t I>
struct uniform_sequence_gen
{
Chao Liu's avatar
Chao Liu committed
199
    struct F
Chao Liu's avatar
Chao Liu committed
200
201
202
203
    {
        __host__ __device__ constexpr index_t operator()(index_t) const { return I; }
    };

Chao Liu's avatar
Chao Liu committed
204
    using type = typename sequence_gen<NSize, F>::type;
Chao Liu's avatar
Chao Liu committed
205
206
207
};

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

Chao Liu's avatar
Chao Liu committed
211
212
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
213
{
Chao Liu's avatar
Chao Liu committed
214
    using old_scan = typename sequence_reverse_inclusive_scan<Sequence<Is...>, Reduce, Init>::type;
Chao Liu's avatar
Chao Liu committed
215
216
217

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

Chao Liu's avatar
Chao Liu committed
218
    using type = typename sequence_merge<Sequence<new_reduce>, old_scan>::type;
Chao Liu's avatar
Chao Liu committed
219
220
};

Chao Liu's avatar
Chao Liu committed
221
222
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
223
{
Chao Liu's avatar
Chao Liu committed
224
    using type = Sequence<Reduce{}(I, Init)>;
Chao Liu's avatar
Chao Liu committed
225
226
};

Chao Liu's avatar
Chao Liu committed
227
228
template <class Reduce, index_t Init>
struct sequence_reverse_inclusive_scan<Sequence<>, Reduce, Init>
Chao Liu's avatar
Chao Liu committed
229
{
Chao Liu's avatar
Chao Liu committed
230
    using type = Sequence<>;
Chao Liu's avatar
Chao Liu committed
231
232
};

Chao Liu's avatar
Chao Liu committed
233
// split sequence
Chao Liu's avatar
Chao Liu committed
234
235
236
237
238
template <class Seq, index_t I>
struct sequence_split
{
    static constexpr index_t NSize = Seq{}.GetSize();

Chao Liu's avatar
Chao Liu committed
239
240
    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
241

Chao Liu's avatar
Chao Liu committed
242
243
    using SeqType0 = decltype(Seq::Extract(range0{}));
    using SeqType1 = decltype(Seq::Extract(range1{}));
Chao Liu's avatar
Chao Liu committed
244
245
};

Chao Liu's avatar
Chao Liu committed
246
// reverse sequence
Chao Liu's avatar
Chao Liu committed
247
248
249
250
251
252
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
253
254
255
    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
256
257
258
259
260
};

template <index_t I>
struct sequence_reverse<Sequence<I>>
{
Chao Liu's avatar
Chao Liu committed
261
    using type = Sequence<I>;
Chao Liu's avatar
Chao Liu committed
262
263
264
265
266
};

template <index_t I0, index_t I1>
struct sequence_reverse<Sequence<I0, I1>>
{
Chao Liu's avatar
Chao Liu committed
267
    using type = Sequence<I1, I0>;
Chao Liu's avatar
Chao Liu committed
268
};
Chao Liu's avatar
Chao Liu committed
269

Chao Liu's avatar
Chao Liu committed
270
271
272
template <class Seq>
struct is_valid_sequence_map
{
Chao Liu's avatar
Chao Liu committed
273
    // not implemented yet, always return true
Chao Liu's avatar
Chao Liu committed
274
    static constexpr integral_constant<bool, true> value = integral_constant<bool, true>{};
Chao Liu's avatar
Chao Liu committed
275
276
277

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

Chao Liu's avatar
Chao Liu committed
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
template <class X2Y, class WorkingY2X, index_t XBegin, index_t XRemain>
struct sequence_map_inverse_impl
{
    private:
    static constexpr auto new_y2x = WorkingY2X::Modify(X2Y{}[XBegin], XBegin);

    public:
    using type =
        typename sequence_map_inverse_impl<X2Y, decltype(new_y2x), XBegin + 1, XRemain - 1>::type;
};

template <class X2Y, class WorkingY2X, index_t XBegin>
struct sequence_map_inverse_impl<X2Y, WorkingY2X, XBegin, 0>
{
    using type = WorkingY2X;
};

template <class X2Y>
struct sequence_map_inverse
{
    using type =
        typename sequence_map_inverse_impl<X2Y,
                                           typename uniform_sequence_gen<X2Y::GetSize(), 0>::type,
                                           0,
                                           X2Y::GetSize()>::type;
};

Chao Liu's avatar
Chao Liu committed
309
template <index_t... Xs, index_t... Ys>
Chao Liu's avatar
Chao Liu committed
310
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
311
312
313
314
315
316
317
{
    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
318
__host__ __device__ constexpr auto operator-(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
319
320
321
322
323
324
325
{
    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
326
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
327
328
329
330
331
332
333
{
    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
334
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
335
336
337
338
339
340
341
{
    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
342
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
Chao Liu's avatar
Chao Liu committed
343
344
345
346
347
348
349
{
    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
350
__host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
351
{
Chao Liu's avatar
Chao Liu committed
352
    return Sequence<(Xs + Y)...>{};
Chao Liu's avatar
Chao Liu committed
353
354
355
}

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

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
362
__host__ __device__ constexpr auto operator*(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
363
{
Chao Liu's avatar
Chao Liu committed
364
    return Sequence<(Xs * Y)...>{};
Chao Liu's avatar
Chao Liu committed
365
366
367
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
368
__host__ __device__ constexpr auto operator/(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
369
{
Chao Liu's avatar
Chao Liu committed
370
    return Sequence<(Xs / Y)...>{};
Chao Liu's avatar
Chao Liu committed
371
372
373
}

template <index_t... Xs, index_t Y>
Chao Liu's avatar
Chao Liu committed
374
__host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
375
{
Chao Liu's avatar
Chao Liu committed
376
    return Sequence<(Xs % Y)...>{};
Chao Liu's avatar
Chao Liu committed
377
378
}

Chao Liu's avatar
Chao Liu committed
379
380
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator+(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
381
{
Chao Liu's avatar
Chao Liu committed
382
    return Sequence<(Y + Xs)...>{};
Chao Liu's avatar
Chao Liu committed
383
384
}

Chao Liu's avatar
Chao Liu committed
385
386
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
387
{
Chao Liu's avatar
Chao Liu committed
388
389
390
    constexpr auto seq_x = Sequence<Xs...>{};

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

Chao Liu's avatar
Chao Liu committed
393
394
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator*(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
395
{
Chao Liu's avatar
Chao Liu committed
396
    return Sequence<(Y * Xs)...>{};
Chao Liu's avatar
Chao Liu committed
397
398
}

Chao Liu's avatar
Chao Liu committed
399
400
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator/(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
401
{
Chao Liu's avatar
Chao Liu committed
402
    return Sequence<(Y / Xs)...>{};
Chao Liu's avatar
Chao Liu committed
403
404
}

Chao Liu's avatar
Chao Liu committed
405
406
template <index_t Y, index_t... Xs>
__host__ __device__ constexpr auto operator%(Number<Y>, Sequence<Xs...>)
Chao Liu's avatar
Chao Liu committed
407
{
Chao Liu's avatar
Chao Liu committed
408
    return Sequence<(Y % Xs)...>{};
Chao Liu's avatar
Chao Liu committed
409
410
}

411
412
413
414
415
416
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
417
418
template <class Seq>
__host__ __device__ constexpr auto sequence_pop_back(Seq)
419
{
420
    static_assert(Seq{}.GetSize() > 0, "wrong! cannot pop an empty Sequence!");
Chao Liu's avatar
Chao Liu committed
421
    return sequence_pop_front(Seq{}.Reverse()).Reverse();
422
}
423

Chao Liu's avatar
Chao Liu committed
424
425
426
427
428
429
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
430
template <class F, index_t... Xs, index_t... Ys>
431
__host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>)
432
{
433
    static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize, "Dim not the same");
434
435
436
437

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

438
439
440
441
442
443
444
445
446
447
448
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
449
450
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto reverse_inclusive_scan_sequence(Seq, Reduce, Number<Init>)
451
{
Chao Liu's avatar
Chao Liu committed
452
    return typename sequence_reverse_inclusive_scan<Seq, Reduce, Init>::type{};
453
454
}

Chao Liu's avatar
Chao Liu committed
455
456
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr auto inclusive_scan_sequence(Seq, Reduce, Number<Init>)
457
{
Chao Liu's avatar
Chao Liu committed
458
    return reverse_inclusive_scan_sequence(Seq{}.Reverse(), Reduce{}, Number<Init>{}).Reverse();
459
}
460

Chao Liu's avatar
Chao Liu committed
461
template <index_t... Is>
462
__host__ __device__ constexpr auto Sequence<Is...>::PopFront()
Chao Liu's avatar
Chao Liu committed
463
{
464
    return sequence_pop_front(Type{});
Chao Liu's avatar
Chao Liu committed
465
}
Chao Liu's avatar
Chao Liu committed
466

467
468
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::PopBack()
Chao Liu's avatar
Chao Liu committed
469
{
470
    return sequence_pop_back(Type{});
Chao Liu's avatar
Chao Liu committed
471
472
}

473
474
template <index_t... Is>
__host__ __device__ constexpr auto Sequence<Is...>::Reverse()
Chao Liu's avatar
Chao Liu committed
475
{
Chao Liu's avatar
Chao Liu committed
476
    return typename sequence_reverse<Sequence<Is...>>::type{};
477
478
479
480
481
482
483
484
485
486
487
488
}

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
489
    return seq_left.PushBack(Number<X>{}).PushBack(seq_right);
Chao Liu's avatar
Chao Liu committed
490
}
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526

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...); });
}
527
528
529

} // namespace ck
#endif