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 - CK_TILE_DEVICE auto operator()(ODramWindowTmp& o_dram_window_tmp, - const OAccTile& o_acc_tile, - const DsDramWindows& /* unused */, - void* = nullptr) + template + 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 // Used for weight-only quantization kernel, B would be dequantized to the same data type as A using BTypeToUse = std::conditional_t, ADataType, BDataType>; - using DsDataType = ck_tile::tuple<>; - using DsLayout = ck_tile::tuple<>; using CLayout = remove_cvref_t; static constexpr index_t kMPerXdl = Problem::kMPerXdl; static constexpr index_t kNPerXdl = Problem::kNPerXdl; @@ -185,8 +181,6 @@ struct DefaultGemm2DEpilogue : public Default2DEpilogue static_assert(false, "Unsupported CLayout!"); } } - - CK_TILE_HOST_DEVICE static constexpr auto GetVectorSizeD() { return 1; } }; } // namespace ck_tile