/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/utility/sequence.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/utility/sequence.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/utility/sequence.hpp Source File
sequence.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #ifndef CK_CODE_GEN_RTC
7 #include <ostream>
8 #endif
9 
11 #include "ck/utility/type.hpp"
13 #include "ck/utility/math.hpp"
14 
15 namespace ck {
16 
17 template <index_t, index_t, index_t>
18 struct static_for;
19 
20 template <index_t...>
21 struct Sequence;
22 
23 template <typename Seq, index_t I>
24 struct sequence_split;
25 
26 template <typename>
27 struct sequence_reverse;
28 
29 template <typename>
30 struct sequence_map_inverse;
31 
32 template <typename>
33 struct is_valid_sequence_map;
34 
35 template <index_t I, index_t... Is>
36 __host__ __device__ constexpr auto sequence_pop_front(Sequence<I, Is...>);
37 
38 template <typename Seq>
39 __host__ __device__ constexpr auto sequence_pop_back(Seq);
40 
41 template <index_t... Is>
42 struct Sequence
43 {
44  using Type = Sequence;
45  using data_type = index_t;
46 
47  static constexpr index_t mSize = sizeof...(Is);
48 
49  __host__ __device__ static constexpr auto Size() { return Number<mSize>{}; }
50 
51  __host__ __device__ static constexpr auto GetSize() { return Size(); }
52 
53  __host__ __device__ static constexpr index_t At(index_t I)
54  {
55  // the last dummy element is to prevent compiler complain about empty array, when mSize = 0
56  const index_t mData[mSize + 1] = {Is..., 0};
57  return mData[I];
58  }
59 
60  template <index_t I>
61  __host__ __device__ static constexpr auto At(Number<I>)
62  {
63  static_assert(I < mSize, "wrong! I too large");
64 
65  return Number<At(I)>{};
66  }
67 
68  template <index_t I>
69  __host__ __device__ static constexpr auto Get(Number<I>)
70  {
71  return At(Number<I>{});
72  }
73 
74  template <typename I>
75  __host__ __device__ constexpr auto operator[](I i) const
76  {
77  return At(i);
78  }
79 
80  template <index_t... IRs>
81  __host__ __device__ static constexpr auto ReorderGivenNew2Old(Sequence<IRs...> /*new2old*/)
82  {
83  static_assert(sizeof...(Is) == sizeof...(IRs),
84  "wrong! reorder map should have the same size as Sequence to be rerodered");
85 
86  static_assert(is_valid_sequence_map<Sequence<IRs...>>::value, "wrong! invalid reorder map");
87 
88  return Sequence<Type::At(Number<IRs>{})...>{};
89  }
90 
91  // MapOld2New is Sequence<...>
92  template <typename MapOld2New>
93  __host__ __device__ static constexpr auto ReorderGivenOld2New(MapOld2New)
94  {
95  static_assert(MapOld2New::Size() == Size(),
96  "wrong! reorder map should have the same size as Sequence to be rerodered");
97 
98  static_assert(is_valid_sequence_map<MapOld2New>::value, "wrong! invalid reorder map");
99 
101  }
102 
103  __host__ __device__ static constexpr auto Reverse()
104  {
105  return typename sequence_reverse<Type>::type{};
106  }
107 
108  __host__ __device__ static constexpr auto Front()
109  {
110  static_assert(mSize > 0, "wrong!");
111  return At(Number<0>{});
112  }
113 
114  __host__ __device__ static constexpr auto Back()
115  {
116  static_assert(mSize > 0, "wrong!");
117  return At(Number<mSize - 1>{});
118  }
119 
120  __host__ __device__ static constexpr auto PopFront() { return sequence_pop_front(Type{}); }
121 
122  __host__ __device__ static constexpr auto PopBack() { return sequence_pop_back(Type{}); }
123 
124  template <index_t... Xs>
125  __host__ __device__ static constexpr auto PushFront(Sequence<Xs...>)
126  {
127  return Sequence<Xs..., Is...>{};
128  }
129 
130  template <index_t... Xs>
131  __host__ __device__ static constexpr auto PushFront(Number<Xs>...)
132  {
133  return Sequence<Xs..., Is...>{};
134  }
135 
136  template <index_t... Xs>
137  __host__ __device__ static constexpr auto PushBack(Sequence<Xs...>)
138  {
139  return Sequence<Is..., Xs...>{};
140  }
141 
142  template <index_t... Xs>
143  __host__ __device__ static constexpr auto PushBack(Number<Xs>...)
144  {
145  return Sequence<Is..., Xs...>{};
146  }
147 
148  template <index_t... Ns>
149  __host__ __device__ static constexpr auto Extract(Number<Ns>...)
150  {
151  return Sequence<Type::At(Number<Ns>{})...>{};
152  }
153 
154  template <index_t... Ns>
155  __host__ __device__ static constexpr auto Extract(Sequence<Ns...>)
156  {
157  return Sequence<Type::At(Number<Ns>{})...>{};
158  }
159 
160  template <index_t I, index_t X>
161  __host__ __device__ static constexpr auto Modify(Number<I>, Number<X>)
162  {
163  static_assert(I < Size(), "wrong!");
164 
165  using seq_split = sequence_split<Type, I>;
166  constexpr auto seq_left = typename seq_split::left_type{};
167  constexpr auto seq_right = typename seq_split::right_type{}.PopFront();
168 
169  return seq_left.PushBack(Number<X>{}).PushBack(seq_right);
170  }
171 
172  template <typename F>
173  __host__ __device__ static constexpr auto Transform(F f)
174  {
175  return Sequence<f(Is)...>{};
176  }
177 
178  __host__ __device__ static void Print()
179  {
180  printf("{");
181  printf("size %d, ", index_t{Size()});
182  static_for<0, Size(), 1>{}([&](auto i) { printf("%d ", At(i).value); });
183  printf("}");
184  }
185 };
186 
187 // merge sequence
188 template <typename Seq, typename... Seqs>
190 {
191  using type = typename sequence_merge<Seq, typename sequence_merge<Seqs...>::type>::type;
192 };
193 
194 template <index_t... Xs, index_t... Ys>
195 struct sequence_merge<Sequence<Xs...>, Sequence<Ys...>>
196 {
197  using type = Sequence<Xs..., Ys...>;
198 };
199 
200 template <typename Seq>
201 struct sequence_merge<Seq>
202 {
203  using type = Seq;
204 };
205 
206 // generate sequence
207 template <index_t NSize, typename F>
209 {
210  template <index_t IBegin, index_t NRemain, typename G>
212  {
213  static constexpr index_t NRemainLeft = NRemain / 2;
214  static constexpr index_t NRemainRight = NRemain - NRemainLeft;
215  static constexpr index_t IMiddle = IBegin + NRemainLeft;
216 
217  using type = typename sequence_merge<
220  };
221 
222  template <index_t I, typename G>
223  struct sequence_gen_impl<I, 1, G>
224  {
225  static constexpr index_t Is = G{}(Number<I>{});
227  };
228 
229  template <index_t I, typename G>
230  struct sequence_gen_impl<I, 0, G>
231  {
232  using type = Sequence<>;
233  };
234 
236 };
237 
238 // arithmetic sequence
239 template <index_t IBegin, index_t IEnd, index_t Increment>
241 {
242  struct F
243  {
244  __host__ __device__ constexpr index_t operator()(index_t i) const
245  {
246  return i * Increment + IBegin;
247  }
248  };
249 
250  using type0 = typename sequence_gen<(IEnd - IBegin) / Increment, F>::type;
251  using type1 = Sequence<>;
252 
253  static constexpr bool kHasContent =
254  (Increment > 0 && IBegin < IEnd) || (Increment < 0 && IBegin > IEnd);
255 
257 };
258 
259 // uniform sequence
260 template <index_t NSize, index_t I>
262 {
263  struct F
264  {
265  __host__ __device__ constexpr index_t operator()(index_t) const { return I; }
266  };
267 
269 };
270 
271 // reverse inclusive scan (with init) sequence
272 template <typename, typename, index_t>
274 
275 template <index_t I, index_t... Is, typename Reduce, index_t Init>
276 struct sequence_reverse_inclusive_scan<Sequence<I, Is...>, Reduce, Init>
277 {
278  using old_scan = typename sequence_reverse_inclusive_scan<Sequence<Is...>, Reduce, Init>::type;
279 
280  static constexpr index_t new_reduce = Reduce{}(I, old_scan{}.Front());
281 
283 };
284 
285 template <index_t I, typename Reduce, index_t Init>
287 {
288  using type = Sequence<Reduce{}(I, Init)>;
289 };
290 
291 template <typename Reduce, index_t Init>
293 {
294  using type = Sequence<>;
295 };
296 
297 // split sequence
298 template <typename Seq, index_t I>
300 {
301  static constexpr index_t NSize = Seq{}.Size();
302 
305 
306  using left_type = decltype(Seq::Extract(range0{}));
307  using right_type = decltype(Seq::Extract(range1{}));
308 };
309 
310 // reverse sequence
311 template <typename Seq>
313 {
314  static constexpr index_t NSize = Seq{}.Size();
315 
316  using seq_split = sequence_split<Seq, NSize / 2>;
317  using type = typename sequence_merge<
320 };
321 
322 template <index_t I>
324 {
325  using type = Sequence<I>;
326 };
327 
328 template <index_t I0, index_t I1>
329 struct sequence_reverse<Sequence<I0, I1>>
330 {
332 };
333 
334 #if 1
335 template <typename Reduce, typename Seq, typename... Seqs>
337 {
338  using type = typename sequence_reduce<Reduce,
339  Seq,
340  typename sequence_reduce<Reduce, Seqs...>::type>::type;
341 };
342 
343 template <typename Reduce, index_t... Xs, index_t... Ys>
344 struct sequence_reduce<Reduce, Sequence<Xs...>, Sequence<Ys...>>
345 {
346  using type = Sequence<Reduce{}(Xs, Ys)...>;
347 };
348 
349 template <typename Reduce, typename Seq>
350 struct sequence_reduce<Reduce, Seq>
351 {
352  using type = Seq;
353 };
354 #endif
355 
356 template <typename Values, typename Ids, typename Compare>
358 {
359  template <typename LeftValues,
360  typename LeftIds,
361  typename RightValues,
362  typename RightIds,
363  typename MergedValues,
364  typename MergedIds,
365  typename Comp>
367  {
368  static constexpr bool choose_left = LeftValues::Front() < RightValues::Front();
369 
370  static constexpr index_t chosen_value =
371  choose_left ? LeftValues::Front() : RightValues::Front();
372  static constexpr index_t chosen_id = choose_left ? LeftIds::Front() : RightIds::Front();
373 
374  using new_merged_values = decltype(MergedValues::PushBack(Number<chosen_value>{}));
375  using new_merged_ids = decltype(MergedIds::PushBack(Number<chosen_id>{}));
376 
378  typename conditional<choose_left, decltype(LeftValues::PopFront()), LeftValues>::type;
379  using new_left_ids =
380  typename conditional<choose_left, decltype(LeftIds::PopFront()), LeftIds>::type;
381 
383  typename conditional<choose_left, RightValues, decltype(RightValues::PopFront())>::type;
385  typename conditional<choose_left, RightIds, decltype(RightIds::PopFront())>::type;
386 
388  new_left_ids,
393  Comp>;
394  // this is output
396  using merged_ids = typename merge::merged_ids;
397  };
398 
399  template <typename LeftValues,
400  typename LeftIds,
401  typename MergedValues,
402  typename MergedIds,
403  typename Comp>
404  struct sorted_sequence_merge_impl<LeftValues,
405  LeftIds,
406  Sequence<>,
407  Sequence<>,
408  MergedValues,
409  MergedIds,
410  Comp>
411  {
414  };
415 
416  template <typename RightValues,
417  typename RightIds,
418  typename MergedValues,
419  typename MergedIds,
420  typename Comp>
422  Sequence<>,
423  RightValues,
424  RightIds,
425  MergedValues,
426  MergedIds,
427  Comp>
428  {
431  };
432 
433  template <typename LeftValues,
434  typename LeftIds,
435  typename RightValues,
436  typename RightIds,
437  typename Comp>
439  {
440  using merge = sorted_sequence_merge_impl<LeftValues,
441  LeftIds,
442  RightValues,
443  RightIds,
444  Sequence<>,
445  Sequence<>,
446  Comp>;
447 
449  using merged_ids = typename merge::merged_ids;
450  };
451 
452  static constexpr index_t nsize = Values::Size();
453 
456 
462 
468 
473  Compare>;
474 
477 };
478 
479 template <index_t ValueX, index_t ValueY, index_t IdX, index_t IdY, typename Compare>
480 struct sequence_sort_impl<Sequence<ValueX, ValueY>, Sequence<IdX, IdY>, Compare>
481 {
482  static constexpr bool choose_x = Compare{}(ValueX, ValueY);
483 
487 };
488 
489 template <index_t Value, index_t Id, typename Compare>
490 struct sequence_sort_impl<Sequence<Value>, Sequence<Id>, Compare>
491 {
494 };
495 
496 template <typename Compare>
497 struct sequence_sort_impl<Sequence<>, Sequence<>, Compare>
498 {
501 };
502 
503 template <typename Values, typename Compare>
505 {
506  using unsorted_ids = typename arithmetic_sequence_gen<0, Values::Size(), 1>::type;
508 
509  // this is output
510  using type = typename sort::sorted_values;
512 };
513 
514 template <typename Values, typename Less, typename Equal>
516 {
517  template <typename RemainValues,
518  typename RemainIds,
519  typename UniquifiedValues,
520  typename UniquifiedIds,
521  typename Eq>
523  {
524  static constexpr index_t current_value = RemainValues::Front();
525  static constexpr index_t current_id = RemainIds::Front();
526 
527  static constexpr bool is_unique_value = (current_value != UniquifiedValues::Back());
528 
529  using new_remain_values = decltype(RemainValues::PopFront());
530  using new_remain_ids = decltype(RemainIds::PopFront());
531 
533  typename conditional<is_unique_value,
534  decltype(UniquifiedValues::PushBack(Number<current_value>{})),
535  UniquifiedValues>::type;
536 
538  typename conditional<is_unique_value,
539  decltype(UniquifiedIds::PushBack(Number<current_id>{})),
540  UniquifiedIds>::type;
541 
546  Eq>;
547 
548  // this is output
551  };
552 
553  template <typename UniquifiedValues, typename UniquifiedIds, typename Eq>
555  Sequence<>,
556  UniquifiedValues,
557  UniquifiedIds,
558  Eq>
559  {
560  using uniquified_values = UniquifiedValues;
561  using uniquified_ids = UniquifiedIds;
562  };
563 
564  template <typename SortedValues, typename SortedIds, typename Eq>
566  {
567  using uniquify = sorted_sequence_uniquify_impl<decltype(SortedValues::PopFront()),
568  decltype(SortedIds::PopFront()),
569  Sequence<SortedValues::Front()>,
570  Sequence<SortedIds::Front()>,
571  Eq>;
572 
575  };
576 
578  using sorted_values = typename sort::type;
580 
582 
583  // this is output
586 };
587 
588 template <typename SeqMap>
589 struct is_valid_sequence_map : is_same<typename arithmetic_sequence_gen<0, SeqMap::Size(), 1>::type,
590  typename sequence_sort<SeqMap, math::less<index_t>>::type>
591 {
592 };
593 
594 template <typename SeqMap>
596 {
597  template <typename X2Y, typename WorkingY2X, index_t XBegin, index_t XRemain>
599  {
600  static constexpr auto new_y2x =
601  WorkingY2X::Modify(X2Y::At(Number<XBegin>{}), Number<XBegin>{});
602 
603  using type =
604  typename sequence_map_inverse_impl<X2Y, decltype(new_y2x), XBegin + 1, XRemain - 1>::
605  type;
606  };
607 
608  template <typename X2Y, typename WorkingY2X, index_t XBegin>
609  struct sequence_map_inverse_impl<X2Y, WorkingY2X, XBegin, 0>
610  {
611  using type = WorkingY2X;
612  };
613 
614  using type =
615  typename sequence_map_inverse_impl<SeqMap,
616  typename uniform_sequence_gen<SeqMap::Size(), 0>::type,
617  0,
618  SeqMap::Size()>::type;
619 };
620 
621 template <index_t... Xs, index_t... Ys>
622 __host__ __device__ constexpr bool operator==(Sequence<Xs...>, Sequence<Ys...>)
623 {
624  return ((Xs == Ys) && ...);
625 }
626 
627 template <index_t... Xs, index_t... Ys>
628 __host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
629 {
630  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
631 
632  return Sequence<(Xs + Ys)...>{};
633 }
634 
635 template <index_t... Xs, index_t... Ys>
636 __host__ __device__ constexpr auto operator-(Sequence<Xs...>, Sequence<Ys...>)
637 {
638  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
639 
640  return Sequence<(Xs - Ys)...>{};
641 }
642 
643 template <index_t... Xs, index_t... Ys>
644 __host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
645 {
646  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
647 
648  return Sequence<(Xs * Ys)...>{};
649 }
650 
651 template <index_t... Xs, index_t... Ys>
652 __host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
653 {
654  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
655 
656  return Sequence<(Xs / Ys)...>{};
657 }
658 
659 template <index_t... Xs, index_t... Ys>
660 __host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
661 {
662  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
663 
664  return Sequence<(Xs % Ys)...>{};
665 }
666 
667 template <index_t... Xs, index_t Y>
668 __host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
669 {
670  return Sequence<(Xs + Y)...>{};
671 }
672 
673 template <index_t... Xs, index_t Y>
674 __host__ __device__ constexpr auto operator-(Sequence<Xs...>, Number<Y>)
675 {
676  return Sequence<(Xs - Y)...>{};
677 }
678 
679 template <index_t... Xs, index_t Y>
680 __host__ __device__ constexpr auto operator*(Sequence<Xs...>, Number<Y>)
681 {
682  return Sequence<(Xs * Y)...>{};
683 }
684 
685 template <index_t... Xs, index_t Y>
686 __host__ __device__ constexpr auto operator/(Sequence<Xs...>, Number<Y>)
687 {
688  return Sequence<(Xs / Y)...>{};
689 }
690 
691 template <index_t... Xs, index_t Y>
692 __host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
693 {
694  return Sequence<(Xs % Y)...>{};
695 }
696 
697 template <index_t Y, index_t... Xs>
698 __host__ __device__ constexpr auto operator+(Number<Y>, Sequence<Xs...>)
699 {
700  return Sequence<(Y + Xs)...>{};
701 }
702 
703 template <index_t Y, index_t... Xs>
704 __host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
705 {
706  return Sequence<(Y - Xs)...>{};
707 }
708 
709 template <index_t Y, index_t... Xs>
710 __host__ __device__ constexpr auto operator*(Number<Y>, Sequence<Xs...>)
711 {
712  return Sequence<(Y * Xs)...>{};
713 }
714 
715 template <index_t Y, index_t... Xs>
716 __host__ __device__ constexpr auto operator/(Number<Y>, Sequence<Xs...>)
717 {
718  return Sequence<(Y / Xs)...>{};
719 }
720 
721 template <index_t Y, index_t... Xs>
722 __host__ __device__ constexpr auto operator%(Number<Y>, Sequence<Xs...>)
723 {
724  return Sequence<(Y % Xs)...>{};
725 }
726 
727 template <index_t I, index_t... Is>
728 __host__ __device__ constexpr auto sequence_pop_front(Sequence<I, Is...>)
729 {
730  return Sequence<Is...>{};
731 }
732 
733 template <typename Seq>
734 __host__ __device__ constexpr auto sequence_pop_back(Seq)
735 {
736  static_assert(Seq::Size() > 0, "wrong! cannot pop an empty Sequence!");
737  return sequence_pop_front(Seq::Reverse()).Reverse();
738 }
739 
740 template <typename... Seqs>
741 __host__ __device__ constexpr auto merge_sequences(Seqs...)
742 {
743  return typename sequence_merge<Seqs...>::type{};
744 }
745 
746 template <typename F, index_t... Xs>
747 __host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>)
748 {
749  return Sequence<f(Xs)...>{};
750 }
751 
752 template <typename F, index_t... Xs, index_t... Ys>
753 __host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>)
754 {
755  static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize, "Dim not the same");
756 
757  return Sequence<f(Xs, Ys)...>{};
758 }
759 
760 template <typename F, index_t... Xs, index_t... Ys, index_t... Zs>
761 __host__ __device__ constexpr auto
763 {
766  "Dim not the same");
767 
768  return Sequence<f(Xs, Ys, Zs)...>{};
769 }
770 
771 template <typename Seq, typename Reduce, index_t Init>
772 __host__ __device__ constexpr auto reverse_inclusive_scan_sequence(Seq, Reduce, Number<Init>)
773 {
775 }
776 
777 template <typename Seq, typename Reduce, index_t Init>
778 __host__ __device__ constexpr auto reverse_exclusive_scan_sequence(Seq, Reduce, Number<Init>)
779 {
780  return reverse_inclusive_scan_sequence(Seq::PopFront(), Reduce{}, Number<Init>{})
781  .PushBack(Number<Init>{});
782 }
783 
784 template <typename Seq, typename Reduce, index_t Init>
785 __host__ __device__ constexpr auto inclusive_scan_sequence(Seq, Reduce, Number<Init>)
786 {
787  return reverse_inclusive_scan_sequence(Seq{}.Reverse(), Reduce{}, Number<Init>{}).Reverse();
788 }
789 
790 template <typename Seq, index_t... Is>
791 __host__ __device__ constexpr auto pick_sequence_elements_by_ids(Seq, Sequence<Is...> /* ids */)
792 {
793  return Sequence<Seq::At(Number<Is>{})...>{};
794 }
795 
796 #if 1
797 namespace detail {
798 template <typename WorkSeq, typename RemainSeq, typename RemainMask>
800 {
801  using new_work_seq = typename conditional<RemainMask::Front(),
802  decltype(WorkSeq::PushBack(RemainSeq::Front())),
803  WorkSeq>::type;
804 
805  using type =
807  decltype(RemainSeq::PopFront()),
808  decltype(RemainMask::PopFront())>::type;
809 };
810 
811 template <typename WorkSeq>
813 {
814  using type = WorkSeq;
815 };
816 
817 } // namespace detail
818 
819 template <typename Seq, typename Mask>
820 __host__ __device__ constexpr auto pick_sequence_elements_by_mask(Seq, Mask)
821 {
822  static_assert(Seq::Size() == Mask::Size(), "wrong!");
823 
824  return typename detail::pick_sequence_elements_by_mask_impl<Sequence<>, Seq, Mask>::type{};
825 }
826 
827 namespace detail {
828 template <typename WorkSeq, typename RemainValues, typename RemainIds>
830 {
831  using new_work_seq = decltype(WorkSeq::Modify(RemainIds::Front(), RemainValues::Front()));
832 
833  using type =
835  decltype(RemainValues::PopFront()),
836  decltype(RemainIds::PopFront())>::type;
837 };
838 
839 template <typename WorkSeq>
841 {
842  using type = WorkSeq;
843 };
844 } // namespace detail
845 
846 template <typename Seq, typename Values, typename Ids>
847 __host__ __device__ constexpr auto modify_sequence_elements_by_ids(Seq, Values, Ids)
848 {
849  static_assert(Values::Size() == Ids::Size() && Seq::Size() >= Values::Size(), "wrong!");
850 
852 }
853 #endif
854 
855 template <typename Seq, typename Reduce, index_t Init>
856 __host__ __device__ constexpr index_t
857 reduce_on_sequence(Seq, Reduce f, Number<Init> /*initial_value*/)
858 {
859  index_t result = Init;
860 
861  for(index_t i = 0; i < Seq::Size(); ++i)
862  {
863  result = f(result, Seq::At(i));
864  }
865 
866  return result;
867 }
868 
869 // TODO: a generic any_of for any container
870 template <typename Seq, typename F>
871 __host__ __device__ constexpr bool sequence_any_of(Seq, F f)
872 {
873  bool flag = false;
874 
875  for(index_t i = 0; i < Seq::Size(); ++i)
876  {
877  flag = flag || f(Seq::At(i));
878  }
879 
880  return flag;
881 }
882 
883 // TODO: a generic all_of for any container
884 template <typename Seq, typename F>
885 __host__ __device__ constexpr bool sequence_all_of(Seq, F f)
886 {
887  bool flag = true;
888 
889  for(index_t i = 0; i < Seq::Size(); ++i)
890  {
891  flag = flag && f(Seq::At(i));
892  }
893 
894  return flag;
895 }
896 
897 template <typename Sx, typename Sy>
899 
900 template <index_t NSize, index_t I>
902 
903 } // namespace ck
904 
905 #ifndef CK_CODE_GEN_RTC
906 template <ck::index_t... Is>
907 std::ostream& operator<<(std::ostream& os, const ck::Sequence<Is...>)
908 {
909  using S = ck::Sequence<Is...>;
910  os << "{";
911  ck::static_for<0, S::Size() - ck::Number<1>{}, 1>{}(
912  [&](auto i) { os << S::At(i).value << ", "; });
913  os << S::At(S::Size() - ck::Number<1>{}).value << "}";
914  return os;
915 }
916 #endif
Definition: ck.hpp:264
__host__ constexpr __device__ auto operator+(const MultiIndex< NSize > &a, const T &b)
Definition: array_multi_index.hpp:50
typename uniform_sequence_gen< NSize, I >::type uniform_sequence_gen_t
Definition: sequence.hpp:901
__host__ constexpr __device__ auto operator/(integral_constant< TX, X >, integral_constant< TY, Y >)
Definition: integral_constant.hpp:38
__host__ constexpr __device__ index_t reduce_on_sequence(Seq, Reduce f, Number< Init >)
Definition: sequence.hpp:857
__host__ constexpr __device__ auto reverse_exclusive_scan_sequence(Seq, Reduce, Number< Init >)
Definition: sequence.hpp:778
__host__ constexpr __device__ auto modify_sequence_elements_by_ids(Seq, Values, Ids)
Definition: sequence.hpp:847
__host__ constexpr __device__ auto merge_sequences(Seqs...)
Definition: sequence.hpp:741
__host__ constexpr __device__ auto pick_sequence_elements_by_mask(Seq, Mask)
Definition: sequence.hpp:820
__host__ constexpr __device__ auto reverse_inclusive_scan_sequence(Seq, Reduce, Number< Init >)
Definition: sequence.hpp:772
__host__ constexpr __device__ auto operator-(const MultiIndex< NSize > &a, const T &b)
Definition: array_multi_index.hpp:60
__host__ constexpr __device__ bool sequence_any_of(Seq, F f)
Definition: sequence.hpp:871
__host__ constexpr __device__ bool operator==(Sequence< Xs... >, Sequence< Ys... >)
Definition: sequence.hpp:622
__host__ constexpr __device__ auto sequence_pop_back(Seq)
Definition: sequence.hpp:734
__host__ constexpr __device__ auto sequence_pop_front(Sequence< I, Is... >)
Definition: sequence.hpp:728
typename sequence_merge< Sx, Sy >::type sequence_merge_t
Definition: sequence.hpp:898
__host__ constexpr __device__ bool sequence_all_of(Seq, F f)
Definition: sequence.hpp:885
__host__ constexpr __device__ auto transform_sequences(F f, Sequence< Xs... >)
Definition: sequence.hpp:747
__host__ constexpr __device__ auto operator*(const MultiIndex< NSize > &a, const T &b)
Definition: array_multi_index.hpp:70
int32_t index_t
Definition: ck.hpp:289
__host__ constexpr __device__ auto pick_sequence_elements_by_ids(Seq, Sequence< Is... >)
Definition: sequence.hpp:791
__host__ constexpr __device__ auto inclusive_scan_sequence(Seq, Reduce, Number< Init >)
Definition: sequence.hpp:785
__host__ constexpr __device__ auto operator%(integral_constant< TX, X >, integral_constant< TY, Y >)
Definition: integral_constant.hpp:45
Definition: sequence.hpp:43
__host__ static constexpr __device__ auto PushBack(Sequence< Xs... >)
Definition: sequence.hpp:137
__host__ static constexpr __device__ auto GetSize()
Definition: sequence.hpp:51
index_t data_type
Definition: sequence.hpp:45
__host__ static constexpr __device__ auto Front()
Definition: sequence.hpp:108
__host__ static constexpr __device__ auto Get(Number< I >)
Definition: sequence.hpp:69
__host__ static constexpr __device__ auto PopBack()
Definition: sequence.hpp:122
__host__ static constexpr __device__ auto ReorderGivenOld2New(MapOld2New)
Definition: sequence.hpp:93
__host__ static constexpr __device__ auto Extract(Number< Ns >...)
Definition: sequence.hpp:149
__host__ static constexpr __device__ auto At(Number< I >)
Definition: sequence.hpp:61
__host__ static constexpr __device__ auto Back()
Definition: sequence.hpp:114
__host__ static constexpr __device__ index_t At(index_t I)
Definition: sequence.hpp:53
__host__ static constexpr __device__ auto PushBack(Number< Xs >...)
Definition: sequence.hpp:143
__host__ static constexpr __device__ auto ReorderGivenNew2Old(Sequence< IRs... >)
Definition: sequence.hpp:81
__host__ constexpr __device__ auto operator[](I i) const
Definition: sequence.hpp:75
__host__ static constexpr __device__ auto Reverse()
Definition: sequence.hpp:103
__host__ static constexpr __device__ auto PopFront()
Definition: sequence.hpp:120
__host__ static constexpr __device__ auto Modify(Number< I >, Number< X >)
Definition: sequence.hpp:161
static constexpr index_t mSize
Definition: sequence.hpp:47
__host__ static constexpr __device__ auto PushFront(Number< Xs >...)
Definition: sequence.hpp:131
__host__ static constexpr __device__ auto Extract(Sequence< Ns... >)
Definition: sequence.hpp:155
__host__ static constexpr __device__ auto Size()
Definition: sequence.hpp:49
__host__ static constexpr __device__ auto Transform(F f)
Definition: sequence.hpp:173
__host__ static constexpr __device__ auto PushFront(Sequence< Xs... >)
Definition: sequence.hpp:125
__host__ static __device__ void Print()
Definition: sequence.hpp:178
Definition: sequence.hpp:243
__host__ constexpr __device__ index_t operator()(index_t i) const
Definition: sequence.hpp:244
Definition: sequence.hpp:241
static constexpr bool kHasContent
Definition: sequence.hpp:253
typename sequence_gen<(IEnd - IBegin)/Increment, F >::type type0
Definition: sequence.hpp:250
typename conditional< kHasContent, type0, type1 >::type type
Definition: sequence.hpp:256
Definition: functional.hpp:100
decltype(WorkSeq::Modify(RemainIds::Front(), RemainValues::Front())) new_work_seq
Definition: sequence.hpp:831
typename modify_sequence_elements_by_ids_impl< new_work_seq, decltype(RemainValues::PopFront()), decltype(RemainIds::PopFront())>::type type
Definition: sequence.hpp:836
typename pick_sequence_elements_by_mask_impl< new_work_seq, decltype(RemainSeq::PopFront()), decltype(RemainMask::PopFront())>::type type
Definition: sequence.hpp:808
typename conditional< RemainMask::Front(), decltype(WorkSeq::PushBack(RemainSeq::Front())), WorkSeq >::type new_work_seq
Definition: sequence.hpp:803
Definition: integral_constant.hpp:10
Definition: type.hpp:177
Definition: sequence.hpp:591
Definition: sequence.hpp:212
static constexpr index_t NRemainRight
Definition: sequence.hpp:214
static constexpr index_t IMiddle
Definition: sequence.hpp:215
static constexpr index_t NRemainLeft
Definition: sequence.hpp:213
typename sequence_merge< typename sequence_gen_impl< IBegin, NRemainLeft, G >::type, typename sequence_gen_impl< IMiddle, NRemainRight, G >::type >::type type
Definition: sequence.hpp:219
Definition: sequence.hpp:209
typename sequence_gen_impl< 0, NSize, F >::type type
Definition: sequence.hpp:235
static constexpr auto new_y2x
Definition: sequence.hpp:600
typename sequence_map_inverse_impl< X2Y, decltype(new_y2x), XBegin+1, XRemain - 1 >::type type
Definition: sequence.hpp:605
Definition: sequence.hpp:596
typename sequence_map_inverse_impl< SeqMap, typename uniform_sequence_gen< SeqMap::Size(), 0 >::type, 0, SeqMap::Size()>::type type
Definition: sequence.hpp:618
Seq type
Definition: sequence.hpp:203
Definition: sequence.hpp:190
typename sequence_merge< Seq, typename sequence_merge< Seqs... >::type >::type type
Definition: sequence.hpp:191
Seq type
Definition: sequence.hpp:352
Definition: sequence.hpp:337
typename sequence_reduce< Reduce, Seq, typename sequence_reduce< Reduce, Seqs... >::type >::type type
Definition: sequence.hpp:340
typename sequence_merge< Sequence< new_reduce >, old_scan >::type type
Definition: sequence.hpp:282
typename sequence_reverse_inclusive_scan< Sequence< Is... >, Reduce, Init >::type old_scan
Definition: sequence.hpp:278
Definition: sequence.hpp:273
Definition: sequence.hpp:313
typename sequence_merge< typename sequence_reverse< typename seq_split::right_type >::type, typename sequence_reverse< typename seq_split::left_type >::type >::type type
Definition: sequence.hpp:319
static constexpr index_t NSize
Definition: sequence.hpp:314
typename sequence_merge< MergedValues, LeftValues >::type merged_values
Definition: sequence.hpp:412
typename sequence_merge< MergedValues, RightValues >::type merged_values
Definition: sequence.hpp:429
typename merge::merged_values merged_values
Definition: sequence.hpp:395
typename merge::merged_ids merged_ids
Definition: sequence.hpp:396
static constexpr bool choose_left
Definition: sequence.hpp:368
decltype(MergedValues::PushBack(Number< chosen_value >{})) new_merged_values
Definition: sequence.hpp:374
typename conditional< choose_left, decltype(LeftValues::PopFront()), LeftValues >::type new_left_values
Definition: sequence.hpp:378
decltype(MergedIds::PushBack(Number< chosen_id >{})) new_merged_ids
Definition: sequence.hpp:375
typename conditional< choose_left, decltype(LeftIds::PopFront()), LeftIds >::type new_left_ids
Definition: sequence.hpp:380
typename conditional< choose_left, RightValues, decltype(RightValues::PopFront())>::type new_right_values
Definition: sequence.hpp:383
static constexpr index_t chosen_value
Definition: sequence.hpp:370
typename conditional< choose_left, RightIds, decltype(RightIds::PopFront())>::type new_right_ids
Definition: sequence.hpp:385
static constexpr index_t chosen_id
Definition: sequence.hpp:372
typename merge::merged_ids merged_ids
Definition: sequence.hpp:449
typename merge::merged_values merged_values
Definition: sequence.hpp:448
typename conditional< choose_x, Sequence< ValueX, ValueY >, Sequence< ValueY, ValueX > >::type sorted_values
Definition: sequence.hpp:485
typename conditional< choose_x, Sequence< IdX, IdY >, Sequence< IdY, IdX > >::type sorted_ids
Definition: sequence.hpp:486
Definition: sequence.hpp:358
typename left_sort::sorted_values left_sorted_values
Definition: sequence.hpp:460
typename right_sort::sorted_values right_sorted_values
Definition: sequence.hpp:466
typename left_sort::sorted_ids left_sorted_ids
Definition: sequence.hpp:461
typename split_unsorted_values::left_type left_unsorted_values
Definition: sequence.hpp:457
static constexpr index_t nsize
Definition: sequence.hpp:452
typename merged_sorted::merged_values sorted_values
Definition: sequence.hpp:475
typename split_unsorted_ids::left_type left_unsorted_ids
Definition: sequence.hpp:458
typename right_sort::sorted_ids right_sorted_ids
Definition: sequence.hpp:467
typename split_unsorted_ids::right_type right_unsorted_ids
Definition: sequence.hpp:464
typename split_unsorted_values::right_type right_unsorted_values
Definition: sequence.hpp:463
typename merged_sorted::merged_ids sorted_ids
Definition: sequence.hpp:476
Definition: sequence.hpp:505
typename sort::sorted_ids sorted2unsorted_map
Definition: sequence.hpp:511
typename sort::sorted_values type
Definition: sequence.hpp:510
typename arithmetic_sequence_gen< 0, Values::Size(), 1 >::type unsorted_ids
Definition: sequence.hpp:506
Definition: sequence.hpp:300
typename arithmetic_sequence_gen< 0, I, 1 >::type range0
Definition: sequence.hpp:303
decltype(Seq::Extract(range1{})) right_type
Definition: sequence.hpp:307
static constexpr index_t NSize
Definition: sequence.hpp:301
typename arithmetic_sequence_gen< I, NSize, 1 >::type range1
Definition: sequence.hpp:304
decltype(Seq::Extract(range0{})) left_type
Definition: sequence.hpp:306
typename uniquify::uniquified_values uniquified_values
Definition: sequence.hpp:549
typename conditional< is_unique_value, decltype(UniquifiedIds::PushBack(Number< current_id >{})), UniquifiedIds >::type new_uniquified_ids
Definition: sequence.hpp:540
typename uniquify::uniquified_ids uniquified_ids
Definition: sequence.hpp:550
decltype(RemainValues::PopFront()) new_remain_values
Definition: sequence.hpp:529
static constexpr index_t current_value
Definition: sequence.hpp:524
decltype(RemainIds::PopFront()) new_remain_ids
Definition: sequence.hpp:530
static constexpr index_t current_id
Definition: sequence.hpp:525
static constexpr bool is_unique_value
Definition: sequence.hpp:527
typename conditional< is_unique_value, decltype(UniquifiedValues::PushBack(Number< current_value >{})), UniquifiedValues >::type new_uniquified_values
Definition: sequence.hpp:535
typename uniquify::uniquified_values uniquified_values
Definition: sequence.hpp:573
typename uniquify::uniquified_ids uniquified_ids
Definition: sequence.hpp:574
Definition: sequence.hpp:516
typename sort::sorted2unsorted_map sorted_ids
Definition: sequence.hpp:579
typename sort::type sorted_values
Definition: sequence.hpp:578
typename uniquify::uniquified_values type
Definition: sequence.hpp:584
typename uniquify::uniquified_ids sorted2unsorted_map
Definition: sequence.hpp:585
Definition: functional2.hpp:31
Definition: sequence.hpp:264
__host__ constexpr __device__ index_t operator()(index_t) const
Definition: sequence.hpp:265
Definition: sequence.hpp:262
typename sequence_gen< NSize, F >::type type
Definition: sequence.hpp:268
std::ostream & operator<<(std::ostream &os, const ck::Sequence< Is... >)
Definition: sequence.hpp:907