From 616932068deb4507e51691c068be98e57ea921d8 Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang@amd.com>
Date: Wed, 13 Mar 2024 18:33:10 +0000
Subject: [PATCH] let more integral_constant->constant, and formating

---
 include/ck_tile/core.hpp                      |  1 +
 .../core/algorithm/coordinate_transform.hpp   | 24 +++++++++----------
 include/ck_tile/core/container/sequence.hpp   |  6 ++---
 include/ck_tile/core/container/tuple.hpp      | 11 +++++++++
 include/ck_tile/core/utility/magic_div.hpp    | 19 ++++-----------
 include/ck_tile/host.hpp                      |  1 +
 include/ck_tile/ops/common.hpp                |  1 +
 include/ck_tile/ops/epilogue.hpp              |  2 ++
 include/ck_tile/ops/fmha.hpp                  |  1 +
 9 files changed, 36 insertions(+), 30 deletions(-)

diff --git a/include/ck_tile/core.hpp b/include/ck_tile/core.hpp
index f1370bdc00..daf5a12d2d 100644
--- a/include/ck_tile/core.hpp
+++ b/include/ck_tile/core.hpp
@@ -55,3 +55,4 @@
 #include "ck_tile/core/utility/to_sequence.hpp"
 #include "ck_tile/core/utility/transpose_vectors.hpp"
 #include "ck_tile/core/utility/type_traits.hpp"
+
diff --git a/include/ck_tile/core/algorithm/coordinate_transform.hpp b/include/ck_tile/core/algorithm/coordinate_transform.hpp
index d0cd88ffbc..ad7054aabd 100644
--- a/include/ck_tile/core/algorithm/coordinate_transform.hpp
+++ b/include/ck_tile/core/algorithm/coordinate_transform.hpp
@@ -1562,25 +1562,25 @@ CK_TILE_HOST_DEVICE constexpr auto
 make_pad_transform(const LowLength& low_length,
                    const LeftPad& left_pad,
                    const RightPad& right_pad,
-                   integral_constant<bool, SkipIsValidCheck> = integral_constant<bool, false>{})
+                   bool_constant<SkipIsValidCheck> = bool_constant<false>{})
 {
     return pad<LowLength, LeftPad, RightPad, SkipIsValidCheck>{low_length, left_pad, right_pad};
 }
 
 template <typename LowLength, typename LeftPadLength, bool SkipIsValidCheck = false>
-CK_TILE_HOST_DEVICE constexpr auto make_left_pad_transform(
-    const LowLength& low_length,
-    const LeftPadLength& left_pad_,
-    integral_constant<bool, SkipIsValidCheck> = integral_constant<bool, false>{})
+CK_TILE_HOST_DEVICE constexpr auto
+make_left_pad_transform(const LowLength& low_length,
+                        const LeftPadLength& left_pad_,
+                        bool_constant<SkipIsValidCheck> = bool_constant<false>{})
 {
     return left_pad<LowLength, LeftPadLength, SkipIsValidCheck>{low_length, left_pad_};
 }
 
 template <typename LowLength, typename RightPadLength, bool SkipIsValidCheck = false>
-CK_TILE_HOST_DEVICE constexpr auto make_right_pad_transform(
-    const LowLength& low_length,
-    const RightPadLength& right_pad_,
-    integral_constant<bool, SkipIsValidCheck> = integral_constant<bool, false>{})
+CK_TILE_HOST_DEVICE constexpr auto
+make_right_pad_transform(const LowLength& low_length,
+                         const RightPadLength& right_pad_,
+                         bool_constant<SkipIsValidCheck> = bool_constant<false>{})
 {
     return right_pad<LowLength, RightPadLength, SkipIsValidCheck>{low_length, right_pad_};
 }
@@ -1615,9 +1615,9 @@ CK_TILE_HOST_DEVICE constexpr auto make_merge_transform(const LowLengths& low_le
 }
 
 template <typename UpLengths, bool Use24BitIntegerCalculation = false>
-CK_TILE_HOST_DEVICE constexpr auto make_unmerge_transform(
-    const UpLengths& up_lengths,
-    integral_constant<bool, Use24BitIntegerCalculation> = integral_constant<bool, false>{})
+CK_TILE_HOST_DEVICE constexpr auto
+make_unmerge_transform(const UpLengths& up_lengths,
+                       bool_constant<Use24BitIntegerCalculation> = bool_constant<false>{})
 {
     return unmerge<UpLengths, Use24BitIntegerCalculation>{up_lengths};
 }
diff --git a/include/ck_tile/core/container/sequence.hpp b/include/ck_tile/core/container/sequence.hpp
index e10ef40111..acf187cfc8 100644
--- a/include/ck_tile/core/container/sequence.hpp
+++ b/include/ck_tile/core/container/sequence.hpp
@@ -60,7 +60,7 @@ struct sequence
     CK_TILE_HOST_DEVICE static constexpr auto get()
     {
         static_assert(I < size(), "wrong! I too large");
-        return number<impl::at_index_t<I, integral_constant<value_type, Is>...>{}>{};
+        return number<impl::at_index_t<I, constant<Is>...>{}>{};
     }
 
     template <index_t I>
@@ -81,7 +81,7 @@ struct sequence
     CK_TILE_HOST_DEVICE static constexpr auto at()
     {
         static_assert(I < size(), "wrong! I too large");
-        return number<impl::at_index_t<I, integral_constant<value_type, Is>...>{}>{};
+        return number<impl::at_index_t<I, constant<Is>...>{}>{};
     }
 
     template <index_t I>
