Update sheduling pattern for IGLP_OPT_STRATEGY=1

fcef1550 · Po-Yen, Chen · 08d12d4f · fcef1550
Commit fcef1550 authored Dec 12, 2022 by Po-Yen, Chen
Show whitespace changes
Inline Side-by-side

Showing with 1 addition and 6 deletions

include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v2.hpp ...k/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v2.hpp +1 -6

No files found.
--- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v2.hpp
+++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v2.hpp
@@ -110,25 +110,20 @@ struct GridwiseGemmPipeline_v2
                __builtin_amdgcn_sched_group_barrier(0x020, 2, 0); // VMEM read
                __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, 1, 0); // MFMA

                __builtin_amdgcn_sched_group_barrier(0x020, 2, 0); // VMEM read
                __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, 1, 0); // MFMA

-                __builtin_amdgcn_sched_group_barrier(0x020, 2, 0); // VMEM read
+                __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read
                __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, 1, 0); // MFMA

-                __builtin_amdgcn_sched_group_barrier(0x020, 2, 0); // VMEM read
                __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, 1, 0); // MFMA
 #elif IGLP_OPT_STRATEGY == 2
                // 16 MFMAs