From d175d5b28ac7e258a8a52351cc42da130d1624a9 Mon Sep 17 00:00:00 2001
From: Po Yen Chen <PoYen.Chen@amd.com>
Date: Wed, 3 Sep 2025 00:51:56 +0800
Subject: [PATCH] [CK_TILE] Fix fmha_fwd_v3() Default2DEpilogue usage (#2765)

* Fix Default2DEpilogue usage

* Fix Default2DEpilogue usage for batch_prefill

[ROCm/composable_kernel commit: 9f35cde374381ba76ea793d0794ac31ced075bb0]
---
 include/ck_tile/ops/fmha/kernel/fmha_batch_prefill_kernel.hpp | 2 +-
 include/ck_tile/ops/fmha/kernel/fmha_fwd_v3_kernel.hpp        | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/include/ck_tile/ops/fmha/kernel/fmha_batch_prefill_kernel.hpp b/include/ck_tile/ops/fmha/kernel/fmha_batch_prefill_kernel.hpp
index 2850ce3379..fcd512056d 100644
--- a/include/ck_tile/ops/fmha/kernel/fmha_batch_prefill_kernel.hpp
+++ b/include/ck_tile/ops/fmha/kernel/fmha_batch_prefill_kernel.hpp
@@ -1143,7 +1143,7 @@ struct FmhaBatchPrefillWithPagedKVCacheKernel
                              make_tuple(number<FmhaPipeline::kM0>{}, number<FmhaPipeline::kN1>{}),
                              {i_m0, i_n1});
 
-        EpiloguePipeline{}(o_dram_window, o_acc_tile);
+        EpiloguePipeline{}(o_dram_window, o_acc_tile, nullptr);
     }
 };
 
diff --git a/include/ck_tile/ops/fmha/kernel/fmha_fwd_v3_kernel.hpp b/include/ck_tile/ops/fmha/kernel/fmha_fwd_v3_kernel.hpp
index be14a36353..87021354aa 100644
--- a/include/ck_tile/ops/fmha/kernel/fmha_fwd_v3_kernel.hpp
+++ b/include/ck_tile/ops/fmha/kernel/fmha_fwd_v3_kernel.hpp
@@ -513,7 +513,7 @@ struct FmhaFwdV3Kernel
                              make_tuple(number<FmhaPipeline::kM0>{}, number<FmhaPipeline::kN1>{}),
                              {i_m0, i_n1});
 
-        EpiloguePipeline{}(o_dram_window, o_acc_tile);
+        EpiloguePipeline{}(o_dram_window, o_acc_tile, nullptr);
     }
 };
 } // namespace ck_tile