-
Notifications
You must be signed in to change notification settings - Fork 263
[CK_BUILDER] Convert convolution traits to a struct with factory functions #3547
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: develop
Are you sure you want to change the base?
Conversation
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Pull request overview
This PR refactors the convolution traits system from a templated class-based design to a struct-based design with explicit factory functions. The previous design used template class instances with concepts for unification, which did not scale well. The new approach uses a non-templated ConvTraits struct populated by device-specific factory functions, making the conversion from instance traits to ConvTraits explicit and maintainable.
Changes:
- Converted
ConvTraitsfrom a templated class to a pure data struct - Added factory functions (
instance_to_conv_traits) for each device kernel variant - Split helper utilities into separate files for better organization
- Updated all tests to use the new factory function approach
Reviewed changes
Copilot reviewed 17 out of 17 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
conv_traits.hpp |
Converted from templated class to pure data struct |
conv_traits_helpers.hpp |
New file containing helper functions for trait extraction |
instance_to_conv_traits.hpp |
New umbrella header for all factory functions |
conv_traits_device_grouped_conv_fwd_multiple_*.hpp |
Factory functions for each device kernel variant |
instance_traits_device_grouped_conv_fwd_multiple_*.hpp |
Added device kernel tag types |
unit_instance_to_conv_traits_features.cpp |
Unit tests for individual helper functions |
unit_instance_to_conv_traits_instances.cpp |
Unit tests for factory functions per device kernel |
test_conv_traits.cpp |
Updated to use new factory function API |
| Other files | Updated references and imports |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| /// members. | ||
| /// @details This struct can hold the data from any ConvTraitsImpl class, allowing runtime storage | ||
| /// and manipulation of convolution configuration information. | ||
| struct ConvTraits |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Is there a reason to not directly construct ConvDescription in instance_to_conv_traits? It seems like this structure and ConvSignatureInfo and ConvAlgorithmInfo hold the same exact fields. Except for the instance_string function, but that should be trivial to add there too.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I've added documentation that should explain the separation of responsibilites. Additionally, note that the include dependencies are simpler with this separation.
...de/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/reflect/conv_traits_helpers.hpp
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/reflect/conv_traits_helpers.hpp
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/reflect/conv_traits_helpers.hpp
Show resolved
Hide resolved
aosewski
left a comment
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I really like those changes! Just have few ideas and future potential problems. See comment
...de/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
Show resolved
Hide resolved
This is a lot of cleanup on tests to have verbose coverage of feature extraction, explicit tests for each supported device kernel, and simple, readable test code.
526938c to
815cb5f
Compare
The convolution traits for forward XDL kernels were complicated and the design wasn't going to scale.
The previous design had separate convolution traits template class instances for each device kernel instance, and used concepts to try to unify the instantiations.
This PR switches to a non-templated composite struct for ConvTraits, with a separate factory function for each device kernel instance. That way the conversion from instance traits and ConvTraits is explicit and can be handle any special cases an instance might have.
We add explicit unit tests for the helper function and also add a unit test for each factory function.