Update 16-MFMA scheduling pattern

e39a4794 · Po-Yen, Chen · 686a0c20 · e39a4794
Commit e39a4794 authored Dec 12, 2022 by Po-Yen, Chen
Hide whitespace changes
Inline Side-by-side

Showing with 1 addition and 14 deletions

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

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
@@ -131,49 +131,36 @@ 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
                // cluster #2
-                __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(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
 #endif
 #endif