Report timestamp of gridwise gemm

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

@@ -54,8 +54,14 @@ struct GridwiseGemmPipeline_v1<1>
                                const BBlockTransferStep& b_block_copy_step,
                                const BlockwiseGemm& blockwise_gemm,
                                CThreadBuffer& c_thread_buf,
-                               index_t num_loop)
+                               index_t num_loop,
+                               long& loop_start,
+                               long& loop_end)
     {
+        __builtin_amdgcn_sched_barrier(0);
+        asm volatile("; [POYENC] pipeline start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         // preload data into LDS
         a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf);
         b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf);
@@ -69,6 +75,11 @@ struct GridwiseGemmPipeline_v1<1>
         a_blockwise_copy.RunWrite(a_block_desc, a_block_buf);
         b_blockwise_copy.RunWrite(b_block_desc, b_block_buf);
 
+        __builtin_amdgcn_sched_barrier(0);
+        loop_start = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] hot-loop start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         // main body
         if constexpr(HasMainLoop)
         {
@@ -102,6 +113,15 @@ struct GridwiseGemmPipeline_v1<1>
 
             blockwise_gemm.Run(a_block_buf, b_block_buf, c_thread_buf);
         }
+
+        __builtin_amdgcn_sched_barrier(0);
+        loop_end = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] hot-loop end" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
+        __builtin_amdgcn_sched_barrier(0);
+        asm volatile("; [POYENC] pipeline end" ::);
+        __builtin_amdgcn_sched_barrier(0);
     }
 };
 
@@ -152,8 +172,14 @@ struct GridwiseGemmPipeline_v1<2>
                                const BBlockTransferStep& b_block_copy_step,
                                const BlockwiseGemm& blockwise_gemm,
                                CThreadBuffer& c_thread_buf,
-                               index_t num_loop)
+                               index_t num_loop,
+                               long& loop_start,
+                               long& loop_end)
     {
+        __builtin_amdgcn_sched_barrier(0);
+        asm volatile("; [POYENC] pipeline start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         // preload data into LDS
         {
             // Read 0
@@ -172,6 +198,11 @@ struct GridwiseGemmPipeline_v1<2>
         // Initialize C
         c_thread_buf.Clear();
 
+        __builtin_amdgcn_sched_barrier(0);
+        loop_start = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] hot-loop start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         // main body
         if constexpr(HasMainLoop)
         {
@@ -250,6 +281,15 @@ struct GridwiseGemmPipeline_v1<2>
             // Gemm num_loop - 1
             blockwise_gemm.Run(a_block_buf, b_block_buf, c_thread_buf);
         }
+
+        __builtin_amdgcn_sched_barrier(0);
+        loop_end = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] hot-loop end" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
+        __builtin_amdgcn_sched_barrier(0);
+        asm volatile("; [POYENC] pipeline end" ::);
+        __builtin_amdgcn_sched_barrier(0);
     }
 };
 
@@ -295,8 +335,14 @@ struct GridwiseGemmPipelineInterwave_v1<1>
                                const BBlockTransferStep& b_block_copy_step,
                                const BlockwiseGemm& blockwise_gemm,
                                CThreadBuffer& c_thread_buf,
-                               index_t num_loop)
+                               index_t num_loop,
+                               long& loop_start,
+                               long& loop_end)
     {
+        __builtin_amdgcn_sched_barrier(0);
+        asm volatile("; [POYENC] pipeline start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         // preload data into LDS
         a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf);
         b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf);
@@ -310,6 +356,11 @@ struct GridwiseGemmPipelineInterwave_v1<1>
         a_blockwise_copy.RunWrite(a_block_desc, a_block_buf);
         b_blockwise_copy.RunWrite(b_block_desc, b_block_buf);
 
+        __builtin_amdgcn_sched_barrier(0);
+        loop_start = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] hot-loop start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         // main body
         if constexpr(HasMainLoop)
         {
@@ -343,6 +394,15 @@ struct GridwiseGemmPipelineInterwave_v1<1>
 
             blockwise_gemm.Run(a_block_buf, b_block_buf, c_thread_buf);
         }
+
+        __builtin_amdgcn_sched_barrier(0);
+        loop_end = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] hot-loop end" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
+        __builtin_amdgcn_sched_barrier(0);
+        asm volatile("; [POYENC] pipeline end" ::);
+        __builtin_amdgcn_sched_barrier(0);
     }
 };
 

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

@@ -49,8 +49,14 @@ struct GridwiseGemmPipeline_v2
                                const BBlockTransferStep& b_block_copy_step,
                                const BlockwiseGemm& blockwise_gemm,
                                CThreadBuffer& c_thread_buf,
-                               index_t num_loop)
+                               index_t num_loop,
+                               long& loop_start,
+                               long& loop_end)
     {
+        __builtin_amdgcn_sched_barrier(0);
+        asm volatile("; [POYENC] pipeline start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         // global read 0
         a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf);
         b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf);
