diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5.hpp
index 108438131e..44b4b32aa8 100644
--- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5.hpp
+++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5.hpp
@@ -507,8 +507,6 @@ struct GemmPipelineAgBgCrCompV5 : public BaseGemmPipelineAgBgCrCompV5<Problem>
                 } while(i < (num_loop - PrefetchStages));
             }
 
-            // <<< ================== TODO WORK ================ >>>
-
             // tail
             auto ReadWriteCompFunc = [&]() {
                 static_for<0, KRepeat, 1>{}([&](auto k0) {
@@ -549,140 +547,74 @@ struct GemmPipelineAgBgCrCompV5 : public BaseGemmPipelineAgBgCrCompV5<Problem>
                     // Local prefetch 3
                     Base::LocalPrefetch(a_block_tile0, a_lds_ld_window0);
                     Base::LocalPrefetch(b_block_tile0, b_lds_ld_window0);
+                });
 
-                    HotLoopScheduler();
+                HotLoopScheduler();
             };
+
             auto ReadCompFunc = [&]() {
-                    vector_type<ComputeDataType, KPack> a_thread_vec;
-                    vector_type<ComputeDataType, KPack> b_thread_vec;
+                block_gemm(c_block_tile, a_block_tile0, b_block_tile0);
 
-                    static_for<0, KRepeat - 1, 1>{}([&](auto k0) {
-                        static_for<0, MRepeat, 1>{}([&](auto m0) {
-                            static_for<0, NRepeat, 1>{}([&](auto n0) {
-                                static_for<0, KPack, 1>{}([&](auto ik) {
-                                    a_thread_vec.template AsType<ComputeDataType>()(ik) =
-                                        a_thread_buf[Number<a_thread_desc_.CalculateOffset(
-                                            make_tuple(m0, I0, I0, ik))>{}];
-                                });
-                                static_for<0, KPack, 1>{}([&](auto ik) {
-                                    b_thread_vec.template AsType<ComputeDataType>()(ik) =
-                                        b_thread_buf[Number<b_thread_desc_.CalculateOffset(
-                                            make_tuple(n0, I0, I0, ik))>{}];
-                                });
+                // Local prefetch 4
+                Base::LocalPrefetch(a_block_tile0, a_lds_ld_window0);
+                Base::LocalPrefetch(b_block_tile0, b_lds_ld_window0);
 
-                                using mfma_input_type =
-                                    typename vector_type<ComputeDataType,
-                                                         xdlops_gemm.K1PerXdlops>::type;
+                // TODO verify?
+                block_gemm(c_block_tile, a_block_tile0, b_block_tile0);
 
-                                constexpr index_t c_offset =
-                                    c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
-
-                                xdlops_gemm.Run(
-                                    a_thread_vec.template AsType<mfma_input_type>(),
-                                    b_thread_vec.template AsType<mfma_input_type>(),
-                                    c_thread_buf.GetVectorTypeReference(Number<c_offset>{}));
-                            });
-
-                            a_thread_copy_.Run(
-                                a_block_desc_m0_m1_m2_k,
-                                make_tuple(m0, I0, I0, Number<(k0 + 1) % KRepeat * AMmaKStride>{}),
-                                a_block_buf,
-                                a_thread_desc_,
-                                make_tuple(m0, I0, I0, I0),
-                                a_thread_buf);
-                        });
-
-                        static_for<0, NRepeat, 1>{}([&](auto n0) {
-                            b_thread_copy_.Run(
-                                b_block_desc_n0_n1_n2_k,
-                                make_tuple(n0, I0, I0, Number<(k0 + 1) % KRepeat * BMmaKStride>{}),
-                                b_block_buf,
-                                b_thread_desc_,
-                                make_tuple(n0, I0, I0, I0),
-                                b_thread_buf);
-                        });
-                    });
-
-                    static_for<0, MRepeat, 1>{}([&](auto m0) {
-                        static_for<0, NRepeat, 1>{}([&](auto n0) {
-                            static_for<0, KPack, 1>{}([&](auto ik) {
-                                a_thread_vec.template AsType<ComputeDataType>()(ik) =
-                                    a_thread_buf[Number<a_thread_desc_.CalculateOffset(
-                                        make_tuple(m0, I0, I0, ik))>{}];
-                            });
-                            static_for<0, KPack, 1>{}([&](auto ik) {
-                                b_thread_vec.template AsType<ComputeDataType>()(ik) =
-                                    b_thread_buf[Number<b_thread_desc_.CalculateOffset(
-                                        make_tuple(n0, I0, I0, ik))>{}];
-                            });
-
-                            using mfma_input_type =
-                                typename vector_type<ComputeDataType,
-                                                     xdlops_gemm.K1PerXdlops>::type;
-
-                            constexpr index_t c_offset =
-                                c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
-
-                            xdlops_gemm.Run(
-                                a_thread_vec.template AsType<mfma_input_type>(),
-                                b_thread_vec.template AsType<mfma_input_type>(),
-                                c_thread_buf.GetVectorTypeReference(Number<c_offset>{}));
-                        });
-                    });
-
-                    HotLoopScheduler();
+                HotLoopScheduler();
             };
 
             if constexpr(TailNum == TailNumber::Odd)
             {
-                    ReadWriteCompFunc(I0);
-                    ReadWriteCompFunc(I1);
-                    ReadCompFunc();
+                ReadWriteCompFunc(I0);
+                ReadWriteCompFunc(I1);
+                ReadCompFunc();
             }
             else if constexpr(TailNum == TailNumber::Even)
             {
-                    ReadWriteCompFunc(I0);
-                    ReadCompFunc();
+                ReadWriteCompFunc(I0);
+                ReadCompFunc();
             }
 
             return c_block_tile;
