From e907e1bdf148e161c135e410c2e812ee9267b4eb Mon Sep 17 00:00:00 2001 From: Enrico Degregori Date: Wed, 28 Jan 2026 16:08:18 +0000 Subject: [PATCH] Rename pipeline (not limited to fp4 any more) --- .../38_block_scale_gemm/run_gemm_quant_example.inc | 2 +- include/ck_tile/ops/gemm_quant.hpp | 6 +++--- ..._bg_cr_base.hpp => gemm_mx_pipeline_ag_bg_cr_base.hpp} | 2 +- ...cr_policy.hpp => gemm_mx_pipeline_ag_bg_cr_policy.hpp} | 2 +- ...e_ag_bg_cr_v3.hpp => gemm_mx_pipeline_ag_bg_cr_v3.hpp} | 8 ++++---- .../ck_tile/gemm_block_scale/test_gemm_quant_fixtures.hpp | 2 +- 6 files changed, 11 insertions(+), 11 deletions(-) rename include/ck_tile/ops/gemm_quant/pipeline/{gemm_mxfp4_pipeline_ag_bg_cr_base.hpp => gemm_mx_pipeline_ag_bg_cr_base.hpp} (96%) rename include/ck_tile/ops/gemm_quant/pipeline/{gemm_mxfp4_pipeline_ag_bg_cr_policy.hpp => gemm_mx_pipeline_ag_bg_cr_policy.hpp} (99%) rename include/ck_tile/ops/gemm_quant/pipeline/{gemm_mxfp4_pipeline_ag_bg_cr_v3.hpp => gemm_mx_pipeline_ag_bg_cr_v3.hpp} (99%) diff --git a/example/ck_tile/38_block_scale_gemm/run_gemm_quant_example.inc b/example/ck_tile/38_block_scale_gemm/run_gemm_quant_example.inc index ae4acfeb0d..37c69db30e 100644 --- a/example/ck_tile/38_block_scale_gemm/run_gemm_quant_example.inc +++ b/example/ck_tile/38_block_scale_gemm/run_gemm_quant_example.inc @@ -164,7 +164,7 @@ float gemm_calc_quant(const ck_tile::QuantGemmHostArgs& args, const ck_tile::str GemmConfig::PreshuffleB, ck_tile::WPQuantBPipelineAgBgCrV2, std::conditional_t, - ck_tile::MxFp4GemmPipelineAgBgCrCompV3, + ck_tile::MxGemmPipelineAgBgCrCompV3, ck_tile::BQuantGemmPipelineAgBgCrCompV3>>; using ABQuantPipeline = diff --git a/include/ck_tile/ops/gemm_quant.hpp b/include/ck_tile/ops/gemm_quant.hpp index 696de378aa..a7edb66f15 100644 --- a/include/ck_tile/ops/gemm_quant.hpp +++ b/include/ck_tile/ops/gemm_quant.hpp @@ -21,9 +21,9 @@ #include "ck_tile/ops/gemm_quant/pipeline/gemm_bquant_pipeline_ag_bg_cr_policy.hpp" #include "ck_tile/ops/gemm_quant/pipeline/gemm_bquant_pipeline_ag_bg_cr_v3.hpp" #include "ck_tile/ops/gemm_quant/pipeline/gemm_group_quant_utils.hpp" -#include "ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_base.hpp" -#include "ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_policy.hpp" -#include "ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_v3.hpp" +#include "ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_base.hpp" +#include "ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_policy.hpp" +#include "ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp" #include "ck_tile/ops/gemm_quant/pipeline/gemm_quant_pipeline_problem.hpp" #include "ck_tile/ops/gemm_quant/pipeline/gemm_wp_abquant_pipeline_ag_bg_cr_base_policy.hpp" #include "ck_tile/ops/gemm_quant/pipeline/gemm_wp_abquant_pipeline_ag_bg_cr_v2.hpp" diff --git a/include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_base.hpp b/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_base.hpp similarity index 96% rename from include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_base.hpp rename to include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_base.hpp index facec252a3..d8b055bd2e 100644 --- a/include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_base.hpp +++ b/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_base.hpp @@ -10,7 +10,7 @@ namespace ck_tile { template -struct GemmMxFp4PipelineAgBgCrImplBase : public GemmPipelineAgBgCrImplBase +struct GemmMxPipelineAgBgCrImplBase : public GemmPipelineAgBgCrImplBase { using Base = GemmPipelineAgBgCrImplBase; using ADataType = typename Base::ADataType; diff --git a/include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_policy.hpp b/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_policy.hpp similarity index 99% rename from include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_policy.hpp rename to include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_policy.hpp index 309585304d..64fff27fa1 100644 --- a/include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_policy.hpp +++ b/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_policy.hpp @@ -9,7 +9,7 @@ namespace ck_tile { -struct GemmMxFp4PipelineAgBgCrPolicy : public UniversalGemmPipelineAgBgCrPolicy +struct GemmMxPipelineAgBgCrPolicy : public UniversalGemmPipelineAgBgCrPolicy { using Base = UniversalGemmPipelineAgBgCrPolicy; using Base::I0; diff --git a/include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_v3.hpp b/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp similarity index 99% rename from include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_v3.hpp rename to include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp index a007190d4c..03381f569a 100644 --- a/include/ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_v3.hpp +++ b/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp @@ -9,7 +9,7 @@ #include "ck_tile/core.hpp" #include "ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp" #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp" -#include "ck_tile/ops/gemm_quant/pipeline/gemm_mxfp4_pipeline_ag_bg_cr_base.hpp" +#include "ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_base.hpp" #include "ck_tile/host/concat.hpp" namespace ck_tile { @@ -18,11 +18,11 @@ namespace ck_tile { // B Tile Window: global memory // C Distributed tensor: register -template -struct MxFp4GemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 +template +struct MxGemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 { using Base = BaseGemmPipelineAgBgCrCompV3; - using PipelineImplBase = GemmMxFp4PipelineAgBgCrImplBase; + using PipelineImplBase = GemmMxPipelineAgBgCrImplBase; using ADataType = remove_cvref_t; using BDataType = remove_cvref_t; diff --git a/test/ck_tile/gemm_block_scale/test_gemm_quant_fixtures.hpp b/test/ck_tile/gemm_block_scale/test_gemm_quant_fixtures.hpp index 68aba343db..701bd2410d 100644 --- a/test/ck_tile/gemm_block_scale/test_gemm_quant_fixtures.hpp +++ b/test/ck_tile/gemm_block_scale/test_gemm_quant_fixtures.hpp @@ -877,7 +877,7 @@ class TestCkTileGemmBQuant : public TestCkTileGemmQuantBase, - ck_tile::MxFp4GemmPipelineAgBgCrCompV3, + ck_tile::MxGemmPipelineAgBgCrCompV3, ck_tile::BQuantGemmPipelineAgBgCrCompV3>, ck_tile::WPQuantBPipelineAgBgCrV2>;