Extract quantization post operation into another file

include/ck/tensor_operation/gpu/element/element_wise_operation.hpp

@@ -7,6 +7,7 @@
 #include "ck/utility/math_v2.hpp"
 #include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp"
 #include "ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp"
+#include "ck/tensor_operation/gpu/element/quantization_operation.hpp"
 
 namespace ck {
 namespace tensor_operation {

include/ck/tensor_operation/gpu/element/quantization_operation.hpp

+#pragma once
+
+#include "ck/utility/data_type.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace element_wise {
+
+// For Activation function which is piecewise linear function, such as relu, leaky relu ...etc
+template <typename Activation>
+struct Activation_Mul_Clamp
+{
+    Activation_Mul_Clamp(float multiplier, Activation activationOp)
+        : multiplier_(multiplier), activationOp_(activationOp)
+    {
+    }
+
+    __host__ __device__ constexpr void operator()(int8_t& y, const int32_t& x) const
+    {
+        float x_fp32 = ck::type_convert<float>(x);
+        activationOp_(x_fp32, x_fp32);
+        float y_fp32 = math::clamp(multiplier_ * x_fp32, -128.f, 127.f);
+        y            = ck::type_convert<int8_t>(y_fp32);
+    }
+
+    __host__ __device__ constexpr void operator()(float& y, const int32_t& x) const
+    {
+        // We might type_convert to int8 after lambda in someplace
+        float x_fp32 = ck::type_convert<float>(x);
+        activationOp_(x_fp32, x_fp32);
+        y = math::clamp(multiplier_ * x_fp32, -128.f, 127.f);
+    }
+
+    float multiplier_;
+    Activation activationOp_;
+};
+
+// For Activation function which is piecewise linear function, such as relu, leaky relu ...etc
+template <typename Activation>
+struct Add_Activation_Mul_Clamp
+{
+    Add_Activation_Mul_Clamp(float multiplier, Activation activationOp)
+        : multiplier_(multiplier), activationOp_(activationOp)
+    {
+    }
+
+    __host__ __device__ constexpr void
+    operator()(int8_t& y, const int32_t& x1, const int32_t& x2) const
+    {
+        float y_fp32 = ck::type_convert<float>(x1 + x2);
+        activationOp_(y_fp32, y_fp32);
+        y_fp32 = math::clamp(multiplier_ * y_fp32, -128.f, 127.f);
+        y      = ck::type_convert<int8_t>(y_fp32);
+    }
+
+    float multiplier_;
+    Activation activationOp_;
+};
+
+} // namespace element_wise
+} // namespace tensor_operation
+} // namespace ck

include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp

@@ -248,57 +248,6 @@ struct Sigmoid
     int32_t divider_ = 1;
 };
 
-// For Activation function which is piecewise linear function, such as relu, leaky relu ...etc
-template <typename Activation>
-struct Activation_Mul_Clamp
-{
-    Activation_Mul_Clamp(float multiplier, Activation activationOp)
-        : multiplier_(multiplier), activationOp_(activationOp)
-    {
-    }
-
-    __host__ __device__ constexpr void operator()(int8_t& y, const int32_t& x) const
-    {
-        float x_fp32 = ck::type_convert<float>(x);
-        activationOp_(x_fp32, x_fp32);
-        float y_fp32 = math::clamp(multiplier_ * x_fp32, -128.f, 127.f);
-        y            = ck::type_convert<int8_t>(y_fp32);
-    }
-
-    __host__ __device__ constexpr void operator()(float& y, const int32_t& x) const
-    {
-        // We might type_convert to int8 after lambda in someplace
-        float x_fp32 = ck::type_convert<float>(x);
-        activationOp_(x_fp32, x_fp32);
-        y = math::clamp(multiplier_ * x_fp32, -128.f, 127.f);
-    }
-
-    float multiplier_;
-    Activation activationOp_;
-};
-
-// For Activation function which is piecewise linear function, such as relu, leaky relu ...etc
-template <typename Activation>
-struct Add_Activation_Mul_Clamp
-{
-    Add_Activation_Mul_Clamp(float multiplier, Activation activationOp)
-        : multiplier_(multiplier), activationOp_(activationOp)
-    {
-    }
-
-    __host__ __device__ constexpr void
-    operator()(int8_t& y, const int32_t& x1, const int32_t& x2) const
-    {
-        float y_fp32 = ck::type_convert<float>(x1 + x2);
-        activationOp_(y_fp32, y_fp32);
-        y_fp32 = math::clamp(multiplier_ * y_fp32, -128.f, 127.f);
-        y      = ck::type_convert<int8_t>(y_fp32);
-    }
-
-    float multiplier_;
-    Activation activationOp_;
-};
-
 } // namespace element_wise
 } // namespace tensor_operation
 } // namespace ck