From f85c70b31e276bbbcb63c13d9c30f55174afbe61 Mon Sep 17 00:00:00 2001
From: Thomas Ning <Thomas.Ning@amd.com>
Date: Tue, 17 Jun 2025 22:43:05 -0700
Subject: [PATCH] Revert "Fix default epilogue  (#2358)" (#2364)

This reverts commit b29e3830a6b7676cac3610ace84fcc8589c993a7.

[ROCm/composable_kernel commit: 64a2fda713a7723e63562f4be80f0cc123baa724]
---
 include/ck_tile/ops/epilogue/default_2d_epilogue.hpp | 12 +++---------
 1 file changed, 3 insertions(+), 9 deletions(-)

diff --git a/include/ck_tile/ops/epilogue/default_2d_epilogue.hpp b/include/ck_tile/ops/epilogue/default_2d_epilogue.hpp
index 623433c1dc..ab3c0df88d 100644
--- a/include/ck_tile/ops/epilogue/default_2d_epilogue.hpp
+++ b/include/ck_tile/ops/epilogue/default_2d_epilogue.hpp
@@ -71,11 +71,9 @@ struct Default2DEpilogue
 
     // TODO: this function assume store out vector size is the same as OAccTile last dimension size
     //       how do we fix this ?
-    template <typename ODramWindowTmp, typename OAccTile, typename DsDramWindows>
-    CK_TILE_DEVICE auto operator()(ODramWindowTmp& o_dram_window_tmp,
-                                   const OAccTile& o_acc_tile,
-                                   const DsDramWindows& /* unused */,
-                                   void* = nullptr)
+    template <typename ODramWindowTmp, typename OAccTile>
+    CK_TILE_DEVICE auto
+    operator()(ODramWindowTmp& o_dram_window_tmp, const OAccTile& o_acc_tile, void* = nullptr)
     {
 
         // TODO: this is ugly
@@ -116,8 +114,6 @@ struct DefaultGemm2DEpilogue : public Default2DEpilogue<Problem_, Policy_>
     // Used for weight-only quantization kernel, B would be dequantized to the same data type as A
     using BTypeToUse =
         std::conditional_t<std::is_same_v<BDataType, pk_int4_t>, ADataType, BDataType>;
-    using DsDataType                       = ck_tile::tuple<>;
-    using DsLayout                         = ck_tile::tuple<>;
     using CLayout                          = remove_cvref_t<typename Problem::CLayout>;
     static constexpr index_t kMPerXdl      = Problem::kMPerXdl;
     static constexpr index_t kNPerXdl      = Problem::kNPerXdl;
@@ -185,8 +181,6 @@ struct DefaultGemm2DEpilogue : public Default2DEpilogue<Problem_, Policy_>
             static_assert(false, "Unsupported CLayout!");
         }
     }
-
-    CK_TILE_HOST_DEVICE static constexpr auto GetVectorSizeD() { return 1; }
 };
 
 } // namespace ck_tile