-
Notifications
You must be signed in to change notification settings - Fork 267
[CK_BUILDER] conv bwd weight testing #3618
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: rvoetter/conv-testing-reference-old-ck
Are you sure you want to change the base?
[CK_BUILDER] conv bwd weight testing #3618
Conversation
In order to prepare for bwd of conv testing, this commit moves some files and types around so that we can reuse ckt::Args for both forward and backwards convolution.
This will allow us to more easily include fwd.hpp from backwards definitions, which is required for initializing bwd values.
Turns out that the supplied layout isn't actually supported...
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 adds backwards weight testing integration for convolution operations, extending the existing ckt::run infrastructure to support backwards weight algorithms alongside the existing forward convolution support.
Changes:
- Introduced
RunResultstructure to replace tuple return types fromckt::run, providing clearer success/failure semantics - Reorganized convolution testing headers into a
testing/conv/directory structure - Added backwards weight support for CK, CK Tile, and reference implementations
Reviewed changes
Copilot reviewed 26 out of 27 changed files in this pull request and generated 12 comments.
Show a summary per file
| File | Description |
|---|---|
experimental/builder/include/ck_tile/builder/testing/testing.hpp |
Adds RunResult struct and updates run() signature to return it |
experimental/builder/include/ck_tile/builder/testing/conv/args.hpp |
Extracts reusable Args<> structure, removing direction-specific code |
experimental/builder/include/ck_tile/builder/testing/conv/fwd.hpp |
Defines forward-specific Inputs, Outputs, and initialization |
experimental/builder/include/ck_tile/builder/testing/conv/bwd_weight.hpp |
Defines backwards weight-specific Inputs, Outputs, and initialization |
experimental/builder/include/ck_tile/builder/testing/conv/reference.hpp |
Implements reference kernel support for both forward and backwards weight |
experimental/builder/include/ck_tile/builder/testing/conv/ck_tile.hpp |
Implements CK Tile support for both forward and backwards weight |
experimental/builder/include/ck_tile/builder/testing/conv/fwd_ck.hpp |
Implements old CK forward convolution support |
experimental/builder/include/ck_tile/builder/testing/conv/bwd_weight_ck.hpp |
Implements old CK backwards weight convolution support |
experimental/builder/test/testing_utils.hpp |
Adds SuccessfulRun() matcher and enhances validation error reporting |
experimental/builder/test/testing_utils.cpp |
Implements RunResultMatcher for testing ckt::run results |
| Various test files | Updates to use new header locations and SuccessfulRun() matcher |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
experimental/builder/include/ck_tile/builder/testing/conv/bwd_weight.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/bwd_weight.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/bwd_weight.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/bwd_weight.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/reference.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/fwd_ck.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/ck_tile.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/ck_tile.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/bwd_weight_ck.hpp
Outdated
Show resolved
Hide resolved
experimental/builder/include/ck_tile/builder/testing/conv/bwd_weight_ck.hpp
Outdated
Show resolved
Hide resolved
84ef211 to
10fb64c
Compare
This type is more convenient than std::tuple, as it will allow us to use google test matchers with this in the future.
Using EXPECT_THAT(..., SuccessfulRun()) will generate a check and a nice error message about how and why running an algorithm failed.
10fb64c to
86058ae
Compare
Proposed changes
NOTE: Targetting #3604, merge that one first and update the target branch (GitHub should implement proper stacked pull requests already).
This adds
ckt::runintegration for backwards weight algorithms, for the reference kernel, ck, and ck tile.Design
Args<>structure is reusable between fwd/bwd wt/bwd data conv, I've split conv_fwd.hpp out into args.hpp and fwd.hpp. The former just containsArgs<>, the latter containsInputs,Outputs, and direction-specific functionality.testing/conv/directory, I figured that'd make more sense than havingconv/in all names.ckt::runimplementations for fwd/bwd wt/bwd data, which are relatively large, I put them in separate filesfwd_ck.hppandbwd_weight_ck.hpp.reference.hppandck_tile.hppckt::runwas changed to a tuple in [CK_BUILDER] Add grouped conv fwd ck tile profiler #3518. We discussed that throwing exceptions inckt::runis not desirable, but I don't really like the tuple either. I've written a customRunResultstruct instead, and testing integration so that we can easily check if a run was correctly, for example by doingEXPECT_THAT(ckt::run(...), SuccessfulRun()). If theres a problem then a message is printed.