opt2

include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v4.hpp

@@ -321,24 +321,22 @@ struct BlockwiseGemmXdlops_pipeline_v4
             ignore = i;
             static_for<0, 2*num_ds_read_inst/num_issue, 1>{}([&](auto ir) {
                 ignore = ir;
-                __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
                 __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read
+                __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
             });
-            __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
             __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read
-            __builtin_amdgcn_sched_group_barrier(0x008, num_mfma_inst/num_buffer_load_inst-2*num_ds_read_inst/num_issue -1, 0); // MFMA
+            __builtin_amdgcn_sched_group_barrier(0x008, num_mfma_inst/num_buffer_load_inst-2*num_ds_read_inst/num_issue, 0); // MFMA
         });
 
         static_for<0, num_issue/2, 1>{}([&](auto i) {
             ignore = i;
             static_for<0, 2*num_ds_write_inst/num_issue, 1>{}([&](auto iw) {
                 ignore = iw;
-                __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
                 __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS write
-            });
-            __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
+                __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
+           });
             __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read
-            __builtin_amdgcn_sched_group_barrier(0x008, num_mfma_inst/num_buffer_load_inst-2*num_ds_write_inst/num_issue -1, 0); // MFMA
+            __builtin_amdgcn_sched_group_barrier(0x008, num_mfma_inst/num_buffer_load_inst-2*num_ds_write_inst/num_issue, 0); // MFMA
         });
     }