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 +struct __integer_sequence; + +template +struct __integer_sequence +{ + using seq_type = Sequence; +}; +} // namespace impl + +template +using make_index_sequence = + typename __make_integer_seq::seq_type; + // merge sequence template 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 +__host__ __device__ constexpr auto generate_tuple_for(F&& f, Sequence) +{ + return make_tuple(f(Number{})...); +} + template __host__ __device__ constexpr auto generate_tuple(F&& f, Number) { - 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{}); } template 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 }; } // namespace impl +template +CK_TILE_HOST_DEVICE constexpr auto generate_tuple_for(F&& f, sequence) +{ + return make_tuple(f(number{})...); +} + template CK_TILE_HOST_DEVICE constexpr auto generate_tuple(F&& f, number) { - 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{}); } template