adjust for fp16

composable_kernel/include/tensor_operation/gridwise_gemm_xdlops_v2r4.hpp

@@ -439,9 +439,10 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_v2r4
                                                 b_block_slice_copy_step,
                                                 b_k0_n_k1_grid_move_slice_window_step_hack);
 
+            a_blockwise_copy.RunRead(a_b_k0_m_k1_grid_desc, a_grid_buf, a_k0_m_k1_grid_step_hacks);
+
             block_sync_lds();
 
-            a_blockwise_copy.RunRead(a_b_k0_m_k1_grid_desc, a_grid_buf, a_k0_m_k1_grid_step_hacks);
             b_blockwise_copy.RunRead(b_b_k0_n_k1_grid_desc, b_grid_buf, b_k0_n_k1_grid_step_hacks);
 
             blockwise_gemm.Run(a_block_buf, b_block_buf, c_thread_buf);

host/driver_offline/include/device_convolution_backward_weight_implicit_gemm_v4r4r5_xdlops_atomic_nhwc_kyxc_nhwk.hpp

@@ -223,7 +223,7 @@ void device_convolution_backward_weight_implicit_gemm_v4r4r5_xdlops_atomic_nhwc_
             GemmABlockTransferThreadSliceLengths_GemmK0_GemmM_GemmK1,
             GemmABlockTransferThreadClusterLengths_GemmK0_GemmM_GemmK1,
             Sequence<0, 1, 2, 3>,
-            Sequence<0, 1, 2, 3>,
+            Sequence<0, 1, 3, 2>,
             2,
             GemmABlockTransferSrcScalarPerVector_GemmM,
             GemmABlockTransferDstScalarPerVector_GemmK1,
@@ -231,7 +231,7 @@ void device_convolution_backward_weight_implicit_gemm_v4r4r5_xdlops_atomic_nhwc_
             GemmBBlockTransferThreadSliceLengths_GemmK0_GemmN_GemmK1,
             GemmBBlockTransferThreadClusterLengths_GemmK0_GemmN_GemmK1,
             Sequence<0, 1, 2, 3>,
-            Sequence<0, 1, 2, 3>,
+            Sequence<0, 1, 3, 2>,
             2,
             GemmBBlockTransferSrcScalarPerVector_GemmN,
             GemmBBlockTransferDstScalarPerVector_GemmK1,