fixed MPerBlock=96

composable_kernel/include/tensor_operation/threadwise_tensor_slice_transfer.hpp

@@ -165,7 +165,7 @@ struct ThreadwiseTensorSliceTransfer_v1r3
                 static_for<1, nDim, 1>{}([&](auto i) {
                     index_t tmp = ordered_access_idx[I0];
 
-                    static_for<0, i, 1>{}([&](auto j) {
+                    static_for<1, i, 1>{}([&](auto j) {
                         tmp = tmp * ordered_access_lengths[j] + ordered_access_idx[j];
                     });
 

device_operation/include/device_gemm_splitk_xdl.hpp

@@ -426,7 +426,6 @@ struct DeviceGemmSplitKXdl
             float ave_time = 0;
 
             const auto Run = [&](const auto& kernel) {
-#if CK_RUN_KERNEL_AND_TIME
                 ave_time = launch_and_time_kernel(kernel,
                                                   nrepeat,
                                                   dim3(grid_size),
@@ -442,23 +441,29 @@ struct DeviceGemmSplitKXdl
                                                   arg.b_element_op_,
                                                   arg.c_element_op_,
                                                   arg.block_2_ctile_map_);
-#else
-                nrepeat++;
-                launch_kernel(kernel,
-                              dim3(grid_size),
-                              dim3(BlockSize),
-                              0,
-                              arg.p_a_grid_,
-                              arg.p_b_grid_,
-                              arg.p_c_grid_,
-                              arg.a_grid_desc_kbatch_k0_m_k1_,
-                              arg.b_grid_desc_kbatch_k0_n_k1_,
-                              arg.c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2_,
-                              arg.a_element_op_,
-                              arg.b_element_op_,
-                              arg.c_element_op_,
-                              arg.block_2_ctile_map_);
-#endif
+                if(kbatch > 1)
+                {
+                    hipGetErrorString(
+                        hipMemset(arg.p_c_grid_,
+                                  0,
+                                  arg.c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2_.GetElementSpaceSize() *
+                                      sizeof(CDataType)));
+
+                    launch_kernel(kernel,
+                                  dim3(grid_size),
+                                  dim3(BlockSize),
+                                  0,
+                                  arg.p_a_grid_,
+                                  arg.p_b_grid_,
+                                  arg.p_c_grid_,
+                                  arg.a_grid_desc_kbatch_k0_m_k1_,
+                                  arg.b_grid_desc_kbatch_k0_n_k1_,
+                                  arg.c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2_,
+                                  arg.a_element_op_,
+                                  arg.b_element_op_,
+                                  arg.c_element_op_,
+                                  arg.block_2_ctile_map_);
+                }
             };
             if(has_main_k0_block_loop)
             {