From e5be038a5e8e64ec6eb56369b910bdccb8df6fb1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Fri, 16 Jan 2026 05:32:29 -0500 Subject: [PATCH] Fix GEMM pipeline concept usage. --- .../include/ck_tile/builder/factory/README.md | 7 +++---- .../ck_tile/builder/factory/conv_algorithms.hpp | 13 ++++++------- .../builder/test/impl/conv_algorithm_types.hpp | 4 ---- 3 files changed, 9 insertions(+), 15 deletions(-) diff --git a/experimental/builder/include/ck_tile/builder/factory/README.md b/experimental/builder/include/ck_tile/builder/factory/README.md index 4c30348f9a..f82b5f9ad6 100644 --- a/experimental/builder/include/ck_tile/builder/factory/README.md +++ b/experimental/builder/include/ck_tile/builder/factory/README.md @@ -49,14 +49,13 @@ XDL and WMMA algorithms share a common base, while DL algorithms have their own Both XDL and WMMA algorithms share the following foundational concepts: ``` -ConvAlgorithm (Base Concept) +ConvWarpGemmAlgorithm (Base Concept) │ │ Requirements: │ • ConvAlgorithmDescriptor │ • SpecifiesThreadBlock │ • SpecifiesTileTransferParameters (ThreadClusters, LdsTransfer, AccessOrders) │ • SpecifiesWarpGemm -│ • SpecifiesGemmPipeline │ ├─── FwdAlgorithm (Forward Convolution) │ │ @@ -64,7 +63,7 @@ ConvAlgorithm (Base Concept) │ │ │ └─── FwdAlgorithmV3 │ │ -│ │ Additional: SpecifiesPipelineV3 +│ │ Additional: SpecifiesPipelineV3 + SpecifiesGemmPipeline │ │ │ └─── BwdAlgorithm (Backward Weight Convolution) @@ -73,7 +72,7 @@ ConvAlgorithm (Base Concept) │ └─── BwdAlgorithmV3 │ - │ Additional: SpecifiesPipelineV3 + │ Additional: SpecifiesPipelineV3 + SpecifiesGemmPipeline │ ``` 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 ad72609f84..0c4d766fff 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -14,23 +14,22 @@ concept SpecifiesTileTransferParameters = // Base algorithm concepts template -concept ConvAlgorithm = ConvAlgorithmDescriptor && SpecifiesThreadBlock && +concept ConvWarpGemmAlgorithm = ConvAlgorithmDescriptor && SpecifiesThreadBlock && SpecifiesTileTransferParameters && - SpecifiesWarpGemm && SpecifiesGemmPipeline; -; + SpecifiesWarpGemm; template -concept FwdAlgorithm = ConvAlgorithm && SpecifiesFwdConvSpecialization; +concept FwdAlgorithm = ConvWarpGemmAlgorithm && SpecifiesFwdConvSpecialization; template -concept FwdAlgorithmV3 = FwdAlgorithm && SpecifiesPipelineV3; +concept FwdAlgorithmV3 = FwdAlgorithm && SpecifiesPipelineV3 && SpecifiesGemmPipeline; template concept BwdAlgorithm = - ConvAlgorithm && SpecifiesBwdWeightConvSpecialization; + ConvWarpGemmAlgorithm && SpecifiesBwdWeightConvSpecialization; template -concept BwdAlgorithmV3 = BwdAlgorithm && SpecifiesPipelineV3; +concept BwdAlgorithmV3 = BwdAlgorithm && SpecifiesPipelineV3 && SpecifiesGemmPipeline; template concept DlAlgorithm = diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp index cf4bc25c0f..6989887a26 100644 --- a/experimental/builder/test/impl/conv_algorithm_types.hpp +++ b/experimental/builder/test/impl/conv_algorithm_types.hpp @@ -543,8 +543,6 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle = WarpGemm_, InputOutputTileTransfer_<4>, ConvSpecializationBwdWeight_, - GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm - // concept. TransposeParams_, AlgorithmSpecialization_<>>; @@ -594,8 +592,6 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle = ConvAlgorithmTemplate, - GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm - // concept. ConvSpecializationBwdWeight_, AlgorithmSpecialization_>;