diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v3.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v3.hpp index 462a1727a9..72d520cc12 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v3.hpp @@ -487,6 +487,7 @@ struct BlockwiseGemmXdlops_pipeline_bpreshuffle_v3{}([&](auto k0) { static_for<0, KGroup, 1>{}([&](auto kg0) { a_thread_copy_.Run( @@ -625,6 +626,7 @@ struct BlockwiseGemmXdlops_pipeline_bpreshuffle_v3{}([&](auto k0) { static_for<0, KGroup, 1>{}([&](auto kg0) { a_thread_copy_.Run( 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 e102864db9..c1b91760c1 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 @@ -392,7 +392,6 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3_b_preshuffle __host__ __device__ static auto MakeBGridDescriptor_Preshuffled(index_t N0, index_t K0) { - constexpr index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl); constexpr index_t WaveSize = BlockSize / (MWave * NWave); constexpr index_t NkSwizzleNumber = Number{}; return make_naive_tensor_descriptor( @@ -901,7 +900,6 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3_b_preshuffle __device__ static constexpr auto GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1() { - constexpr index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl); constexpr index_t WaveSize = BlockSize / (MWave * NWave); // A matrix in LDS memory, dst of blockwise copy if constexpr(ABlockLdsExtraM) @@ -2202,11 +2200,13 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3_b_preshuffle const Block2CTileMap& block_2_ctile_map) { ignore = b_element_op; + index_t BN0Shuffled = CalculateBN0Shuffled(problem.N); + index_t BK0Shuffled = CalculateBK0Shuffled(problem.K); const auto a_grid_desc_ak0_m_ak1 = MakeAGridDescriptor_AK0_M_AK1( problem.M, problem.MPadded, problem.K, problem.KPadded, problem.StrideA, problem.AK0); const auto b_grid_desc_bpreshuffled = - MakeBGridDescriptor_Preshuffled(problem.BN0Shuffled, problem.BK0Shuffled); + MakeBGridDescriptor_Preshuffled(BN0Shuffled, BK0Shuffled); const auto c_grid_desc_m_n = MakeCGridDescriptor_M_N( problem.M, problem.MPadded, problem.N, problem.NPadded, problem.StrideC);