From 26724086f3a60e809377f6abfb17280c67f2bb1b Mon Sep 17 00:00:00 2001
From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
Date: Tue, 8 Apr 2025 09:00:51 -0700
Subject: [PATCH] simplify generate_tuple (#2043)

[ROCm/composable_kernel commit: 6ce0797dadfc6d0c6cdde3e01532e90137fc5b0c]
---
 include/ck/utility/sequence.hpp          | 15 +++++++++++++++
 include/ck/utility/tuple_helper.hpp      |  9 +++++++--
 include/ck_tile/core/container/tuple.hpp |  9 +++++++--
 3 files changed, 29 insertions(+), 4 deletions(-)

diff --git a/include/ck/utility/sequence.hpp b/include/ck/utility/sequence.hpp
index 99935a6d8d..497625f7e2 100644
--- a/include/ck/utility/sequence.hpp
+++ b/include/ck/utility/sequence.hpp
@@ -184,6 +184,21 @@ struct Sequence
     }
 };
 
+namespace impl {
+template <typename T, T... Ints>
+struct __integer_sequence;
+
+template <index_t... Ints>
+struct __integer_sequence<index_t, Ints...>
+{
+    using seq_type = Sequence<Ints...>;
+};
+} // namespace impl
+
+template <index_t N>
+using make_index_sequence =
+    typename __make_integer_seq<impl::__integer_sequence, index_t, N>::seq_type;
+
 // merge sequence
 template <typename Seq, typename... Seqs>
 struct sequence_merge
diff --git a/include/ck/utility/tuple_helper.hpp b/include/ck/utility/tuple_helper.hpp
index b4f1545aa9..b1a0c1fc5d 100644
--- a/include/ck/utility/tuple_helper.hpp
+++ b/include/ck/utility/tuple_helper.hpp
@@ -11,11 +11,16 @@
 
 namespace ck {
 
+template <typename F, index_t... ids>
+__host__ __device__ constexpr auto generate_tuple_for(F&& f, Sequence<ids...>)
+{
+    return make_tuple(f(Number<ids>{})...);
+}
+
 template <typename F, index_t N>
 __host__ __device__ constexpr auto generate_tuple(F&& f, Number<N>)
 {
-    return unpack([&f](auto&&... xs) { return make_tuple(f(xs)...); },
-                  typename arithmetic_sequence_gen<0, N, 1>::type{});
+    return generate_tuple_for(f, make_index_sequence<N>{});
 }
 
 template <typename F, index_t N>
diff --git a/include/ck_tile/core/container/tuple.hpp b/include/ck_tile/core/container/tuple.hpp
index fd02177e25..3700d348e7 100644
--- a/include/ck_tile/core/container/tuple.hpp
+++ b/include/ck_tile/core/container/tuple.hpp
@@ -396,11 +396,16 @@ struct tuple_array_impl<T, 1>
 };
 } // namespace impl
 
+template <typename F, index_t... ids>
+CK_TILE_HOST_DEVICE constexpr auto generate_tuple_for(F&& f, sequence<ids...>)
+{
+    return make_tuple(f(number<ids>{})...);
+}
+
 template <typename F, index_t N>
 CK_TILE_HOST_DEVICE constexpr auto generate_tuple(F&& f, number<N>)
 {
-    return unpack([&f](auto&&... is) { return make_tuple(f(is)...); },
-                  typename arithmetic_sequence_gen<0, N, 1>::type{});
+    return generate_tuple_for(f, make_index_sequence<N>{});
 }
 
 template <typename F, index_t N>