Skip to content

Conversation

@Snektron
Copy link
Contributor

@Snektron Snektron commented Jan 20, 2026

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::run integration for backwards weight algorithms, for the reference kernel, ck, and ck tile.

Design

  • Since the 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 contains Args<>, the latter contains Inputs, Outputs, and direction-specific functionality.
  • I've moved the headers around a bit to group everything in the testing/conv/ directory, I figured that'd make more sense than having conv/ in all names.
  • Since CK requires separate ckt::run implementations for fwd/bwd wt/bwd data, which are relatively large, I put them in separate files fwd_ck.hpp and bwd_weight_ck.hpp.
  • The reference and CK tile implementations are identical and easy to abstract, so both fwd and bwd weight are both in reference.hpp and ck_tile.hpp
  • The return type of ckt::run was changed to a tuple in [CK_BUILDER] Add grouped conv fwd ck tile profiler #3518. We discussed that throwing exceptions in ckt::run is not desirable, but I don't really like the tuple either. I've written a custom RunResult struct instead, and testing integration so that we can easily check if a run was correctly, for example by doing EXPECT_THAT(ckt::run(...), SuccessfulRun()). If theres a problem then a message is printed.

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...
Copy link
Contributor

Copilot AI left a 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 RunResult structure to replace tuple return types from ckt::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.

Using EXPECT_THAT(..., SuccessfulRun()) will generate a check and a nice error
message about how and why running an algorithm failed.
@Snektron Snektron force-pushed the rvoetter/conv-testing-bwd-wt branch from 10fb64c to 86058ae Compare January 21, 2026 09:45
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants