Refactor conv algorithms into more categorized form.

2026-07-11 09:40:51 +00:00 · 2026-01-14 08:29:49 -05:00
parent 3f7b250d33
commit 096592eb99
3 changed files with 268 additions and 53 deletions
--- a/experimental/builder/include/ck_tile/builder/factory/README.md
+++ b/experimental/builder/include/ck_tile/builder/factory/README.md
@@ -29,3 +29,231 @@ using Factory = decltype(make_conv_instance<signature, algorithm, "v1">());
 ```

 The dispatcher automatically selects the appropriate factory following explicit logic.
+
+# Convolution Algorithm Hierarchy
+
+This section illustrates the hierarchy of convolution algorithm concepts defined in `conv_algorithms.hpp`.
+
+## Overview
+
+The convolution algorithms are organized into three main categories:
+
+1. **XDL Algorithms** - GPU matrix multiplication using XDL (matrix core instructions)
+2. **WMMA Algorithms** - GPU matrix multiplication using WMMA (Wave Matrix Multiply-Accumulate)
+3. **DL Algorithms** - Special vectorized dot-product kernels optimized for specific data layouts with separate implementation.
+
+XDL and WMMA algorithms share a common base, while DL algorithms have their own independent base.
+
+## Common Base Hierarchy (XDL & WMMA)
+
+Both XDL and WMMA algorithms share the following foundational concepts:
+
+```
+ConvAlgorithm (Base Concept)
+│
+│  Requirements:
+│  • ConvAlgorithmDescriptor
+│  • SpecifiesThreadBlock
+│  • SpecifiesTileTransferParameters (ThreadClusters, LdsTransfer, AccessOrders)
+│  • SpecifiesWarpGemm
+│  • SpecifiesGemmPipeline
+│
+├─── FwdAlgorithm (Forward Convolution)
+│    │
+│    │  Additional: SpecifiesFwdConvSpecialization
+│    │
+│    └─── FwdAlgorithmV3
+│         │
+│         │  Additional: SpecifiesPipelineV3
+│         │
+│
+└─── BwdAlgorithm (Backward Weight Convolution)
+     │
+     │  Additional: SpecifiesBwdWeightConvSpecialization
+     │
+     └─── BwdAlgorithmV3
+          │
+          │  Additional: SpecifiesPipelineV3
+          │
+```
+
+---
+
+## XDL Algorithm Hierarchy
+
+### Forward XDL Algorithms
+
+```
+FwdAlgorithm + SpecifiesXdl
+│
+├─── FwdXdlAlgorithmBase
+     │
+     ├─── FwdXdlAlgorithm
+     │    │
+     │    └─ Requirements: Base + SpecifiesGenericInstance
+     │
+     ├─── LargeTensorAlgorithm
+     │    │
+     │    └─ Requirements: Base + SpecifiesLargeTensorSupport
+     │
+     └─── FwdXdlV3Algorithm
+          │
+          └─ Based on: FwdAlgorithmV3 + SpecifiesXdl
+```
+
+### Backward XDL Algorithms
+
+```
+BwdAlgorithm + SpecifiesXdl
+│
+├─── BwdXdlAlgorithmBase (ThreadClusterRank=4)
+│    │
+│    ├─── BwdXdlAlgorithm
+│    │    │
+│    │    └─ Requirements: Base + SpecifiesTransposeTransfer + SpecifiesGenericInstance
+│    │
+│    └─── BwdMultiDXdlAlgorithm
+│         │
+│         └─ Requirements: Base + SpecifiesMultipleDSupport
+│
+└─── BwdXdlV3AlgorithmBase
+     │
+     ├─── BwdXdlV3Algorithm
+     │    │
+     │    └─ Requirements: Base
+     │
+     └─── BwdTwoStageXdlAlgorithm
+          │
+          └─ Requirements: Base + SpecifiesTransposeTransfer + SpecifiesTwoStageSupport
+```
+
+**Valid XDL Algorithms:**
+- FwdXdlAlgorithm
+- FwdXdlV3Algorithm
+- LargeTensorAlgorithm
+- BwdXdlAlgorithm
+- BwdXdlV3Algorithm
+- BwdTwoStageXdlAlgorithm
+- BwdMultiDXdlAlgorithm
+
+---
+
+## WMMA Algorithm Hierarchy
+
+### Forward WMMA Algorithms
+
+```
+FwdAlgorithm + SpecifiesWmma
+│
+└─── FwdWmmaAlgorithm
+     │
+     └─ Requirements: Base  + SpecifiesWmma
+```
+
+### Backward WMMA Algorithms
+
+```
+BwdAlgorithm + SpecifiesWmma
+│
+├─── BwdWmmaAlgorithmBase (ThreadClusterRank=3)
+│    │
+│    └─── BwdWmmaAlgorithm
+│         │
+│         └─ Requirements: Base + SpecifiesNumPrefetchStages + SpecifiesGemmPipeline + SpecifiesGenericInstance
+│
+└─── BwdWmmaV3AlgorithmBase (Based on BwdAlgorithmV3)
+     │
+     ├─── BwdMultiDWmmaV3Algorithm
+     │    │
+     │    └─ Requirements: Base + SpecifiesMultipleDSupport
+     │
+     ├─── BwdWmmaV3Algorithm
+     │    │
+     │    └─ Requirements: Base + SpecifiesTransposeTransfer + SpecifiesGenericInstance
+     │
+     └─── BwdTwoStageWmmaV3Algorithm
+          │
+          └─ Requirements: Base + SpecifiesTransposeTransfer + SpecifiesTwoStageSupport
+```
+
+**Valid WMMA Algorithms:**
+- FwdWmmaAlgorithm
+- BwdWmmaAlgorithm
+- BwdWmmaV3Algorithm
+- BwdTwoStageWmmaV3Algorithm
+- BwdMultiDWmmaV3Algorithm
+
+---
+
+## DL Algorithm Hierarchy
+
+DL algorithms have a separate base and do not share the common hierarchy with XDL/WMMA algorithms.
+
+```
+DlAlgorithm
+│
+│  Requirements:
+│  • ConvAlgorithmDescriptor
+│  • SpecifiesThreadBlock
+│  • SpecifiesDlThreadConfig
+│  • SpecifiesDlThreadCluster
+│  • SpecifiesDlEpilogue
+│
+├─── FwdDlAlgorithmBase
+│    │
+│    │  Requirements: Base + SpecifiesFwdConvSpecialization + SpecifiesDlFwdBlockTransfer + SpecifiesGemmSpecialization
+│    │
+│    └─── FwdDlAlgorithm
+│
+└─── BwdDlAlgorithm
+     │
+     └─ Requirements: Base + SpecifiesBwdWeightConvSpecialization + SpecifiesDlBwdBlockTransfer
+```
+
+**Valid DL Algorithms:**
+- FwdDlAlgorithm
+- BwdDlAlgorithm
+
+---
+
+
+## Reference Algorithms
+
+```
+ReferenceAlgorithm
+│
+└─ Requirements: ConvAlgorithmDescriptor
+                + SpecifiesReferenceAlgorithm
+```
+
+Used for reference implementations and testing.
+
+## CK Tile Algorithms
+
+```
+TileAlgorithm
+│
+└─ Requirements: ConvAlgorithmDescriptor
+                + SpecifiesTileThreadBlock
+                + SpecifiesTileTransfer
+                + SpecifiesTileConvSpecialization
+                + SpecifiesTileBlockGemm
+                + SpecifiesTileOptimizations
+```
+
+The CK Tile algorithms are applicable to foward convolution as well as backwards convolution (weight and data).
+
+---
+
+## Summary for XDL/WMMA/DL algorithms
+
+| Category | Algorithm Type | Forward Variants | Backward Variants |
+|----------|---------------|------------------|-------------------|
+| **XDL** | Base | FwdXdlAlgorithmBase | BwdXdlAlgorithmBase, BwdXdlV3AlgorithmBase |
+| | Concrete | • FwdXdlAlgorithm<br>• FwdXdlV3Algorithm<br>• LargeTensorAlgorithm | • BwdXdlAlgorithm<br>• BwdXdlV3Algorithm<br>• BwdTwoStageXdlAlgorithm<br>• BwdMultiDXdlAlgorithm |
+| **WMMA** | Base | FwdAlgorithm | BwdWmmaAlgorithmBase, BwdWmmaV3AlgorithmBase |
+| | Concrete | • FwdWmmaAlgorithm | • BwdWmmaAlgorithm<br>• BwdWmmaV3Algorithm<br>• BwdTwoStageWmmaV3Algorithm<br>• BwdMultiDWmmaV3Algorithm |
+| **DL** | Base | FwdDlAlgorithmBase | DlAlgorithm |
+| | Concrete | • FwdDlAlgorithm | • BwdDlAlgorithm |
+
+---
--- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp
+++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp
@@ -7,52 +7,53 @@

 namespace ck_tile::builder::factory {

-// Base algorithm concepts
 template <typename T, size_t ThreadClusterRank = 3>
-concept TileTransferParameters =
+concept SpecifiesTileTransferParameters =
    SpecifiesThreadClusters<T, ThreadClusterRank> && SpecifiesLdsTransfer<T> &&
    SpecifiesThreadClusterAccessOrder<T> && SpecifiesSourceAccessOrder<T>;

-template <typename T>
-concept SpecifiesTileTransferParameters3D = TileTransferParameters<T, 3>;
+// Base algorithm concepts
+template <typename T, size_t ThreadClusterRank = 3>
+concept ConvAlgorithm = 
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && 
+    SpecifiesTileTransferParameters<T, ThreadClusterRank> && 
+    SpecifiesWarpGemm<T> && SpecifiesGemmPipeline<T>;;

 template <typename T>
-concept SpecifiesTileTransferParameters4D = TileTransferParameters<T, 4>;
+concept FwdAlgorithm = ConvAlgorithm<T, 3> && SpecifiesFwdConvSpecialization<T>; 
+
+template <typename T>
+concept FwdAlgorithmV3 = FwdAlgorithm<T> && SpecifiesPipelineV3<T>;
+
+template <typename T, size_t ThreadClusterRank = 3>
+concept BwdAlgorithm = ConvAlgorithm<T, ThreadClusterRank> &&  SpecifiesBwdWeightConvSpecialization<T>;
+
+template <typename T>
+concept BwdAlgorithmV3 = BwdAlgorithm<T, 3> && SpecifiesPipelineV3<T>;
+
+template <typename T>
+concept DlAlgorithm = 
+    ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && 
+    SpecifiesDlThreadConfig<T> && SpecifiesDlThreadCluster<T> && SpecifiesDlEpilogue<T>;
+
+template <typename T>
+concept FwdDlAlgorithmBase = DlAlgorithm<T> && SpecifiesFwdConvSpecialization<T> && 
+    SpecifiesDlFwdBlockTransfer<T> && SpecifiesGemmSpecialization<T>;

 template <auto Value>
-concept FwdXdlAlgorithmBase =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> && 
-    SpecifiesTileTransferParameters3D<decltype(Value)> && 
-    SpecifiesWarpGemm<decltype(Value)> && SpecifiesFwdConvSpecialization<decltype(Value)> && 
-    SpecifiesGemmPipeline<decltype(Value)> && SpecifiesXdl<Value>;
+concept FwdXdlAlgorithmBase = FwdAlgorithm<decltype(Value)> && SpecifiesXdl<Value>;

 template <auto Value>
-concept BwdXdlAlgorithmBase =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> && 
-    SpecifiesTileTransferParameters4D<decltype(Value)> && 
-    SpecifiesWarpGemm<decltype(Value)> && SpecifiesBwdWeightConvSpecialization<decltype(Value)> && 
-    SpecifiesXdl<Value>;
+concept BwdXdlAlgorithmBase = BwdAlgorithm<decltype(Value), 4> && SpecifiesXdl<Value>;

 template <auto Value>
-concept BwdXdlV3AlgorithmBase =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> && 
-    SpecifiesTileTransferParameters3D<decltype(Value)> && 
-    SpecifiesWarpGemm<decltype(Value)> && SpecifiesBwdWeightConvSpecialization<decltype(Value)> && 
-    SpecifiesGemmPipeline<decltype(Value)> && SpecifiesXdl<Value> && SpecifiesPipelineV3<decltype(Value)>;
+concept BwdXdlV3AlgorithmBase = BwdAlgorithmV3<decltype(Value)> && SpecifiesXdl<Value>;

 template <auto Value>
-concept BwdWmmaAlgorithmBase =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> && 
-    SpecifiesTileTransferParameters3D<decltype(Value)> && 
-    SpecifiesWarpGemm<decltype(Value)> && SpecifiesBwdWeightConvSpecialization<decltype(Value)> && 
-    SpecifiesWmma<Value>;
+concept BwdWmmaAlgorithmBase = BwdAlgorithm<decltype(Value), 3> && SpecifiesWmma<Value>;

 template <auto Value>
-concept BwdWmmaV3AlgorithmBase =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> && 
-    SpecifiesTileTransferParameters3D<decltype(Value)> && 
-    SpecifiesWarpGemm<decltype(Value)> && SpecifiesBwdWeightConvSpecialization<decltype(Value)> && 
-    SpecifiesGemmPipeline<decltype(Value)> && SpecifiesWmma<Value> && SpecifiesPipelineV3<decltype(Value)>;
+concept BwdWmmaV3AlgorithmBase = BwdAlgorithmV3<decltype(Value)> && SpecifiesWmma<Value>;

 // Reference algorithm concept
 template <auto Value>
@@ -72,28 +73,16 @@ template <auto Value>
 concept LargeTensorAlgorithm = FwdXdlAlgorithmBase<Value> && SpecifiesLargeTensorSupport<decltype(Value)>;

 template <auto Value>
-concept FwdXdlV3Algorithm =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> && 
-    SpecifiesTileTransferParameters3D<decltype(Value)> && 
-    SpecifiesWarpGemm<decltype(Value)> && SpecifiesFwdConvSpecialization<decltype(Value)> && 
-    SpecifiesGemmPipeline<decltype(Value)> && SpecifiesXdl<Value> && SpecifiesPipelineV3<decltype(Value)>;
+concept FwdXdlV3Algorithm = FwdAlgorithmV3<decltype(Value)> && SpecifiesXdl<Value>;

 // FWD WMMA algorithm concepts
 template <auto Value>
-concept FwdWmmaAlgorithm =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> && 
-    SpecifiesTileTransferParameters3D<decltype(Value)> && 
-    SpecifiesWarpGemm<decltype(Value)> && SpecifiesFwdConvSpecialization<decltype(Value)> && 
-    SpecifiesGemmPipeline<decltype(Value)> && SpecifiesWmma<Value>;
+concept FwdWmmaAlgorithm = FwdAlgorithm<decltype(Value)> && SpecifiesWmma<Value>;

 // FWD DL algorithms
 template <auto Value>
-concept FwdDlAlgorithm =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> && 
-    SpecifiesFwdConvSpecialization<decltype(Value)> && SpecifiesGemmSpecialization<decltype(Value)> && 
-    SpecifiesDlThreadConfig<decltype(Value)> && SpecifiesDlThreadCluster<decltype(Value)> &&
-    SpecifiesDlFwdBlockTransfer<decltype(Value)> && SpecifiesDlEpilogue<decltype(Value)>;
-
+concept FwdDlAlgorithm = FwdDlAlgorithmBase<decltype(Value)>;
+   
 // BWD weight XDL algorithm concepts
 template <auto Value>
 concept BwdXdlAlgorithm =
@@ -130,11 +119,8 @@ concept BwdTwoStageWmmaV3Algorithm = BwdWmmaV3AlgorithmBase<Value> && SpecifiesT

 // BWD weight DL algorithms
 template <auto Value>
-concept BwdDlAlgorithm =
-    ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesThreadBlock<decltype(Value)> &&
-    SpecifiesBwdWeightConvSpecialization<decltype(Value)> && SpecifiesDlThreadConfig<decltype(Value)> &&
-    SpecifiesDlThreadCluster<decltype(Value)> && SpecifiesDlBwdBlockTransfer<decltype(Value)> && 
-    SpecifiesDlEpilogue<decltype(Value)>;
+concept BwdDlAlgorithm = DlAlgorithm<decltype(Value)> && SpecifiesBwdWeightConvSpecialization<decltype(Value)> && 
+                        SpecifiesDlBwdBlockTransfer<decltype(Value)>;

 // Concepts for valid XDL/WMMA algorithms
 template <auto Value>
@@ -154,7 +140,6 @@ concept SpecifiesValidBwdWmmaAlgorithm =
    BwdWmmaAlgorithm<Value> || BwdWmmaV3Algorithm<Value> || 
    BwdTwoStageWmmaV3Algorithm<Value> || BwdMultiDWmmaV3Algorithm<Value>;

-
 template <auto Value>
 concept FwdWarpGemmOrDL = SpecifiesValidWarpGemm<Value> || FwdDlAlgorithm<Value>;

--- a/experimental/builder/test/impl/conv_algorithm_types.hpp
+++ b/experimental/builder/test/impl/conv_algorithm_types.hpp
@@ -42,7 +42,7 @@ struct GemmPipeline
 {
    size_t num_gemm_k_prefetch_stages{1};
    size_t num_conv_groups_to_merge{1};
-    PipelineVersion pipeline_version;
+    PipelineVersion pipeline_version{PipelineVersion::V1};
    PipelineScheduler scheduler{PipelineScheduler::DEFAULT};
 };
 static_assert(ckb::GemmPipelineDescriptor<GemmPipeline>);
@@ -543,6 +543,7 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle =
                          WarpGemm_,
                          InputOutputTileTransfer_<4>,
                          ConvSpecializationBwdWeight_,
+                          GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm concept.
                          TransposeParams_,
                          AlgorithmSpecialization_<>>;

@@ -592,6 +593,7 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle =
    ConvAlgorithmTemplate<ThreadBlock_,
                          WarpGemm_,
                          InputOutputTileTransfer_<4>,
+                          GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm concept.
                          ConvSpecializationBwdWeight_,
                          AlgorithmSpecialization_<MULTIPLE_D>>;