From 62b1f2e830c1fb7f28693148e184d9c5203216dc Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Wed, 18 Dec 2024 02:00:39 +0000 Subject: [PATCH] roll back splitkv [ROCm/composable_kernel commit: 60113859faf4bf4899b89b87091a237977d91a98] --- .../fmha/pipeline/block_fmha_fwd_splitkv_combine_pipeline.hpp | 2 -- .../pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs.hpp | 4 ---- 2 files changed, 6 deletions(-) diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_combine_pipeline.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_combine_pipeline.hpp index 4b16b1fc81..7ac86e6d12 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_combine_pipeline.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_combine_pipeline.hpp @@ -95,8 +95,6 @@ struct BlockFmhaFwdSplitKVCombinePipeline { constexpr std::array occupancy{2, 2, 2, 2, 2, 1}; return occupancy[detail::log2::value - 2]; - } else if constexpr(kHeadDimV <= 512) { - return 1; } } }(); diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs.hpp index 01a7bd36f4..04aa85644d 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs.hpp @@ -96,10 +96,6 @@ struct BlockFmhaFwdSplitKVPipelineQRKSVS { return 1; } - else if constexpr(kQKHeaddim <= 512) - { - return 1; - } } }();