From 66d6f152a71dd3f01c7a2d8227a884df7268aeb1 Mon Sep 17 00:00:00 2001 From: Jin Zhou Date: Mon, 8 Sep 2025 02:42:15 +0000 Subject: [PATCH] refine --- ...ce_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp | 10 ---------- ...dwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp | 5 ++--- 2 files changed, 2 insertions(+), 13 deletions(-) diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp index 97ab45da5a..55e8929237 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp @@ -299,7 +299,6 @@ struct DeviceGemmMultiD_Xdl_CShuffle_V3_BPreshuffle else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v2 || BlkGemmPipelineVer == BlockGemmPipelineVersion::v3) { - #if 0 if(arg.KBatch > 1) { if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd) @@ -325,10 +324,7 @@ struct DeviceGemmMultiD_Xdl_CShuffle_V3_BPreshuffle Run(kernel); } } - else - #endif { - #if 0 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd) { const auto kernel = @@ -340,8 +336,6 @@ struct DeviceGemmMultiD_Xdl_CShuffle_V3_BPreshuffle TailNumber::Odd>; Run(kernel); } - else - #endif { const auto kernel = kernel_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle_2lds< @@ -359,12 +353,10 @@ struct DeviceGemmMultiD_Xdl_CShuffle_V3_BPreshuffle throw std::runtime_error("todo: only v1 v2 and v3 support now"); } } -#if 0 else { if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v1) { -#if 0 if(arg.KBatch > 1) { if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd) @@ -411,7 +403,6 @@ struct DeviceGemmMultiD_Xdl_CShuffle_V3_BPreshuffle Run(kernel); } } -#endif if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd) { const auto kernel = kernel_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle< @@ -491,7 +482,6 @@ struct DeviceGemmMultiD_Xdl_CShuffle_V3_BPreshuffle throw std::runtime_error("todo: only v3 support now"); } } -#endif return ave_time; } diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp index 999f7db331..21e8fec05f 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp @@ -304,8 +304,8 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3_b_preshuffle make_tuple(make_merge_transform_v3_division_mod(make_tuple(Number{}, Number{})), make_unmerge_transform(make_tuple( Number{}, Number{}, Number{}))), - make_tuple(Sequence<0, 2>{}, Sequence<1>{}), // old 0 2 1 - make_tuple(Sequence<3>{}, Sequence<0, 1, 2>{})); // new 3 0 1 2 + make_tuple(Sequence<0, 2>{}, Sequence<1>{}), + make_tuple(Sequence<3>{}, Sequence<0, 1, 2>{})); } __host__ __device__ static auto MakeAGridDescriptor_AK0_M_AK1( @@ -783,7 +783,6 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3_b_preshuffle MBlock{CalculateMBlock(M_)}, NBlock{CalculateNBlock(N_)} { - Print(); } __host__ void Print() const