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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/utility/static_counter.hpp Source File
static_counter.hpp
Go to the documentation of this file.
1 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
2 // SPDX-License-Identifier: MIT
3 
4 #pragma once
5 
6 #include "ck_tile/core/config.hpp"
7 
8 namespace ck_tile {
9 
10 template <typename Context, index_t Start = 0, index_t Step = 1>
11 struct static_counter
12 {
13  public:
14  template <typename Unique>
15  static constexpr index_t next()
16  {
17  return next<Unique>(0) * Step + Start;
18  }
19 
20  template <unsigned long long>
21  static constexpr index_t next()
22  {
23  struct Unique
24  {
25  };
26  return next<Unique>(0) * Step + Start;
27  }
28 
29  template <typename Unique>
30  static constexpr index_t current()
31  {
32  return current<Unique>(0) * Step + Start;
33  }
34 
35  template <unsigned long long>
36  static constexpr index_t current()
37  {
38  struct Unique
39  {
40  };
41  return current<Unique>(0) * Step + Start;
42  }
43 
44  private:
45  template <index_t I>
46  struct slot
47  {
48  _Pragma("GCC diagnostic push");
49  _Pragma("GCC diagnostic ignored \"-Wundefined-internal\"");
50  friend constexpr bool slot_allocated(slot<I>);
51  _Pragma("GCC diagnostic pop");
52  };
53 
54  template <index_t I>
55  struct allocate_slot
56  {
57  friend constexpr bool slot_allocated(slot<I>) { return true; }
58  enum
59  {
60  value = I
61  };
62  };
63 
64  // If slot_allocated(slot<I>) has NOT been defined, then SFINAE will keep this function out of
65  // the overload set...
66  template <typename Unique, index_t I = 0, bool = slot_allocated(slot<I>())>
67  static constexpr index_t next(index_t)
68  {
69  return next<Unique, I + 1>(0);
70  }
71 
72  // ...And this function will be used, instead, which will define slot_allocated(slot<I>) via
73  // allocate_slot<I>.
74  template <typename Unique, index_t I = 0>
75  static constexpr index_t next(double)
76  {
77  return allocate_slot<I>::value;
78  }
79 
80  // If slot_allocated(slot<I>) has NOT been defined, then SFINAE will keep this function out of
81  // the overload set...
82  template <typename Unique, index_t I = Start, bool = slot_allocated(slot<I>())>
83  static constexpr index_t current(index_t)
84  {
85  return current<Unique, I + 1>(0);
86  }
87 
88  // ...And this function will be used, instead, which will return the current counter, or assert
89  // in case next() hasn't been called yet.
90  template <typename Unique, index_t I = Start>
91  static constexpr index_t current(double)
92  {
93  static_assert(I != 0, "You must invoke next() first");
94 
95  return I - 1;
96  }
97 };
98 
99 namespace impl {
100 template <int I>
101 struct static_counter_uniq_;
102 }
103 
104 // clang-format off
105 #define MAKE_SC() \
106  _Pragma("clang diagnostic push") \
107  _Pragma("clang diagnostic ignored \"-Wpre-c2y-compat\"") \
108  _Pragma("clang diagnostic ignored \"-Wc2y-extensions\"") \
109  ck_tile::static_counter<ck_tile::impl::static_counter_uniq_<__COUNTER__>>{} \
110  _Pragma("clang diagnostic pop")
111 #define MAKE_SC_WITH(start_, step_) \
112  _Pragma("clang diagnostic push") \
113  _Pragma("clang diagnostic ignored \"-Wpre-c2y-compat\"") \
114  _Pragma("clang diagnostic ignored \"-Wc2y-extensions\"") ck_tile:: \
115  static_counter<ck_tile::impl::static_counter_uniq_<__COUNTER__>, start_, step_>{} \
116  _Pragma("clang diagnostic pop")
117 #define NEXT_SC(c_) \
118  _Pragma("clang diagnostic push") \
119  _Pragma("clang diagnostic ignored \"-Wpre-c2y-compat\"") \
120  _Pragma("clang diagnostic ignored \"-Wc2y-extensions\"") c_.next<__COUNTER__>() \
121  _Pragma("clang diagnostic pop")
122 #define NEXT_SCI(c_, static_i_) \
123  _Pragma("clang diagnostic push") \
124  _Pragma("clang diagnostic ignored \"-Wpre-c2y-compat\"") \
125  _Pragma("clang diagnostic ignored \"-Wc2y-extensions\"") \
126  c_.next<__COUNTER__ + static_i_>() _Pragma("clang diagnostic pop")
127 // clang-format on
128 
129 // Usage:
130 // constexpr auto c = MAKE_SC()
131 // NEXT_SC(c) // -> constexpr 0
132 // NEXT_SC(c) // -> constexpr 1
133 // NEXT_SC(c) // -> constexpr 2
134 } // namespace ck_tile
ck_tile
Definition: cluster_descriptor.hpp:13
ck_tile::index_t
int32_t index_t
Definition: integer.hpp:9
value
const GenericPointer< typename T::ValueType > T2 value
Definition: pointer.h:1697
ck_tile::impl::static_counter_uniq_
Definition: static_counter.hpp:101
ck_tile::static_counter
Definition: static_counter.hpp:12
ck_tile::static_counter::current
static constexpr index_t current()
Definition: static_counter.hpp:30
ck_tile::static_counter::next
static constexpr index_t next()
Definition: static_counter.hpp:15