From c5cdd51ce4cf5749ca158d466cf3647283d7c4be Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <>
Date: Wed, 7 Jan 2026 05:41:22 -0500
Subject: [PATCH] Fix factory for regular WMMA conv bwd weight.

---
 .../builder/factory/conv_bwd_weight_wmma_factory.hpp      | 8 ++++----
 .../conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp    | 2 +-
 experimental/builder/test/impl/conv_algorithm_types.hpp   | 2 +-
 .../builder/test/utils/conv_algorithm_type_utils.hpp      | 2 +-
 4 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp
index 817432081b..32161a234a 100644
--- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp
+++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp
@@ -51,13 +51,13 @@ struct ConvBwdWeightWmmaFactory
     static_assert(InputVectorTransferLimits<A_BLOCK_TRANSFER>, "Invalid A block transfer config");
     static_assert(InputVectorTransferLimits<B_BLOCK_TRANSFER>, "Invalid B block transfer config");
     static_assert(OutputVectorTransferLimits<C_BLOCK_TRANSFER>, "Invalid C block transfer config");
-    static_assert(AccessOrderLimits4D<A_BLOCK_TRANSFER.thread_cluster_order>,
+    static_assert(AccessOrderLimits3D<A_BLOCK_TRANSFER.thread_cluster_order>,
                   "Invalid A thread cluster access order");
-    static_assert(AccessOrderLimits4D<B_BLOCK_TRANSFER.thread_cluster_order>,
+    static_assert(AccessOrderLimits3D<B_BLOCK_TRANSFER.thread_cluster_order>,
                   "Invalid B thread cluster access order");
-    static_assert(AccessOrderLimits4D<A_BLOCK_TRANSFER.src_access_order>,
+    static_assert(AccessOrderLimits3D<A_BLOCK_TRANSFER.src_access_order>,
                   "Invalid A source access order");
-    static_assert(AccessOrderLimits4D<B_BLOCK_TRANSFER.src_access_order>,
+    static_assert(AccessOrderLimits3D<B_BLOCK_TRANSFER.src_access_order>,
                   "Invalid B source access order");
 
     // The forward convolution kernel class instance.
diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp
index 0981ea6c11..ff350ac804 100644
--- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp
+++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp
@@ -22,7 +22,7 @@ constexpr auto SIGNATURE = ckt::ConvSignature{.spatial_dim = 3,
 constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle{}
                                .with_thread_block(cku::ThreadBlock_64_32x32x32)
                                .with_gemm_config(cku::GemmParams_Wmma_16x16_2x1_per_wave)
-                               .with_transfer(cku::BwdTransfer_4x64x1)
+                               .with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3)
                                .with_bwd_specialization(ckb::ConvSpecialization::DEFAULT)
                                .with_prefetch_config(1, ckb::PipelineScheduler::DEFAULT)
                                .with_gridwise_gemm_pipeline(ckb::PipelineVersion::V1);
diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp
index 797f440bdc..27ba1ec3b6 100644
--- a/experimental/builder/test/impl/conv_algorithm_types.hpp
+++ b/experimental/builder/test/impl/conv_algorithm_types.hpp
@@ -669,7 +669,7 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Wmma_CShuffle_V3 =
 using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle =
     ConvAlgorithmTemplate<ThreadBlock_,
                           WmmaGemm_,
-                          Transfer_<4>,
+                          Transfer_<>,
                           ConvSpecializationBwdWeight_,
                           GridGemm_,
                           Prefetch_>;
diff --git a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp
index b705ab1e47..d80d6a1b8c 100644
--- a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp
+++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp
@@ -428,7 +428,7 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuf
 {
     std::ostringstream oss;
     oss << to_string(static_cast<ThreadBlock_>(t)) << "," << to_string(static_cast<WmmaGemm_>(t))
-        << "," << to_string(static_cast<Transfer_<4>>(t));
+        << "," << to_string(static_cast<Transfer_<>>(t));
     return oss.str();
 }