From a8f1d44078e13a90df052b4e5af96f1946b6bad3 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <>
Date: Tue, 13 Jan 2026 04:17:06 -0500
Subject: [PATCH] Make BlockTransferDescriptor concept parametrized. Introduce
 a common TileTransferParameters concept for conv algorithms.

---
 .../builder/conv_algorithm_concepts.hpp       | 21 +++-----
 .../builder/factory/conv_algorithms.hpp       | 49 ++++++++++---------
 .../test/impl/conv_algorithm_types.hpp        |  4 +-
 3 files changed, 37 insertions(+), 37 deletions(-)

diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp
index 2b0f63296b..791924ccd4 100644
--- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp
+++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp
@@ -56,7 +56,7 @@ concept GridwiseWmmaGemmDescriptor = requires(T t) {
 
 // Concept for vectorized data transfer for convolution input tensors.
 template <typename T>
-concept BlockTransferDescriptor = requires(T t) {
+concept BlockTransferDescriptor3D = requires(T t) {
     { t.k0 } -> SizeType;
     { t.m_n } -> SizeType;
     { t.k1 } -> SizeType;
@@ -70,6 +70,10 @@ concept BlockTransferDescriptor4D = requires(T t) {
     { t.k_batch_size } -> SizeType;
 };
 
+template <typename T, size_t ThreadClusterRank>
+concept BlockTransferDescriptor = (ThreadClusterRank == 3 && BlockTransferDescriptor3D<T>) ||
+                                  (ThreadClusterRank == 4 && BlockTransferDescriptor4D<T>);
+
 // Concept for thread cluster dimensions for GEMM output tensor.
 template <typename T>
 concept ThreadClusterDescriptor = requires(T t) {
@@ -202,19 +206,10 @@ concept SpecifiesGridwiseWmmaGemm = requires(T t) {
 };
 
 // Concept to check if a struct specifies convolution input and output block transfer info.
-template <typename T>
+template <typename T, size_t ThreadClusterRank = 3>
 concept SpecifiesBlockTransfer = requires(T t) {
-    { T::transfer.a.block_transfer } -> BlockTransferDescriptor;
-    { T::transfer.b.block_transfer } -> BlockTransferDescriptor;
-    { T::transfer.c.thread_cluster_dims } -> ThreadClusterDescriptor;
-};
-
-// Concept to check if a struct specifies convolution input and output block transfer info
-// for 4D thread slices.
-template <typename T>
-concept SpecifiesBlockTransfer4D = requires(T t) {
-    { T::transfer.a.block_transfer } -> BlockTransferDescriptor4D;
-    { T::transfer.b.block_transfer } -> BlockTransferDescriptor4D;
+    { T::transfer.a.block_transfer } -> BlockTransferDescriptor<ThreadClusterRank>;
+    { T::transfer.b.block_transfer } -> BlockTransferDescriptor<ThreadClusterRank>;
     { T::transfer.c.thread_cluster_dims } -> ThreadClusterDescriptor;
 };
 
diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp
index ffd45efe49..93979440ca 100644
--- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp
+++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp
@@ -8,39 +8,46 @@
 namespace ck_tile::builder::factory {
 
 // Base algorithm concepts
+template <typename T, size_t ThreadClusterRank = 3>
+concept TileTransferParameters = 
+    SpecifiesBlockTransfer<T, ThreadClusterRank> && SpecifiesLdsTransfer<T> &&
+    SpecifiesThreadClusterAccessOrder<T> && SpecifiesSourceAccessOrder<T>;
+
+template <typename T>
+concept SpecifiesTileTransferParameters3D = TileTransferParameters<T, 3>;
+
+template <typename T>
+concept SpecifiesTileTransferParameters4D = TileTransferParameters<T, 4>;
+
+
 template <typename T>
 concept FwdXdlAlgorithmBase =
-    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesBlockTransfer<T> &&
-    SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
-    SpecifiesSourceAccessOrder<T> && SpecifiesGridwiseFwdXdlGemm<T> &&
-    SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> &&
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters3D<T> && 
+    SpecifiesGridwiseFwdXdlGemm<T> && SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> &&
     SpecifiesNumPrefetchStages<T> && SpecifiesNumGroupsToMerge<T> && SpecifiesLoopScheduler<T>;
 
 template <typename T>
 concept BwdXdlAlgorithmBase =
-    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesBlockTransfer4D<T> &&
-    SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
-    SpecifiesSourceAccessOrder<T> && SpecifiesGridwiseBwdXdlGemm<T> &&
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && 
+    SpecifiesTileTransferParameters4D<T> && SpecifiesGridwiseBwdXdlGemm<T> &&
     SpecifiesBwdWeightConvSpecialization<T>;
 
 template <typename T>
 concept BwdXdlV3AlgorithmBase =
-    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesBlockTransfer<T> &&
-    SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
-    SpecifiesSourceAccessOrder<T> && SpecifiesGridwiseBwdXdlGemm<T> &&
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && 
+    SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseBwdXdlGemm<T> &&
     SpecifiesBwdWeightConvSpecialization<T> && SpecifiesBlockGemm<T>;
+
 template <typename T>
 concept BwdWmmaAlgorithmBase =
-    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesBlockTransfer<T> &&
-    SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
-    SpecifiesSourceAccessOrder<T> && SpecifiesGridwiseWmmaGemm<T> &&
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && 
+    SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseWmmaGemm<T> &&
     SpecifiesBwdWeightConvSpecialization<T>;
 
 template <typename T>
 concept BwdWmmaV3AlgorithmBase =
-    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesBlockTransfer<T> &&
-    SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
-    SpecifiesSourceAccessOrder<T> && SpecifiesGridwiseWmmaGemm<T> &&
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && 
+    SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseWmmaGemm<T> &&
     SpecifiesBwdWeightConvSpecialization<T> && SpecifiesBlockGemm<T>;
 
 // Reference algorithm concept
@@ -62,17 +69,15 @@ concept LargeTensorAlgorithm = FwdXdlAlgorithmBase<T> && SpecifiesLargeTensorSup
 
 template <typename T>
 concept FwdXdlV3Algorithm =
-    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesBlockTransfer<T> &&
-    SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
-    SpecifiesSourceAccessOrder<T> && SpecifiesGridwiseFwdXdlGemm<T> &&
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && 
+    SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseFwdXdlGemm<T> &&
     SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> && SpecifiesBlockGemm<T>;
 
 // FWD WMMA algorithm concepts
 template <typename T>
 concept FwdWmmaAlgorithm =
-    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesBlockTransfer<T> &&
-    SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
-    SpecifiesSourceAccessOrder<T> && SpecifiesGridwiseWmmaGemm<T> &&
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && 
+    SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseWmmaGemm<T> &&
     SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> &&
     SpecifiesNumPrefetchStages<T> && SpecifiesLoopScheduler<T> && SpecifiesGridwiseGemmPipeline<T>;
 
diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp
index 27ba1ec3b6..617686fda1 100644
--- a/experimental/builder/test/impl/conv_algorithm_types.hpp
+++ b/experimental/builder/test/impl/conv_algorithm_types.hpp
@@ -90,8 +90,8 @@ struct BlockTransfer<3>
     size_t m_n;
     size_t k1;
 };
-static_assert(ckb::BlockTransferDescriptor<BlockTransfer<>>);
-static_assert(ckb::BlockTransferDescriptor<BlockTransfer<4>>);
+static_assert(ckb::BlockTransferDescriptor<BlockTransfer<3>, 3>);
+static_assert(ckb::BlockTransferDescriptor<BlockTransfer<4>, 4>);
 
 // Describe C block transfer thread cluster lengths.
 struct ThreadCluster