@@ -72,6 +78,11 @@ struct GridwiseGemmPipeline_v2
         // global Read 1
         b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf);
 
+        __builtin_amdgcn_sched_barrier(0);
+        loop_start = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] hot-loop start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         // main body
         if constexpr(HasMainLoop)
         {
@@ -122,6 +133,15 @@ struct GridwiseGemmPipeline_v2
             // GEMM num_loop - 1
             blockwise_gemm.Run(a_block_buf, b_block_buf, c_thread_buf);
         }
+
+        __builtin_amdgcn_sched_barrier(0);
+        loop_end = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] hot-loop end" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
+        __builtin_amdgcn_sched_barrier(0);
+        asm volatile("; [POYENC] pipeline end" ::);
+        __builtin_amdgcn_sched_barrier(0);
     }
 };
 

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

@@ -679,6 +679,11 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
                                void* __restrict__ p_shared,
                                const Argument& karg)
     {
+        __builtin_amdgcn_sched_barrier(0);
+        const long kernel_start = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] kernel start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
 #define CREATE_DESCS_ON_HOST 1
 #if CREATE_DESCS_ON_HOST
         const auto a_grid_desc_ak0_m_ak1 = karg.a_grid_desc_ak0_m_ak1;
@@ -850,6 +855,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
             (a_grid_desc_ak0_m_ak1.GetLength(I0) * a_grid_desc_ak0_m_ak1.GetLength(I2)) /
             KPerBlock);
 
+        long loop_start = 0, loop_end = 0;
         gridwise_gemm_pipeline.template Run<HasMainKBlockLoop>(a_grid_desc_ak0_m_ak1,
                                                                a_block_desc_ak0_m_ak1,
                                                                a_blockwise_copy,
@@ -864,7 +870,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
                                                                b_block_slice_copy_step,
                                                                blockwise_gemm,
                                                                c_thread_buf,
-                                                               num_k_block_main_loop);
+                                                               num_k_block_main_loop, loop_start, loop_end);
 
         // shuffle C and write out
         {
@@ -1061,6 +1067,16 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
                         c_grid_desc_mblock_mperblock_nblock_nperblock, c_global_step);
                 }
             });
+
+            __builtin_amdgcn_sched_barrier(0);
+            const long kernel_end = __builtin_readcyclecounter();
+            asm volatile("; [POYENC] kernel end" ::);
+            __builtin_amdgcn_sched_barrier(0);
+
+            if (blockIdx.x == 0 && threadIdx.x == 0) {
+               printf("[POYENC] prolog: %ld, hot-loop: %ld, epilog: %ld\n",
+                   loop_start - kernel_start, loop_end - loop_start, kernel_end - loop_end);
+            }
         }
     }
 };

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

@@ -492,6 +492,11 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_v2r3
                                void* __restrict__ p_shared,
                                const Argument& karg)
     {
+        __builtin_amdgcn_sched_barrier(0);
+        const long kernel_start = __builtin_readcyclecounter();
+        asm volatile("; [POYENC] kernel start" ::);
+        __builtin_amdgcn_sched_barrier(0);
+
         const auto a_grid_desc_k0_m_k1 =
             MakeAGridDescriptor_K0_M_K1(karg.M, karg.MPadded, karg.K, karg.StrideA);
         const auto b_grid_desc_k0_n_k1 =
@@ -646,6 +651,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_v2r3
         // gridwise GEMM pipeline
         const index_t num_k_block_main_loop = __builtin_amdgcn_readfirstlane(K0 / K0PerBlock);
 
+        long loop_start = 0, loop_end = 0;
         GridwiseGemmPipe::template Run<HasMainKBlockLoop>(a_grid_desc_k0_m_k1,
                                                           a_block_desc_k0_m_k1,
                                                           a_blockwise_copy,
@@ -660,7 +666,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_v2r3
                                                           b_block_slice_copy_step,
                                                           blockwise_gemm,
                                                           c_thread_buf,
-                                                          num_k_block_main_loop);
+                                                          num_k_block_main_loop, loop_start, loop_end);
 
         // output: register to global memory
         {
@@ -738,6 +744,16 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_v2r3
                               c_thread_buf,
                               c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2,
                               c_grid_buf);
+
+            __builtin_amdgcn_sched_barrier(0);
+            const long kernel_end = __builtin_readcyclecounter();
+            asm volatile("; [POYENC] kernel end" ::);
+            __builtin_amdgcn_sched_barrier(0);
+
+            if (blockIdx.x == 0 && threadIdx.x == 0) {
+               printf("[POYENC] prolog: %ld, hot-loop: %ld, epilog: %ld\n",
+                   loop_start - kernel_start, loop_end - loop_start, kernel_end - loop_end);
+            }
         }
     }
 };