From 52c018b7456f38ab4ba6989cdea495ef2d4f4902 Mon Sep 17 00:00:00 2001 From: mtgu0705 Date: Thu, 13 Feb 2025 11:12:57 +0800 Subject: [PATCH] fix bug, function pass. --- .../block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v1.hpp | 4 ++-- .../gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v1.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v1.hpp index 0e758bbbee..73c892382d 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v1.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v1.hpp @@ -365,10 +365,10 @@ struct BlockwiseGemmXdlops_pipeline_bpreshuffle_v1VGPR dequant b_thread_dequant_copy.Run(b_block_desc_n0_n1_k0_k1, b_block_origin_idx, - b_thread_bufs(mfma_reg_buf), + b_thread_bufs(local_read_buf), b_thread_desc_, make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(mfma_reg_buf)); + b_thread_dequant_bufs(local_read_buf)); HotLoopScheduler(); __builtin_amdgcn_sched_barrier(0); diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp index 1e4ceebeeb..5f15f24288 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp @@ -1620,7 +1620,7 @@ struct GridwiseGemm_xdl_cshuffle_v3_b_preshuffle decltype(b_block_desc_bk0_n_bk1), tensor_operation::element_wise::PassThrough, Sequence{}, I1, Number{}, Number{}>, - Sequence<0, 1, 2, 3>, + Sequence<1, 2, 0, 3>, 3, BK1Number>(b_element_op);