From dc6f8f1067f018d35ec447a9f48012fe0c214ab3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Thu, 15 Jan 2026 05:27:22 -0500 Subject: [PATCH] Fix WMMA conv algorithms hierarchy. --- .../include/ck_tile/builder/conv_algorithm_concepts.hpp | 5 ----- .../builder/include/ck_tile/builder/factory/README.md | 4 ++-- .../include/ck_tile/builder/factory/conv_algorithms.hpp | 7 +++---- ...test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp | 2 +- 4 files changed, 6 insertions(+), 12 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 ce60bf4a94..ffe265b0cc 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -260,11 +260,6 @@ concept SpecifiesGemmSpecialization = requires { { T::gemm_specialization } -> std::convertible_to; }; -template -concept SpecifiesNumPrefetchStages = requires { - { T::num_gemm_k_prefetch_stages } -> SizeType; -}; - template concept SpecifiesNumGroupsToMerge = requires { { T::num_conv_groups_to_merge } -> SizeType; diff --git a/experimental/builder/include/ck_tile/builder/factory/README.md b/experimental/builder/include/ck_tile/builder/factory/README.md index 4c30348f9a..ab3fe73755 100644 --- a/experimental/builder/include/ck_tile/builder/factory/README.md +++ b/experimental/builder/include/ck_tile/builder/factory/README.md @@ -159,7 +159,7 @@ BwdAlgorithm + SpecifiesWmma │ │ │ └─── BwdWmmaAlgorithm │ │ -│ └─ Requirements: Base + SpecifiesNumPrefetchStages + SpecifiesGemmPipeline + SpecifiesGenericInstance +│ └─ Requirements: Base + SpecifiesGemmPipeline + SpecifiesGenericInstance │ └─── BwdWmmaV3AlgorithmBase (Based on BwdAlgorithmV3) │ @@ -169,7 +169,7 @@ BwdAlgorithm + SpecifiesWmma │ ├─── BwdWmmaV3Algorithm │ │ - │ └─ Requirements: Base + SpecifiesTransposeTransfer + SpecifiesGenericInstance + │ └─ Requirements: Base + SpecifiesTransposeTransfer │ └─── BwdTwoStageWmmaV3Algorithm │ 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..8327b7a2ad 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -108,8 +108,8 @@ concept BwdTwoStageXdlAlgorithm = // BWD weight WMMA algorithm concepts template concept BwdWmmaAlgorithm = - BwdWmmaAlgorithmBase && SpecifiesNumPrefetchStages && - SpecifiesGemmPipeline && SpecifiesGenericInstance; + BwdWmmaAlgorithmBase && SpecifiesGemmPipeline && + SpecifiesGenericInstance; template concept BwdMultiDWmmaV3Algorithm = @@ -117,8 +117,7 @@ concept BwdMultiDWmmaV3Algorithm = template concept BwdWmmaV3Algorithm = - BwdWmmaV3AlgorithmBase && SpecifiesTransposeTransfer && - SpecifiesGenericInstance; + BwdWmmaV3AlgorithmBase && SpecifiesTransposeTransfer; template concept BwdTwoStageWmmaV3Algorithm = diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp index 9310f5a9a6..9ff27f6f99 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp @@ -19,7 +19,7 @@ constexpr auto SIGNATURE = ckt::ConvSignature{.spatial_dim = 2, .weight = {.config = {.layout = GKYXC}}, .output = {.config = {.layout = NGKHW}}}; -constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle{} +constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle_V3{} .with_thread_block(cku::ThreadBlock_64_32x32x32) .with_gemm_config(cku::GemmParams_Wmma_16x16_2x1_per_wave) .with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3)