From cf08db04a6a6eb2761040cf32335e4c1eeb0b33c Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Tue, 1 Apr 2025 12:06:25 -0700 Subject: [PATCH] add a fast compilation path for static for (0..N) (#2005) * add a fast compilation path for static for (0..N) * Update functional2.hpp add comment and put range applier into detail namespace * Update functional.hpp ditto for ck-tile * prettify * prettify more * add comment * clang-format [ROCm/composable_kernel commit: c59a8bb206d4dc763d07e16f730e563849e68cb6] --- include/ck/utility/functional2.hpp | 24 +++++++++++++++++++++ include/ck_tile/core/utility/functional.hpp | 24 +++++++++++++++++++++ 2 files changed, 48 insertions(+) diff --git a/include/ck/utility/functional2.hpp b/include/ck/utility/functional2.hpp index 99c65f4eb8..a11963cb47 100644 --- a/include/ck/utility/functional2.hpp +++ b/include/ck/utility/functional2.hpp @@ -46,4 +46,28 @@ struct static_for } }; +namespace detail { + +template +struct applier +{ + template + __host__ __device__ constexpr void operator()(F f) const + { + // tweak -fbracket-depth if compilation fails. Clang default limit is 256 + (f(Number{}), ...); + } +}; + +template // == sizeof...(Is) +using make_applier = __make_integer_seq; + +} // namespace detail + +template +struct static_for<0, N, 1> : detail::make_applier +{ + using detail::make_applier::operator(); +}; + } // namespace ck diff --git a/include/ck_tile/core/utility/functional.hpp b/include/ck_tile/core/utility/functional.hpp index 2cdce94063..fd0252d3ca 100644 --- a/include/ck_tile/core/utility/functional.hpp +++ b/include/ck_tile/core/utility/functional.hpp @@ -58,6 +58,30 @@ struct static_for } }; +namespace detail { + +template +struct applier +{ + template + CK_TILE_HOST_DEVICE constexpr void operator()(F f) const + { + // tweak -fbracket-depth if compilation fails. Clang default limit is 256 + (f(number{}), ...); + } +}; + +template // == sizeof...(Is) +using make_applier = __make_integer_seq; + +} // namespace detail + +template +struct static_for<0, N, 1> : detail::make_applier +{ + using detail::make_applier::operator(); +}; + struct identity { template