From 2d642d2737802c8994ffaff3fb12a87c8dad3d49 Mon Sep 17 00:00:00 2001
From: Rostyslav Geyyer <46627076+geyyer@users.noreply.github.com>
Date: Fri, 26 Apr 2024 12:55:45 -0500
Subject: [PATCH] Add element op (#1259)

[ROCm/composable_kernel commit: f044ff71fbc17c479a7c26f522418e2e55a05a56]
---
 .../gpu/element/element_wise_operation.hpp    | 20 +++++++++++++++++++
 1 file changed, 20 insertions(+)

diff --git a/include/ck/tensor_operation/gpu/element/element_wise_operation.hpp b/include/ck/tensor_operation/gpu/element/element_wise_operation.hpp
index 798c5580a2..5499689c9b 100644
--- a/include/ck/tensor_operation/gpu/element/element_wise_operation.hpp
+++ b/include/ck/tensor_operation/gpu/element/element_wise_operation.hpp
@@ -528,6 +528,26 @@ struct UnaryTypeConvert<ck::bhalf_t, float>
     }
 };
 
+struct ConvInvscale
+{
+    /// @brief Op to multiply convolution results by inverted scale factors
+    /// @param e Output after scaling
+    /// @param c Convolution result
+    /// @param d0 Input scale factor
+    /// @param d1 Weights scale factor
+    /// @param d2 Output scale factor
+    template <typename E, typename C, typename D0, typename D1, typename D2>
+    __host__ __device__ void
+    operator()(E& e, const C& c, const D0& d0, const D1& d1, const D2& d2) const;
+
+    template <>
+    __host__ __device__ void operator()<f8_t, float, float, float, float>(
+        f8_t& e, const float& c, const float& d0, const float& d1, const float& d2) const
+    {
+        e = type_convert<f8_t>(c / d0 / d1 / d2);
+    };
+};
+
 } // namespace element_wise
 } // namespace tensor_operation
 } // namespace ck