From 871af334d10d82bc416aa682c2700307399ea578 Mon Sep 17 00:00:00 2001 From: rocking Date: Thu, 24 Oct 2024 20:42:40 +0000 Subject: [PATCH] Refine pipeline name --- .../pipeline/layernorm2d_fwd_pipeline_one_pass.hpp | 4 ++-- .../pipeline/layernorm2d_fwd_pipeline_two_pass.hpp | 4 ++-- .../rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp | 4 ++-- .../rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp | 4 ++-- 4 files changed, 8 insertions(+), 8 deletions(-) diff --git a/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp b/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp index bf002141b8..c767a472a9 100644 --- a/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp +++ b/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp @@ -35,9 +35,9 @@ struct Layernorm2dFwdPipelineOnePass static constexpr const char* name = []() { if constexpr(kNeedCrossWarpSync) - return "bpr"; // block per row + return "bpr_op"; // block per row else - return "wpr"; // warp per row + return "wpr_op"; // warp per row }(); CK_TILE_HOST_DEVICE static constexpr index_t GetSmemSize() diff --git a/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp b/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp index 3347c2cb90..e35d02e707 100644 --- a/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp +++ b/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp @@ -35,9 +35,9 @@ struct Layernorm2dFwdPipelineTwoPass static constexpr const char* name = []() { if constexpr(kNeedCrossWarpSync) - return "bpr"; // block per row + return "bpr_tp"; // block per row else - return "wpr"; // warp per row + return "wpr_tp"; // warp per row }(); CK_TILE_HOST_DEVICE static constexpr index_t GetSmemSize() diff --git a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp index 540ac1750c..8559485038 100644 --- a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp +++ b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp @@ -31,9 +31,9 @@ struct Rmsnorm2dFwdPipelineOnePass static constexpr const char* name = []() { if constexpr(kNeedCrossWarpSync) - return "bpr"; // block per row + return "bpr_op"; // block per row else - return "wpr"; // warp per row + return "wpr_op"; // warp per row }(); CK_TILE_HOST_DEVICE static constexpr index_t GetSmemSize() diff --git a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp index 1fa8ed9738..28e02fe651 100644 --- a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp +++ b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp @@ -31,9 +31,9 @@ struct Rmsnorm2dFwdPipelineTwoPass static constexpr const char* name = []() { if constexpr(kNeedCrossWarpSync) - return "bpr"; // block per row + return "bpr_tp"; // block per row else - return "wpr"; // warp per row + return "wpr_tp"; // warp per row }(); CK_TILE_HOST_DEVICE static constexpr index_t GetSmemSize()