change read order

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_skip_b_lds_v1.hpp

@@ -113,7 +113,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_b_lds_v1
     static constexpr auto I7 = Number<7>{};
 
     static constexpr auto BaseMultK0 = 4;
-    static constexpr auto MultiK0    = BaseMultK0 * 1;
+    static constexpr auto MultiK0    = BaseMultK0 * 2;
 
     // K1 should be Number<...>
     static constexpr auto K1 = Number<K1Value>{};
@@ -571,11 +571,6 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_b_lds_v1
                                               b_thread_3rd_buf);
                         b_threadwise_copy.MoveSrcSliceWindow(b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3,
                                                              b_thread_slice_copy_step);
-
-                        blockwise_gemm.Run(a_block_buf, b_thread_1st_buf, c_thread_buf);
-                        blockwise_gemm.MoveABlockSliceWindow();
-                        s_nop();
-                        // 2nd
                         b_threadwise_copy.Run(b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3,
                                               b_grid_buf,
                                               b_thread_desc_k0_k1_k2_n0_n1_n2_n3_k3,
@@ -583,10 +578,15 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_b_lds_v1
                                               b_thread_4th_buf);
                         b_threadwise_copy.MoveSrcSliceWindow(b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3,
                                                              b_thread_slice_copy_step);
+                        s_nop();
+                        blockwise_gemm.Run(a_block_buf, b_thread_1st_buf, c_thread_buf);
+                        blockwise_gemm.MoveABlockSliceWindow();
+
+                        // 2nd
 
                         blockwise_gemm.Run(a_block_buf, b_thread_2nd_buf, c_thread_buf);
                         blockwise_gemm.MoveABlockSliceWindow();
-                        s_nop();
+
                         // 3rd
                         b_threadwise_copy.Run(b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3,
                                               b_grid_buf,
@@ -595,11 +595,6 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_b_lds_v1
                                               b_thread_1st_buf);
                         b_threadwise_copy.MoveSrcSliceWindow(b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3,
                                                              b_thread_slice_copy_step);
-
-                        blockwise_gemm.Run(a_block_buf, b_thread_3rd_buf, c_thread_buf);
-                        blockwise_gemm.MoveABlockSliceWindow();
-                        s_nop();
-                        // 4th
                         b_threadwise_copy.Run(b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3,
                                               b_grid_buf,
                                               b_thread_desc_k0_k1_k2_n0_n1_n2_n3_k3,
@@ -607,6 +602,11 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_b_lds_v1
                                               b_thread_2nd_buf);
                         b_threadwise_copy.MoveSrcSliceWindow(b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3,
                                                              b_thread_slice_copy_step);
+                        s_nop();
+                        blockwise_gemm.Run(a_block_buf, b_thread_3rd_buf, c_thread_buf);
+                        blockwise_gemm.MoveABlockSliceWindow();
+
+                        // 4th
 
                         blockwise_gemm.Run(a_block_buf, b_thread_4th_buf, c_thread_buf);
                         blockwise_gemm.MoveABlockSliceWindow();

include/ck/utility/synchronization.hpp

@@ -16,7 +16,8 @@ __device__ void block_sync_lds()
     __syncthreads();
 #endif
 }
-__device__ void s_nop(){
+__device__ void s_nop()
+{
     asm volatile("\
     s_nop 0 \n \
     " ::);