From 35e372edca162c8cad64e0d4b0123d3d2af28b78 Mon Sep 17 00:00:00 2001 From: Aleksander Dudek Date: Fri, 20 Jun 2025 04:48:14 -0500 Subject: [PATCH] Change name to avoid conflict --- example/ck_tile/03_gemm/gemm_utils.hpp | 12 ++++++------ example/ck_tile/03_gemm/universal_gemm.cpp | 2 +- ...omp_v5.hpp => gemm_pipeline_ag_bg_cr_comp_v6.hpp} | 10 +++++----- ...emm_pipeline_ag_bg_cr_comp_v6_default_policy.hpp} | 6 +++--- 4 files changed, 15 insertions(+), 15 deletions(-) rename include/ck_tile/ops/gemm/pipeline/{gemm_pipeline_ag_bg_cr_comp_v5.hpp => gemm_pipeline_ag_bg_cr_comp_v6.hpp} (99%) rename include/ck_tile/ops/gemm/pipeline/{gemm_pipeline_ag_bg_cr_comp_v5_default_policy.hpp => gemm_pipeline_ag_bg_cr_comp_v6_default_policy.hpp} (91%) diff --git a/example/ck_tile/03_gemm/gemm_utils.hpp b/example/ck_tile/03_gemm/gemm_utils.hpp index 715f51e735..9e4b3893ba 100644 --- a/example/ck_tile/03_gemm/gemm_utils.hpp +++ b/example/ck_tile/03_gemm/gemm_utils.hpp @@ -14,10 +14,10 @@ #define CK_TILE_PIPELINE_COMPUTE_V3 1 #define CK_TILE_PIPELINE_MEMORY 2 #define CK_TILE_PIPELINE_COMPUTE_V4 3 -#define CK_TILE_PIPELINE_COMPUTE_V5 4 +#define CK_TILE_PIPELINE_COMPUTE_V6 5 #ifndef CK_TILE_PIPELINE_DEFAULT -#define CK_TILE_PIPELINE_DEFAULT CK_TILE_PIPELINE_COMPUTE_V5 +#define CK_TILE_PIPELINE_DEFAULT CK_TILE_PIPELINE_COMPUTE_V6 #endif #if(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_MEMORY) @@ -32,9 +32,9 @@ #define GEMM_PIPELINE ck_tile::GemmPipelineAgBgCrCompV4 #define UNIVERSAL_GEMM_PIPELINE ck_tile::BaseGemmPipelineAgBgCrCompV4 #define GEMM_PIPELINE_SCHEDULER ck_tile::GemmPipelineScheduler::Intrawave -#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V5) -#define GEMM_PIPELINE ck_tile::GemmPipelineAgBgCrCompV5 -#define UNIVERSAL_GEMM_PIPELINE ck_tile::BaseGemmPipelineAgBgCrCompV5 +#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V6) +#define GEMM_PIPELINE ck_tile::GemmPipelineAgBgCrCompV6 +#define UNIVERSAL_GEMM_PIPELINE ck_tile::BaseGemmPipelineAgBgCrCompV6 #define GEMM_PIPELINE_SCHEDULER ck_tile::GemmPipelineScheduler::Intrawave #else #error "unsupported CK_TILE_PIPELINE_DEFAULT value" @@ -89,7 +89,7 @@ struct GemmConfig static constexpr ck_tile::index_t K_Warp_Tile = 16; static constexpr bool DoubleSmemBuffer = true; -#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V5) +#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V6) // Compute friendly for Intrawave scheduler // Using the ping pong reader in the lds level static constexpr ck_tile::index_t M_Tile = 256; diff --git a/example/ck_tile/03_gemm/universal_gemm.cpp b/example/ck_tile/03_gemm/universal_gemm.cpp index 0a98ce3f88..435df1afda 100644 --- a/example/ck_tile/03_gemm/universal_gemm.cpp +++ b/example/ck_tile/03_gemm/universal_gemm.cpp @@ -211,7 +211,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& RunSplitk(ck_tile::bool_constant{}, ck_tile::integral_constant{}); } -#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V5) +#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V6) if(tail_num == ck_tile::TailNumber::Three) { RunSplitk( diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6.hpp similarity index 99% rename from include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5.hpp rename to include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6.hpp index 3cfac30fc9..e9cf0f017c 100644 --- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5.hpp +++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6.hpp @@ -4,7 +4,7 @@ #include "ck_tile/core.hpp" #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp" #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_base.hpp" -#include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5_default_policy.hpp" +#include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6_default_policy.hpp" namespace ck_tile { @@ -12,7 +12,7 @@ namespace ck_tile { // B Tile Window: global memory // C Distributed tensor: register template -struct BaseGemmPipelineAgBgCrCompV5 +struct BaseGemmPipelineAgBgCrCompV6 { static constexpr index_t PrefetchStages = 3; static constexpr index_t PrefillStages = 1; @@ -45,10 +45,10 @@ struct BaseGemmPipelineAgBgCrCompV5 * * @note TODO */ -template -struct GemmPipelineAgBgCrCompV5 : public BaseGemmPipelineAgBgCrCompV5 +template +struct GemmPipelineAgBgCrCompV6 : public BaseGemmPipelineAgBgCrCompV6 { - using Base = BaseGemmPipelineAgBgCrCompV5; + using Base = BaseGemmPipelineAgBgCrCompV6; using PipelineImplBase = GemmPipelineAgBgCrImplBase; using ADataType = remove_cvref_t; diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5_default_policy.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6_default_policy.hpp similarity index 91% rename from include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5_default_policy.hpp rename to include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6_default_policy.hpp index 74aeccdbe5..7307981f5f 100644 --- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5_default_policy.hpp +++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6_default_policy.hpp @@ -9,13 +9,13 @@ #include "ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp" namespace ck_tile { -// Default policy for GemmPipelineAGmemBGmemCregComputeV5, except the block gemm method, it shares +// Default policy for GemmPipelineAGmemBGmemCregComputeV6, except the block gemm method, it shares // the same vector size implementation, SmemSize, Global memory tile distiribution as the // UniversalGemm Pipeline Policy. // Default policy class should not be templated, put template on // member functions instead. -struct GemmPipelineAgBgCrCompV5DefaultPolicy - : public UniversalGemmBasePolicy +struct GemmPipelineAgBgCrCompV6DefaultPolicy + : public UniversalGemmBasePolicy { template CK_TILE_HOST_DEVICE static constexpr auto GetBlockGemm()