From aa5ed1a74983026da0ce8ea5aaf385993c243b9b Mon Sep 17 00:00:00 2001
From: Anton Gorenko <anton@streamhpc.com>
Date: Wed, 17 Jun 2026 15:46:44 +0500
Subject: [PATCH] Add operator() overload to
 GemmPipelineAgBgCrCompAsyncEightWaves to use in QuantGemmKernel

---
 ...ipeline_ag_bg_cr_comp_async_eight_waves.hpp | 18 ++++++++++++++++++
 1 file changed, 18 insertions(+)

diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp
index 7c4e42c700..e106bc6a92 100644
--- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp
+++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp
@@ -237,6 +237,24 @@ struct GemmPipelineAgBgCrCompAsyncEightWaves : public BaseGemmPipelineAgBgCrComp
 
         return Base::TailHandler(RunPipeline, has_hot_loop, tail_number);
     }
+
+    template <typename ADramBlockWindowTmp,
+              typename BDramBlockWindowTmp,
+              typename std::enable_if_t<!is_detected<is_tuple, ADramBlockWindowTmp>::value &&
+                                            !is_detected<is_tuple, BDramBlockWindowTmp>::value,
+                                        bool>* = nullptr>
+    CK_TILE_DEVICE auto operator()(const ADramBlockWindowTmp& a_dram_block_window_tmp,
+                                   const BDramBlockWindowTmp& b_dram_block_window_tmp,
+                                   index_t num_loop,
+                                   void* p_smem) const
+    {
+        return operator()(ck_tile::make_tuple(a_dram_block_window_tmp),
+                          ck_tile::element_wise::PassThrough{},
+                          ck_tile::make_tuple(b_dram_block_window_tmp),
+                          ck_tile::element_wise::PassThrough{},
+                          num_loop,
+                          p_smem);
+    }
 };
 
 } // namespace ck_tile