[CK_BUILDER] conv bwd weight testing (#3618)

* ck-builder: restructure testing conv

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.

* ck-builder: decouple fwd_ck.hpp and fwd_reference.hpp from fwd.hpp

This will allow us to more easily include fwd.hpp from backwards
definitions, which is required for initializing bwd values.

* ck-builder: fix layout of test_ckb_conv_bwd_weight_xdl_cshuffle_v3

Turns out that the supplied layout isn't actually supported...

* ck-builder: ck and reference conv integration for bwd weight

* ck-builder: ck bwd weight execution test

* ck-builder: ckt::run support for ck-tile bwd weight

* ck-builder: ck tile bwd weight execution test

* ck-builder: extra debug printing in MatchesReference

* ck-builder: make ckt::run return RunResult

This type is more convenient than std::tuple, as it will allow us to
use google test matchers with this in the future.

* ck-builder: RunResult matcher

Using EXPECT_THAT(..., SuccessfulRun()) will generate a check and a nice error
message about how and why running an algorithm failed.

* ck-builder: doc fixes

* ck-builder: add missing headers

[ROCm/composable_kernel commit: cc75948d1c]

profiler/include/profiler/grouped_convolution_forward_tile_algs.hpp

@@ -9,8 +9,9 @@
 #include "grouped_convolution_signatures.hpp"
 
 #include "ck_tile/builder/testing/filter_extent.hpp"
-#include "ck_tile/builder/testing/conv_fwd_ck_tile.hpp"
-#include "ck_tile/builder/testing/conv_fwd_reference.hpp"
+#include "ck_tile/builder/testing/conv/fwd.hpp"
+#include "ck_tile/builder/testing/conv/ck_tile.hpp"
+#include "ck_tile/builder/testing/conv/reference.hpp"
 #include "ck_tile/builder/conv_builder.hpp"
 
 namespace ck_tile::builder::profiling {
@@ -113,8 +114,8 @@ run_grouped_conv_forward_tile_algs(const ckt::Args<SIGNATURE>& args,
     auto reference = ckt::alloc_outputs(args);
     using ReferenceInstance =
         typename ckb::ConvBuilder<SIGNATURE, ckt::ConvAlgorithm_Reference{}>::Instance;
-    auto ref_conv = ReferenceInstance{};
-    ckt::run(ref_conv, args, inputs, reference.get());
+    auto ref_conv                    = ReferenceInstance{};
+    [[maybe_unused]] auto ref_result = ckt::run(ref_conv, args, inputs, reference.get());
 
     [[maybe_unused]] auto run_alg = [&](auto&& run_alg_func) {
         std::tie(is_supported, avg_time, op_name) = run_alg_func(args, inputs, outputs, s_conf);

profiler/include/profiler/grouped_convolution_signatures.hpp

@@ -6,7 +6,7 @@
 #include <tuple>
 
 #include "../../experimental/builder/test/impl/conv_signature_types.hpp"
-#include "ck_tile/builder/testing/conv_fwd_ck_tile.hpp"
+#include "ck_tile/builder/testing/conv/ck_tile.hpp"
 
 namespace ck_tile::builder::profiling {