@@ -384,7 +384,7 @@ template <index_t... Ids, index_t... Ns>
 struct seq_reverse<sequence<Ids...>, Ns...>
 {
     template <index_t I>
-    using element = impl::at_index_t<I, integral_constant<index_t, Ns>...>;
+    using element = impl::at_index_t<I, constant<Ns>...>;
     using type    = sequence<element<(sizeof...(Ns) - 1 - Ids)>::value...>;
 };
 } // namespace impl
diff --git a/include/ck_tile/core/container/tuple.hpp b/include/ck_tile/core/container/tuple.hpp
index 1be5a55dce..6692012d0d 100644
--- a/include/ck_tile/core/container/tuple.hpp
+++ b/include/ck_tile/core/container/tuple.hpp
@@ -274,6 +274,17 @@ struct tuple : impl::tuple_base<make_index_sequence<sizeof...(T)>, T...>
 #undef TP_COM_
 };
 
+template <typename>
+struct vector_traits;
+
+// specialization for array
+template <typename... T>
+struct vector_traits<tuple<T...>>
+{
+    using scalar_type                    = __type_pack_element<0, T...>;
+    static constexpr index_t vector_size = sizeof...(T);
+};
+
 // template <class... T>
 // CK_TILE_HOST_DEVICE constexpr
 // tuple<T...>
diff --git a/include/ck_tile/core/utility/magic_div.hpp b/include/ck_tile/core/utility/magic_div.hpp
index fba72865c4..09038ba296 100644
--- a/include/ck_tile/core/utility/magic_div.hpp
+++ b/include/ck_tile/core/utility/magic_div.hpp
@@ -52,8 +52,7 @@ struct magic_division32_bit_range
         constexpr uint32_t multiplier = tmp[number<0>{}];
         constexpr uint32_t shift      = tmp[number<1>{}];
 
-        return make_tuple(integral_constant<uint32_t, multiplier>{},
-                          integral_constant<uint32_t, shift>{});
+        return make_tuple(constant<multiplier>{}, constant<shift>{});
     }
 
     // magic division for uint32_t
@@ -116,25 +115,15 @@ struct magic_division16_bit_range
     }
 
     // integral_constant<uint32_t, .>
-    template <uint32_t Divisor>
-    CK_TILE_HOST_DEVICE static constexpr auto
-        calculate_magic_numbers(integral_constant<uint32_t, Divisor>)
+    template <auto Divisor>
+    CK_TILE_HOST_DEVICE static constexpr auto calculate_magic_numbers(constant<Divisor>)
     {
         constexpr auto tmp = calculate_magic_numbers(uint32_t{Divisor});
 
         constexpr uint32_t multiplier = tmp[number<0>{}];
         constexpr uint32_t shift      = tmp[number<1>{}];
 
-        return make_tuple(integral_constant<uint32_t, multiplier>{},
-                          integral_constant<uint32_t, shift>{});
-    }
-
-    // integral_constant<int32_t, .>
-    template <int32_t Divisor>
-    CK_TILE_HOST_DEVICE static constexpr auto
-        calculate_magic_numbers(integral_constant<int32_t, Divisor>)
-    {
-        return calculate_magic_numbers(integral_constant<uint32_t, Divisor>{});
+        return make_tuple(constant<multiplier>{}, constant<shift>{});
     }
 
     // magic division for uint32_t
diff --git a/include/ck_tile/host.hpp b/include/ck_tile/host.hpp
index 0c4a778226..1bbb4b9539 100644
--- a/include/ck_tile/host.hpp
+++ b/include/ck_tile/host.hpp
@@ -20,3 +20,4 @@
 #include "ck_tile/host/reference/reference_reduce.hpp"
 #include "ck_tile/host/reference/reference_softmax.hpp"
 #include "ck_tile/host/stream_config.hpp"
+
diff --git a/include/ck_tile/ops/common.hpp b/include/ck_tile/ops/common.hpp
index 4363ea1f55..9fc1c0d0c1 100644
--- a/include/ck_tile/ops/common.hpp
+++ b/include/ck_tile/ops/common.hpp
@@ -4,3 +4,4 @@
 #pragma once
 
 #include "ck_tile/ops/common/tensor_layout.hpp"
+
diff --git a/include/ck_tile/ops/epilogue.hpp b/include/ck_tile/ops/epilogue.hpp
index 8e0d3e449b..ab399dbf7a 100644
--- a/include/ck_tile/ops/epilogue.hpp
+++ b/include/ck_tile/ops/epilogue.hpp
@@ -4,3 +4,5 @@
 #pragma once
 
 #include "ck_tile/ops/epilogue/default_2d_epilogue.hpp"
+#include "ck_tile/ops/common/tensor_layout.hpp"
+
diff --git a/include/ck_tile/ops/fmha.hpp b/include/ck_tile/ops/fmha.hpp
index 1e9acc6d7b..f886d470d5 100644
--- a/include/ck_tile/ops/fmha.hpp
+++ b/include/ck_tile/ops/fmha.hpp
@@ -18,3 +18,4 @@
 #include "ck_tile/ops/fmha/pipeline/tile_fmha_shape.hpp"
 #include "ck_tile/ops/fmha/pipeline/tile_fmha_traits.hpp"
 #include "ck_tile/ops/common/tensor_layout.hpp"
+