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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/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 #if !defined(__HIPCC_RTC__) || !defined(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 namespace impl {
188 template <typename T, T... Ints>
190 
191 template <index_t... Ints>
192 struct __integer_sequence<index_t, Ints...>
193 {
194  using seq_type = Sequence<Ints...>;
195 };
196 } // namespace impl
197 
198 template <index_t N>
200  typename __make_integer_seq<impl::__integer_sequence, index_t, N>::seq_type;
201 
202 // merge sequence
203 template <typename Seq, typename... Seqs>
205 {
206  using type = typename sequence_merge<Seq, typename sequence_merge<Seqs...>::type>::type;
207 };
208 
209 template <index_t... Xs, index_t... Ys>
210 struct sequence_merge<Sequence<Xs...>, Sequence<Ys...>>
211 {
212  using type = Sequence<Xs..., Ys...>;
213 };
214 
215 template <typename Seq>
216 struct sequence_merge<Seq>
217 {
218  using type = Seq;
219 };
220 
221 // generate sequence
222 template <index_t NSize, typename F>
224 {
225  template <index_t IBegin, index_t NRemain, typename G>
227  {
228  static constexpr index_t NRemainLeft = NRemain / 2;
229  static constexpr index_t NRemainRight = NRemain - NRemainLeft;
230  static constexpr index_t IMiddle = IBegin + NRemainLeft;
231 
232  using type = typename sequence_merge<
235  };
236 
237  template <index_t I, typename G>
238  struct sequence_gen_impl<I, 1, G>
239  {
240  static constexpr index_t Is = G{}(Number<I>{});
242  };
243 
244  template <index_t I, typename G>
245  struct sequence_gen_impl<I, 0, G>
246  {
247  using type = Sequence<>;
248  };
249 
251 };
252 
253 // arithmetic sequence
254 template <index_t IBegin, index_t IEnd, index_t Increment>
256 {
257  struct F
258  {
259  __host__ __device__ constexpr index_t operator()(index_t i) const
260  {
261  return i * Increment + IBegin;
262  }
263  };
264 
265  using type0 = typename sequence_gen<(IEnd - IBegin) / Increment, F>::type;
266  using type1 = Sequence<>;
267 
268  static constexpr bool kHasContent =
269  (Increment > 0 && IBegin < IEnd) || (Increment < 0 && IBegin > IEnd);
270 
272 };
273 
274 template <index_t IEnd>
275 struct arithmetic_sequence_gen<0, IEnd, 1>
276 {
277  template <typename T, T... Ints>
278  struct WrapSequence
279  {
280  using type = Sequence<Ints...>;
281  };
282  // https://reviews.llvm.org/D13786
283  using type = typename __make_integer_seq<WrapSequence, index_t, IEnd>::type;
284 };
285 
286 // uniform sequence
287 template <index_t NSize, index_t I>
289 {
290  struct F
291  {
292  __host__ __device__ constexpr index_t operator()(index_t) const { return I; }
293  };
294 
296 };
297 
298 // reverse inclusive scan (with init) sequence
299 template <typename, typename, index_t>
301 
302 template <index_t I, index_t... Is, typename Reduce, index_t Init>
303 struct sequence_reverse_inclusive_scan<Sequence<I, Is...>, Reduce, Init>
304 {
305  using old_scan = typename sequence_reverse_inclusive_scan<Sequence<Is...>, Reduce, Init>::type;
306 
307  static constexpr index_t new_reduce = Reduce{}(I, old_scan{}.Front());
308 
310 };
311 
312 template <index_t I, typename Reduce, index_t Init>
314 {
315  using type = Sequence<Reduce{}(I, Init)>;
316 };
317 
318 template <typename Reduce, index_t Init>
320 {
321  using type = Sequence<>;
322 };
323 
324 // split sequence
325 template <typename Seq, index_t I>
327 {
328  static constexpr index_t NSize = Seq{}.Size();
329 
332 
333  using left_type = decltype(Seq::Extract(range0{}));
334  using right_type = decltype(Seq::Extract(range1{}));
335 };
336 
337 // reverse sequence
338 template <typename Seq>
340 {
341  static constexpr index_t NSize = Seq{}.Size();
342 
343  using seq_split = sequence_split<Seq, NSize / 2>;
344  using type = typename sequence_merge<
347 };
348 
349 template <index_t I>
351 {
352  using type = Sequence<I>;
353 };
354 
355 template <index_t I0, index_t I1>
356 struct sequence_reverse<Sequence<I0, I1>>
357 {
359 };
360 
361 #if 1
362 template <typename Reduce, typename Seq, typename... Seqs>
364 {
365  using type = typename sequence_reduce<Reduce,
366  Seq,
367  typename sequence_reduce<Reduce, Seqs...>::type>::type;
368 };
369 
370 template <typename Reduce, index_t... Xs, index_t... Ys>
371 struct sequence_reduce<Reduce, Sequence<Xs...>, Sequence<Ys...>>
372 {
373  using type = Sequence<Reduce{}(Xs, Ys)...>;
374 };
375 
376 template <typename Reduce, typename Seq>
377 struct sequence_reduce<Reduce, Seq>
378 {
379  using type = Seq;
380 };
381 #endif
382 
383 template <typename Values, typename Ids, typename Compare>
385 {
386  template <typename LeftValues,
387  typename LeftIds,
388  typename RightValues,
389  typename RightIds,
390  typename MergedValues,
391  typename MergedIds,
392  typename Comp>
394  {
395  static constexpr bool choose_left = LeftValues::Front() < RightValues::Front();
396 
397  static constexpr index_t chosen_value =
398  choose_left ? LeftValues::Front() : RightValues::Front();
399  static constexpr index_t chosen_id = choose_left ? LeftIds::Front() : RightIds::Front();
400 
401  using new_merged_values = decltype(MergedValues::PushBack(Number<chosen_value>{}));
402  using new_merged_ids = decltype(MergedIds::PushBack(Number<chosen_id>{}));
403 
405  typename conditional<choose_left, decltype(LeftValues::PopFront()), LeftValues>::type;
406  using new_left_ids =
407  typename conditional<choose_left, decltype(LeftIds::PopFront()), LeftIds>::type;
408 
410  typename conditional<choose_left, RightValues, decltype(RightValues::PopFront())>::type;
412  typename conditional<choose_left, RightIds, decltype(RightIds::PopFront())>::type;
413 
415  new_left_ids,
420  Comp>;
421  // this is output
423  using merged_ids = typename merge::merged_ids;
424  };
425 
426  template <typename LeftValues,
427  typename LeftIds,
428  typename MergedValues,
429  typename MergedIds,
430  typename Comp>
431  struct sorted_sequence_merge_impl<LeftValues,
432  LeftIds,
433  Sequence<>,
434  Sequence<>,
435  MergedValues,
436  MergedIds,
437  Comp>
438  {
441  };
442 
443  template <typename RightValues,
444  typename RightIds,
445  typename MergedValues,
446  typename MergedIds,
447  typename Comp>
449  Sequence<>,
450  RightValues,
451  RightIds,
452  MergedValues,
453  MergedIds,
454  Comp>
455  {
458  };
459 
460  template <typename LeftValues,
461  typename LeftIds,
462  typename RightValues,
463  typename RightIds,
464  typename Comp>
466  {
467  using merge = sorted_sequence_merge_impl<LeftValues,
468  LeftIds,
469  RightValues,
470  RightIds,
471  Sequence<>,
472  Sequence<>,
473  Comp>;
474 
476  using merged_ids = typename merge::merged_ids;
477  };
478 
479  static constexpr index_t nsize = Values::Size();
480 
483 
489 
495 
500  Compare>;
501 
504 };
505 
506 template <index_t ValueX, index_t ValueY, index_t IdX, index_t IdY, typename Compare>
507 struct sequence_sort_impl<Sequence<ValueX, ValueY>, Sequence<IdX, IdY>, Compare>
508 {
509  static constexpr bool choose_x = Compare{}(ValueX, ValueY);
510 
514 };
515 
516 template <index_t Value, index_t Id, typename Compare>
518 {
521 };
522 
523 template <typename Compare>
524 struct sequence_sort_impl<Sequence<>, Sequence<>, Compare>
525 {
528 };
529 
530 template <typename Values, typename Compare>
532 {
533  using unsorted_ids = typename arithmetic_sequence_gen<0, Values::Size(), 1>::type;
535 
536  // this is output
537  using type = typename sort::sorted_values;
539 };
540 
541 template <typename Values, typename Less, typename Equal>
543 {
544  template <typename RemainValues,
545  typename RemainIds,
546  typename UniquifiedValues,
547  typename UniquifiedIds,
548  typename Eq>
550  {
551  static constexpr index_t current_value = RemainValues::Front();
552  static constexpr index_t current_id = RemainIds::Front();
553 
554  static constexpr bool is_unique_value = (current_value != UniquifiedValues::Back());
555 
556  using new_remain_values = decltype(RemainValues::PopFront());
557  using new_remain_ids = decltype(RemainIds::PopFront());
558 
560  typename conditional<is_unique_value,
561  decltype(UniquifiedValues::PushBack(Number<current_value>{})),
562  UniquifiedValues>::type;
563 
565  typename conditional<is_unique_value,
566  decltype(UniquifiedIds::PushBack(Number<current_id>{})),
567  UniquifiedIds>::type;
568 
573  Eq>;
574 
575  // this is output
578  };
579 
580  template <typename UniquifiedValues, typename UniquifiedIds, typename Eq>
582  Sequence<>,
583  UniquifiedValues,
584  UniquifiedIds,
585  Eq>
586  {
587  using uniquified_values = UniquifiedValues;
588  using uniquified_ids = UniquifiedIds;
589  };
590 
591  template <typename SortedValues, typename SortedIds, typename Eq>
593  {
594  using uniquify = sorted_sequence_uniquify_impl<decltype(SortedValues::PopFront()),
595  decltype(SortedIds::PopFront()),
596  Sequence<SortedValues::Front()>,
597  Sequence<SortedIds::Front()>,
598  Eq>;
599 
602  };
603 
605  using sorted_values = typename sort::type;
607 
609 
610  // this is output
613 };
614 
615 template <typename SeqMap>
616 struct is_valid_sequence_map : is_same<typename arithmetic_sequence_gen<0, SeqMap::Size(), 1>::type,
617  typename sequence_sort<SeqMap, math::less<index_t>>::type>
618 {
619 };
620 
621 template <typename SeqMap>
623 {
624  template <typename X2Y, typename WorkingY2X, index_t XBegin, index_t XRemain>
626  {
627  static constexpr auto new_y2x =
628  WorkingY2X::Modify(X2Y::At(Number<XBegin>{}), Number<XBegin>{});
629 
630  using type =
631  typename sequence_map_inverse_impl<X2Y, decltype(new_y2x), XBegin + 1, XRemain - 1>::
632  type;
633  };
634 
635  template <typename X2Y, typename WorkingY2X, index_t XBegin>
636  struct sequence_map_inverse_impl<X2Y, WorkingY2X, XBegin, 0>
637  {
638  using type = WorkingY2X;
639  };
640 
641  using type =
642  typename sequence_map_inverse_impl<SeqMap,
643  typename uniform_sequence_gen<SeqMap::Size(), 0>::type,
644  0,
645  SeqMap::Size()>::type;
646 };
647 
648 template <index_t... Xs, index_t... Ys>
649 __host__ __device__ constexpr bool operator==(Sequence<Xs...>, Sequence<Ys...>)
650 {
651  return ((Xs == Ys) && ...);
652 }
653 
654 template <index_t... Xs, index_t... Ys>
655 __host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
656 {
657  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
658 
659  return Sequence<(Xs + Ys)...>{};
660 }
661 
662 template <index_t... Xs, index_t... Ys>
663 __host__ __device__ constexpr auto operator-(Sequence<Xs...>, Sequence<Ys...>)
664 {
665  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
666 
667  return Sequence<(Xs - Ys)...>{};
668 }
669 
670 template <index_t... Xs, index_t... Ys>
671 __host__ __device__ constexpr auto operator*(Sequence<Xs...>, Sequence<Ys...>)
672 {
673  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
674 
675  return Sequence<(Xs * Ys)...>{};
676 }
677 
678 template <index_t... Xs, index_t... Ys>
679 __host__ __device__ constexpr auto operator/(Sequence<Xs...>, Sequence<Ys...>)
680 {
681  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
682 
683  return Sequence<(Xs / Ys)...>{};
684 }
685 
686 template <index_t... Xs, index_t... Ys>
687 __host__ __device__ constexpr auto operator%(Sequence<Xs...>, Sequence<Ys...>)
688 {
689  static_assert(sizeof...(Xs) == sizeof...(Ys), "wrong! inconsistent size");
690 
691  return Sequence<(Xs % Ys)...>{};
692 }
693 
694 template <index_t... Xs, index_t Y>
695 __host__ __device__ constexpr auto operator+(Sequence<Xs...>, Number<Y>)
696 {
697  return Sequence<(Xs + Y)...>{};
698 }
699 
700 template <index_t... Xs, index_t Y>
701 __host__ __device__ constexpr auto operator-(Sequence<Xs...>, Number<Y>)
702 {
703  return Sequence<(Xs - Y)...>{};
704 }
705 
706 template <index_t... Xs, index_t Y>
707 __host__ __device__ constexpr auto operator*(Sequence<Xs...>, Number<Y>)
708 {
709  return Sequence<(Xs * Y)...>{};
710 }
711 
712 template <index_t... Xs, index_t Y>
713 __host__ __device__ constexpr auto operator/(Sequence<Xs...>, Number<Y>)
714 {
715  return Sequence<(Xs / Y)...>{};
716 }
717 
718 template <index_t... Xs, index_t Y>
719 __host__ __device__ constexpr auto operator%(Sequence<Xs...>, Number<Y>)
720 {
721  return Sequence<(Xs % Y)...>{};
722 }
723 
724 template <index_t Y, index_t... Xs>
725 __host__ __device__ constexpr auto operator+(Number<Y>, Sequence<Xs...>)
726 {
727  return Sequence<(Y + Xs)...>{};
728 }
729 
730 template <index_t Y, index_t... Xs>
731 __host__ __device__ constexpr auto operator-(Number<Y>, Sequence<Xs...>)
732 {
733  return Sequence<(Y - Xs)...>{};
734 }
735 
736 template <index_t Y, index_t... Xs>
737 __host__ __device__ constexpr auto operator*(Number<Y>, Sequence<Xs...>)
738 {
739  return Sequence<(Y * Xs)...>{};
740 }
741 
742 template <index_t Y, index_t... Xs>
743 __host__ __device__ constexpr auto operator/(Number<Y>, Sequence<Xs...>)
744 {
745  return Sequence<(Y / Xs)...>{};
746 }
747 
748 template <index_t Y, index_t... Xs>
749 __host__ __device__ constexpr auto operator%(Number<Y>, Sequence<Xs...>)
750 {
751  return Sequence<(Y % Xs)...>{};
752 }
753 
754 template <index_t I, index_t... Is>
755 __host__ __device__ constexpr auto sequence_pop_front(Sequence<I, Is...>)
756 {
757  return Sequence<Is...>{};
758 }
759 
760 template <typename Seq>
761 __host__ __device__ constexpr auto sequence_pop_back(Seq)
762 {
763  static_assert(Seq::Size() > 0, "wrong! cannot pop an empty Sequence!");
764  return sequence_pop_front(Seq::Reverse()).Reverse();
765 }
766 
767 template <typename... Seqs>
768 __host__ __device__ constexpr auto merge_sequences(Seqs...)
769 {
770  return typename sequence_merge<Seqs...>::type{};
771 }
772 
773 template <typename F, index_t... Xs>
774 __host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>)
775 {
776  return Sequence<f(Xs)...>{};
777 }
778 
779 template <typename F, index_t... Xs, index_t... Ys>
780 __host__ __device__ constexpr auto transform_sequences(F f, Sequence<Xs...>, Sequence<Ys...>)
781 {
782  static_assert(Sequence<Xs...>::mSize == Sequence<Ys...>::mSize, "Dim not the same");
783 
784  return Sequence<f(Xs, Ys)...>{};
785 }
786 
787 template <typename F, index_t... Xs, index_t... Ys, index_t... Zs>
788 __host__ __device__ constexpr auto
790 {
793  "Dim not the same");
794 
795  return Sequence<f(Xs, Ys, Zs)...>{};
796 }
797 
798 template <typename Seq, typename Reduce, index_t Init>
799 __host__ __device__ constexpr auto reverse_inclusive_scan_sequence(Seq, Reduce, Number<Init>)
800 {
802 }
803 
804 template <typename Seq, typename Reduce, index_t Init>
805 __host__ __device__ constexpr auto reverse_exclusive_scan_sequence(Seq, Reduce, Number<Init>)
806 {
807  return reverse_inclusive_scan_sequence(Seq::PopFront(), Reduce{}, Number<Init>{})
808  .PushBack(Number<Init>{});
809 }
810 
811 template <typename Seq, typename Reduce, index_t Init>
812 __host__ __device__ constexpr auto inclusive_scan_sequence(Seq, Reduce, Number<Init>)
813 {
814  return reverse_inclusive_scan_sequence(Seq{}.Reverse(), Reduce{}, Number<Init>{}).Reverse();
815 }
816 
817 template <typename Seq, index_t... Is>
818 __host__ __device__ constexpr auto pick_sequence_elements_by_ids(Seq, Sequence<Is...> /* ids */)
819 {
820  return Sequence<Seq::At(Number<Is>{})...>{};
821 }
822 
823 #if 1
824 namespace detail {
825 template <typename WorkSeq, typename RemainSeq, typename RemainMask>
827 {
828  using new_work_seq = typename conditional<RemainMask::Front(),
829  decltype(WorkSeq::PushBack(RemainSeq::Front())),
830  WorkSeq>::type;
831 
832  using type =
834  decltype(RemainSeq::PopFront()),
835  decltype(RemainMask::PopFront())>::type;
836 };
837 
838 template <typename WorkSeq>
840 {
841  using type = WorkSeq;
842 };
843 
844 } // namespace detail
845 
846 template <typename Seq, typename Mask>
847 __host__ __device__ constexpr auto pick_sequence_elements_by_mask(Seq, Mask)
848 {
849  static_assert(Seq::Size() == Mask::Size(), "wrong!");
850 
851  return typename detail::pick_sequence_elements_by_mask_impl<Sequence<>, Seq, Mask>::type{};
852 }
853 
854 namespace detail {
855 template <typename WorkSeq, typename RemainValues, typename RemainIds>
857 {
858  using new_work_seq = decltype(WorkSeq::Modify(RemainIds::Front(), RemainValues::Front()));
859 
860  using type =
862  decltype(RemainValues::PopFront()),
863  decltype(RemainIds::PopFront())>::type;
864 };
865 
866 template <typename WorkSeq>
868 {
869  using type = WorkSeq;
870 };
871 } // namespace detail
872 
873 template <typename Seq, typename Values, typename Ids>
874 __host__ __device__ constexpr auto modify_sequence_elements_by_ids(Seq, Values, Ids)
875 {
876  static_assert(Values::Size() == Ids::Size() && Seq::Size() >= Values::Size(), "wrong!");
877 
879 }
880 #endif
881 
882 template <typename Seq, typename Reduce, index_t Init>
883 __host__ __device__ constexpr index_t
884 reduce_on_sequence(Seq, Reduce f, Number<Init> /*initial_value*/)
885 {
886  index_t result = Init;
887 
888  for(index_t i = 0; i < Seq::Size(); ++i)
889  {
890  result = f(result, Seq::At(i));
891  }
892 
893  return result;
894 }
895 
896 // TODO: a generic any_of for any container
897 template <typename Seq, typename F>
898 __host__ __device__ constexpr bool sequence_any_of(Seq, F f)
899 {
900  bool flag = false;
901 
902  for(index_t i = 0; i < Seq::Size(); ++i)
903  {
904  flag = flag || f(Seq::At(i));
905  }
906 
907  return flag;
908 }
909 
910 // TODO: a generic all_of for any container
911 template <typename Seq, typename F>
912 __host__ __device__ constexpr bool sequence_all_of(Seq, F f)
913 {
914  bool flag = true;
915 
916  for(index_t i = 0; i < Seq::Size(); ++i)
917  {
918  flag = flag && f(Seq::At(i));
919  }
920 
921  return flag;
922 }
923 
924 template <typename Sx, typename Sy>
926 
927 template <index_t NSize, index_t I>
929 
930 } // namespace ck
931 
932 #if !defined(__HIPCC_RTC__) || !defined(CK_CODE_GEN_RTC)
933 template <ck::index_t... Is>
934 std::ostream& operator<<(std::ostream& os, const ck::Sequence<Is...>)
935 {
936  using S = ck::Sequence<Is...>;
937  os << "{";
938  ck::static_for<0, S::Size() - ck::Number<1>{}, 1>{}(
939  [&](auto i) { os << S::At(i).value << ", "; });
940  os << S::At(S::Size() - ck::Number<1>{}).value << "}";
941  return os;
942 }
943 #endif
Represents a JSON value. Use Value for UTF8 encoding and default allocator.
Definition: document.h:668
Definition: ck.hpp:267
typename __make_integer_seq< impl::__integer_sequence, index_t, N >::seq_type make_index_sequence
Definition: sequence.hpp:200
__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:928
__host__ constexpr __device__ auto operator/(integral_constant< TX, X >, integral_constant< TY, Y >)
Definition: integral_constant.hpp:46
__host__ constexpr __device__ index_t reduce_on_sequence(Seq, Reduce f, Number< Init >)
Definition: sequence.hpp:884
__host__ constexpr __device__ auto reverse_exclusive_scan_sequence(Seq, Reduce, Number< Init >)
Definition: sequence.hpp:805
__host__ constexpr __device__ auto modify_sequence_elements_by_ids(Seq, Values, Ids)
Definition: sequence.hpp:874
__host__ constexpr __device__ auto merge_sequences(Seqs...)
Definition: sequence.hpp:768
__host__ constexpr __device__ auto pick_sequence_elements_by_mask(Seq, Mask)
Definition: sequence.hpp:847
__host__ constexpr __device__ auto reverse_inclusive_scan_sequence(Seq, Reduce, Number< Init >)
Definition: sequence.hpp:799
__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:898
__host__ constexpr __device__ bool operator==(Sequence< Xs... >, Sequence< Ys... >)
Definition: sequence.hpp:649
__host__ constexpr __device__ auto sequence_pop_back(Seq)
Definition: sequence.hpp:761
__host__ constexpr __device__ auto sequence_pop_front(Sequence< I, Is... >)
Definition: sequence.hpp:755
typename sequence_merge< Sx, Sy >::type sequence_merge_t
Definition: sequence.hpp:925
__host__ constexpr __device__ bool sequence_all_of(Seq, F f)
Definition: sequence.hpp:912
__host__ constexpr __device__ auto transform_sequences(F f, Sequence< Xs... >)
Definition: sequence.hpp:774
__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:298
__host__ constexpr __device__ auto pick_sequence_elements_by_ids(Seq, Sequence< Is... >)
Definition: sequence.hpp:818
__host__ constexpr __device__ auto inclusive_scan_sequence(Seq, Reduce, Number< Init >)
Definition: sequence.hpp:812
__host__ constexpr __device__ auto operator%(integral_constant< TX, X >, integral_constant< TY, Y >)
Definition: integral_constant.hpp:53
const GenericPointer< typename T::ValueType > T2 value
Definition: pointer.h:1350
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:258
__host__ constexpr __device__ index_t operator()(index_t i) const
Definition: sequence.hpp:259
typename __make_integer_seq< WrapSequence, index_t, IEnd >::type type
Definition: sequence.hpp:283
Definition: sequence.hpp:256
static constexpr bool kHasContent
Definition: sequence.hpp:268
typename sequence_gen<(IEnd - IBegin)/Increment, F >::type type0
Definition: sequence.hpp:265
typename conditional< kHasContent, type0, type1 >::type type
Definition: sequence.hpp:271
Definition: functional.hpp:100
decltype(WorkSeq::Modify(RemainIds::Front(), RemainValues::Front())) new_work_seq
Definition: sequence.hpp:858
typename modify_sequence_elements_by_ids_impl< new_work_seq, decltype(RemainValues::PopFront()), decltype(RemainIds::PopFront())>::type type
Definition: sequence.hpp:863
typename pick_sequence_elements_by_mask_impl< new_work_seq, decltype(RemainSeq::PopFront()), decltype(RemainMask::PopFront())>::type type
Definition: sequence.hpp:835
typename conditional< RemainMask::Front(), decltype(WorkSeq::PushBack(RemainSeq::Front())), WorkSeq >::type new_work_seq
Definition: sequence.hpp:830
Definition: sequence.hpp:189
Definition: integral_constant.hpp:20
Definition: type.hpp:177
Definition: sequence.hpp:618
Definition: sequence.hpp:227
static constexpr index_t NRemainRight
Definition: sequence.hpp:229
static constexpr index_t IMiddle
Definition: sequence.hpp:230
static constexpr index_t NRemainLeft
Definition: sequence.hpp:228
typename sequence_merge< typename sequence_gen_impl< IBegin, NRemainLeft, G >::type, typename sequence_gen_impl< IMiddle, NRemainRight, G >::type >::type type
Definition: sequence.hpp:234
Definition: sequence.hpp:224
typename sequence_gen_impl< 0, NSize, F >::type type
Definition: sequence.hpp:250
static constexpr auto new_y2x
Definition: sequence.hpp:627
typename sequence_map_inverse_impl< X2Y, decltype(new_y2x), XBegin+1, XRemain - 1 >::type type
Definition: sequence.hpp:632
Definition: sequence.hpp:623
typename sequence_map_inverse_impl< SeqMap, typename uniform_sequence_gen< SeqMap::Size(), 0 >::type, 0, SeqMap::Size()>::type type
Definition: sequence.hpp:645
Seq type
Definition: sequence.hpp:218
Definition: sequence.hpp:205
typename sequence_merge< Seq, typename sequence_merge< Seqs... >::type >::type type
Definition: sequence.hpp:206
Seq type
Definition: sequence.hpp:379
Definition: sequence.hpp:364
typename sequence_reduce< Reduce, Seq, typename sequence_reduce< Reduce, Seqs... >::type >::type type
Definition: sequence.hpp:367
typename sequence_merge< Sequence< new_reduce >, old_scan >::type type
Definition: sequence.hpp:309
typename sequence_reverse_inclusive_scan< Sequence< Is... >, Reduce, Init >::type old_scan
Definition: sequence.hpp:305
Definition: sequence.hpp:300
Definition: sequence.hpp:340
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:346
static constexpr index_t NSize
Definition: sequence.hpp:341
typename sequence_merge< MergedValues, LeftValues >::type merged_values
Definition: sequence.hpp:439
typename sequence_merge< MergedValues, RightValues >::type merged_values
Definition: sequence.hpp:456
typename merge::merged_values merged_values
Definition: sequence.hpp:422
typename merge::merged_ids merged_ids
Definition: sequence.hpp:423
static constexpr bool choose_left
Definition: sequence.hpp:395
decltype(MergedValues::PushBack(Number< chosen_value >{})) new_merged_values
Definition: sequence.hpp:401
typename conditional< choose_left, decltype(LeftValues::PopFront()), LeftValues >::type new_left_values
Definition: sequence.hpp:405
decltype(MergedIds::PushBack(Number< chosen_id >{})) new_merged_ids
Definition: sequence.hpp:402
typename conditional< choose_left, decltype(LeftIds::PopFront()), LeftIds >::type new_left_ids
Definition: sequence.hpp:407
typename conditional< choose_left, RightValues, decltype(RightValues::PopFront())>::type new_right_values
Definition: sequence.hpp:410
static constexpr index_t chosen_value
Definition: sequence.hpp:397
typename conditional< choose_left, RightIds, decltype(RightIds::PopFront())>::type new_right_ids
Definition: sequence.hpp:412
static constexpr index_t chosen_id
Definition: sequence.hpp:399
typename merge::merged_ids merged_ids
Definition: sequence.hpp:476
typename merge::merged_values merged_values
Definition: sequence.hpp:475
typename conditional< choose_x, Sequence< ValueX, ValueY >, Sequence< ValueY, ValueX > >::type sorted_values
Definition: sequence.hpp:512
typename conditional< choose_x, Sequence< IdX, IdY >, Sequence< IdY, IdX > >::type sorted_ids
Definition: sequence.hpp:513
Definition: sequence.hpp:385
typename left_sort::sorted_values left_sorted_values
Definition: sequence.hpp:487
typename right_sort::sorted_values right_sorted_values
Definition: sequence.hpp:493
typename left_sort::sorted_ids left_sorted_ids
Definition: sequence.hpp:488
typename split_unsorted_values::left_type left_unsorted_values
Definition: sequence.hpp:484
static constexpr index_t nsize
Definition: sequence.hpp:479
typename merged_sorted::merged_values sorted_values
Definition: sequence.hpp:502
typename split_unsorted_ids::left_type left_unsorted_ids
Definition: sequence.hpp:485
typename right_sort::sorted_ids right_sorted_ids
Definition: sequence.hpp:494
typename split_unsorted_ids::right_type right_unsorted_ids
Definition: sequence.hpp:491
typename split_unsorted_values::right_type right_unsorted_values
Definition: sequence.hpp:490
typename merged_sorted::merged_ids sorted_ids
Definition: sequence.hpp:503
Definition: sequence.hpp:532
typename sort::sorted_ids sorted2unsorted_map
Definition: sequence.hpp:538
typename sort::sorted_values type
Definition: sequence.hpp:537
typename arithmetic_sequence_gen< 0, Values::Size(), 1 >::type unsorted_ids
Definition: sequence.hpp:533
Definition: sequence.hpp:327
typename arithmetic_sequence_gen< 0, I, 1 >::type range0
Definition: sequence.hpp:330
decltype(Seq::Extract(range1{})) right_type
Definition: sequence.hpp:334
static constexpr index_t NSize
Definition: sequence.hpp:328
typename arithmetic_sequence_gen< I, NSize, 1 >::type range1
Definition: sequence.hpp:331
decltype(Seq::Extract(range0{})) left_type
Definition: sequence.hpp:333
typename uniquify::uniquified_values uniquified_values
Definition: sequence.hpp:576
typename conditional< is_unique_value, decltype(UniquifiedIds::PushBack(Number< current_id >{})), UniquifiedIds >::type new_uniquified_ids
Definition: sequence.hpp:567
typename uniquify::uniquified_ids uniquified_ids
Definition: sequence.hpp:577
decltype(RemainValues::PopFront()) new_remain_values
Definition: sequence.hpp:556
static constexpr index_t current_value
Definition: sequence.hpp:551
decltype(RemainIds::PopFront()) new_remain_ids
Definition: sequence.hpp:557
static constexpr index_t current_id
Definition: sequence.hpp:552
static constexpr bool is_unique_value
Definition: sequence.hpp:554
typename conditional< is_unique_value, decltype(UniquifiedValues::PushBack(Number< current_value >{})), UniquifiedValues >::type new_uniquified_values
Definition: sequence.hpp:562
typename uniquify::uniquified_values uniquified_values
Definition: sequence.hpp:600
typename uniquify::uniquified_ids uniquified_ids
Definition: sequence.hpp:601
Definition: sequence.hpp:543
typename sort::sorted2unsorted_map sorted_ids
Definition: sequence.hpp:606
typename sort::type sorted_values
Definition: sequence.hpp:605
typename uniquify::uniquified_values type
Definition: sequence.hpp:611
typename uniquify::uniquified_ids sorted2unsorted_map
Definition: sequence.hpp:612
Definition: functional2.hpp:33
Definition: sequence.hpp:291
__host__ constexpr __device__ index_t operator()(index_t) const
Definition: sequence.hpp:292
Definition: sequence.hpp:289
typename sequence_gen< NSize, F >::type type
Definition: sequence.hpp:295
std::ostream & operator<<(std::ostream &os, const ck::Sequence< Is... >)
Definition: sequence.hpp:934