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()