-            }
-        };
-
-        template <typename ADramBlockWindowTmp,
-                  typename BDramBlockWindowTmp,
-                  typename AElementFunction,
-                  typename BElementFunction>
-        CK_TILE_DEVICE auto operator()(const ADramBlockWindowTmp& a_dram_block_window_tmp,
-                                       const AElementFunction& a_element_func,
-                                       const BDramBlockWindowTmp& b_dram_block_window_tmp,
-                                       const BElementFunction& b_element_func,
-                                       index_t num_loop,
-                                       void* p_smem) const
-        {
-            return PipelineImpl<Scheduler>{}.template operator()<HasHotLoop, TailNum>(
-                a_dram_block_window_tmp,
-                a_element_func,
-                b_dram_block_window_tmp,
-                b_element_func,
-                num_loop,
-                p_smem);
-        }
-
-        public:
-        template <typename ADramBlockWindowTmp, typename BDramBlockWindowTmp>
-        CK_TILE_DEVICE auto operator()(const ADramBlockWindowTmp& a_dram_block_window_tmp,
-                                       const BDramBlockWindowTmp& b_dram_block_window_tmp,
-                                       const index_t num_loop,
-                                       void* __restrict__ p_smem) const
-        {
-            return PipelineImpl<Scheduler>{}.template operator()<HasHotLoop, TailNum>(
-                a_dram_block_window_tmp,
-                [](const ADataType& a) { return a; },
-                b_dram_block_window_tmp,
-                [](const BDataType& b) { return b; },
-                num_loop,
-                p_smem);
         }
     };
+
+    template <typename ADramBlockWindowTmp,
+              typename BDramBlockWindowTmp,
+              typename AElementFunction,
+              typename BElementFunction>
+    CK_TILE_DEVICE auto operator()(const ADramBlockWindowTmp& a_dram_block_window_tmp,
+                                   const AElementFunction& a_element_func,
+                                   const BDramBlockWindowTmp& b_dram_block_window_tmp,
+                                   const BElementFunction& b_element_func,
+                                   index_t num_loop,
+                                   void* p_smem) const
+    {
+        return PipelineImpl<Scheduler>{}.template operator()<HasHotLoop, TailNum>(
+            a_dram_block_window_tmp,
+            a_element_func,
+            b_dram_block_window_tmp,
+            b_element_func,
+            num_loop,
+            p_smem);
+    }
+
+    public:
+    template <typename ADramBlockWindowTmp, typename BDramBlockWindowTmp>
+    CK_TILE_DEVICE auto operator()(const ADramBlockWindowTmp& a_dram_block_window_tmp,
+                                   const BDramBlockWindowTmp& b_dram_block_window_tmp,
+                                   const index_t num_loop,
+                                   void* __restrict__ p_smem) const
+    {
+        return PipelineImpl<Scheduler>{}.template operator()<HasHotLoop, TailNum>(
+            a_dram_block_window_tmp,
+            [](const ADataType& a) { return a; },
+            b_dram_block_window_tmp,
+            [](const BDataType& b) { return b; },
+            num_loop,
+            p_smem);
+    }
+};
 } // namespace ck_tile