Template recursion-free static_ford #2011

tenpercent · 2025-03-25T03:53:07Z

Proposed changes

The current implementation of static_ford uses recursion over the sequence of dimensions and instantiates lambdas at each internal step. This can be avoided and should improve compilation time. The current most time consuming template is static_for(0,1) from static_ford instantiation internals for both old CK and CK-tile

Checklist

Please put an x into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.

I have added tests relevant to the introduced functionality, and the unit tests are passing locally
I have added the test to REGRESSION_TESTS list defined at the top of CMakeLists.txt in tests/CMakeLists.txt, IF the test takes more than 30 seconds to run.
I have added inline documentation which enables the maintainers with understanding the motivation
I have removed the stale documentation which is no longer relevant after this pull request
(If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request
I have run clang-format on all changed files
Any dependent changes have been merged

Discussion

Snapshot -

ck::static_ford<ck::Sequence<14, 14, 16>>{} ([] (auto I) {
        I.Print();
    });

*** Old
 10703 ms: ck::static_for<0, 14, 1>::operator()<(lambda at ../include/ck/utility/functional3.hpp:31:52)> (15 times, avg 713 ms)
 10701 ms: ck::detail::static_for_impl<ck::Sequence<0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13>>::operator()<(lambda at ../include/ck/utility/functional3.hpp:31:52)> (15 times, avg 713 ms)
  5364 ms: ck::static_ford<ck::Sequence<14, 14, 16>>::operator()<(lambda at main.cpp:114:50)> (1 times, avg 5364 ms)
  5358 ms: ck::detail::static_ford_impl<const ck::Sequence<14, 14, 16>, ck::Sequence<0, 1, 2>>::operator()<(lambda at main.cpp:114:50), ck::Sequence<>> (1 times, avg 5358 ms)
  5210 ms: ck::static_for<0, 16, 1>::operator()<(lambda at ../include/ck/utility/functional3.hpp:31:52)> (196 times, avg 26 ms)
  5184 ms: ck::detail::static_for_impl<ck::Sequence<0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15>>::operator()<(lambda at ../include/ck/utility/functional3.hpp:31:52)> (196 times, avg 26 ms)
  1844 ms: ck::static_for<0, 3, 1>::operator()<(lambda at ../include/ck/utility/sequence.hpp:182:36)> (3136 times, avg 0 ms)
   401 ms: ck::detail::static_ford_impl<const ck::Sequence<14, 14, 16>, ck::Sequence<0, 1, 2>>::operator()((lambda at main.cpp:114:50), ck::Sequence<>)::(anonymous class)::operator()<ck::integral_constant<int, 7>> (1 times, avg 401 ms)
   400 ms: ck::detail::static_ford_impl<ck::Sequence<14, 16>, ck::Sequence<0, 1, 2>>::operator()<(lambda at main.cpp:114:50), ck::Sequence<7>> (1 times, avg 400 ms)
   399 ms: ck::detail::static_ford_impl<const ck::Sequence<14, 14, 16>, ck::Sequence<0, 1, 2>>::operator()((lambda at main.cpp:114:50), ck::Sequence<>)::(anonymous class)::operator()<ck::integral_constant<int, 11>> (1 times, avg 399 ms)
   398 ms: ck::detail::static_ford_impl<ck::Sequence<14, 16>, ck::Sequence<0, 1, 2>>::operator()<(lambda at main.cpp:114:50), ck::Sequence<11>> (1 times, avg 398 ms)
...
*** New
  2978 ms: ck::static_for<0, 3, 1>::operator()<(lambda at ../include/ck/utility/sequence.hpp:182:36)> (3136 times, avg 0 ms)
  2639 ms: ck::detail::static_for_impl<ck::Sequence<0, 1, 2>>::operator()<(lambda at ../include/ck/utility/sequence.hpp:182:36)> (3136 times, avg 0 ms)
  1457 ms: ck::static_ford<ck::Sequence<14, 14, 16>>::operator()<(lambda at main.cpp:114:50), ck::static_ford<ck::Sequence<14, 14, 16>>::convert_t> (1 times, avg 1457 ms)
  1457 ms: ck::detail::applier<int, 0, 1, 2, 3, 4, 5, 6, 7... (1 times, avg 1457 ms)
    20 ms: ck::Sequence<9, 12, 13>::Print (1 times, avg 20 ms)
...

…ePlatform/composable_kernel into special-static-for-d

tenpercent · 2025-03-25T19:55:18Z

include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v4r1.hpp

@@ -4,6 +4,7 @@
 #pragma once

 #include "ck/utility/common_header.hpp"
+#include "ck/utility/functional3.hpp"


static_ford is used but not included explicitly

tenpercent · 2025-03-25T20:16:42Z

CMakeLists.txt

@@ -547,6 +547,10 @@ if("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU" AND CMAKE_CXX_COMPILER_VERSION VERS
    add_compile_options(-fdiagnostics-color=always)
 endif()

+# fold expression depth
+# device_grouped_conv2d_fwd_xdl_ngchw_gkyxc_ngkhw_f16_comp_instance.cpp
+add_compile_options(-fbracket-depth=3136)


alternative is to implement some sort of paging for the parameter packs, not sure how to do it yet

…ePlatform/composable_kernel into special-static-for-d

tenpercent · 2025-03-26T19:37:13Z

include/ck/utility/functional3.hpp

    __host__ __device__ constexpr void operator()(F f) const
    {
-        constexpr auto ordered_lengths = Lengths::ReorderGivenNew2Old(Orders{});
-        detail::static_ford_impl<decltype(ordered_lengths), Orders>{}(f, Sequence<>{});
+        base::template operator()<F, IndexTransform>(f);


possible improvement here is to figure out how to bake index transform into the base class, to do using base::operator() and avoid one extra template instantiation

tenpercent · 2025-03-26T19:40:29Z

include/ck/utility/functional3.hpp

+    using TCumProd             = typename make_cumulative_product<Dims...>::type;
+
+    template <ck::index_t flat_idx>
+    using type = decltype((TCumProd{} * ck::Number<flat_idx>{} / ck::Number<Prod>{}) % SDim{});


my initial expectation was this would be fairly cheap but it is not

tenpercent · 2025-03-26T20:58:42Z

include/ck/utility/functional3.hpp

+
+// clang-format off
+template <int32_t... Idims>
+struct make_cumulative_product


recursive template definition seems doable but complex

tenpercent · 2025-03-27T05:40:20Z

include/ck/utility/functional3.hpp

+// clang-format off
+
+template <int32_t IDim0>
+constexpr auto make_cumulative_product(ck::Number<IDim0>)


functions are faster than structs

…template parameter

tenpercent added 10 commits March 25, 2025 00:15

try static-ford replacement

ab09203

try static-ford replacement

bd07e1d

Merge branch 'special-static-for-d' of https://github.com/ROCmSoftwar…

b0e1c86

…ePlatform/composable_kernel into special-static-for-d

handle const sequences and up to 5d

a399463

tweak clang fold expression depth limit

9ffe001

tweak more

ed21968

add 6d for moe

c06d68b

clang-format

f998ca1

pretty template towers

68b33cb

add safeguard for pack constraints

f1c6ab9

tenpercent added the compilation time label Mar 25, 2025

tenpercent changed the title ~~Special static for d~~ Recursion-free static_fordd Mar 25, 2025

tenpercent changed the title ~~Recursion-free static_fordd~~ Template recursion-free static_ford Mar 25, 2025

Merge branch 'develop' into special-static-for-d

a11cc46

tenpercent commented Mar 25, 2025

View reviewed changes

tenpercent added 11 commits March 25, 2025 17:36

compact reuse by const overload

8b25656

Merge branch 'special-static-for-d' of https://github.com/ROCmSoftwar…

a065251

…ePlatform/composable_kernel into special-static-for-d

generalize from ck::sequence

b451ddd

typo

759014c

remove redundant lambda

deae42b

factor out converter type template

2102118

Merge branch 'develop' into special-static-for-d

53815b5

rename template parameter

e6ee84d

Merge branch 'special-static-for-d' of https://github.com/ROCmSoftwar…

86a50c6

…ePlatform/composable_kernel into special-static-for-d

Merge branch 'develop' into special-static-for-d

3ba589f

Merge branch 'special-static-for-d' of https://github.com/ROCmSoftwar…

71f9247

…ePlatform/composable_kernel into special-static-for-d

tenpercent commented Mar 26, 2025

View reviewed changes

factor out applier

ce62a85

tenpercent commented Mar 26, 2025

View reviewed changes

tenpercent marked this pull request as ready for review March 26, 2025 21:24

tenpercent requested review from illsilin, carlushuang, qianfengz, aosewski, poyenc, geyyer, bartekxk, andriy-ca, afagaj and asleepzzz as code owners March 26, 2025 21:24

replace helper struct with function

1ef9738

tenpercent commented Mar 27, 2025

View reviewed changes

tenpercent added 5 commits March 27, 2025 02:16

replace helper struct with function

e15bb4d

add comment

bab0eab

clang-format

aeb2f6b

Merge branch 'develop' into special-static-for-d

b56d6d9

rename applier duplicate to make difference for having the transform …

3cbbf28

…template parameter

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Template recursion-free static_ford #2011

Template recursion-free static_ford #2011

tenpercent commented Mar 25, 2025 •

edited

Loading

tenpercent Mar 25, 2025

tenpercent Mar 25, 2025

tenpercent Mar 26, 2025

tenpercent Mar 26, 2025

tenpercent Mar 26, 2025

tenpercent Mar 27, 2025

Template recursion-free static_ford #2011

Are you sure you want to change the base?

Template recursion-free static_ford #2011

Conversation

tenpercent commented Mar 25, 2025 • edited Loading

Proposed changes

Checklist

Discussion

tenpercent Mar 25, 2025

Choose a reason for hiding this comment

tenpercent Mar 25, 2025

Choose a reason for hiding this comment

tenpercent Mar 26, 2025

Choose a reason for hiding this comment

tenpercent Mar 26, 2025

Choose a reason for hiding this comment

tenpercent Mar 26, 2025

Choose a reason for hiding this comment

tenpercent Mar 27, 2025

Choose a reason for hiding this comment

tenpercent commented Mar 25, 2025 •

edited

Loading