PaddlePaddle
diff --git a/‎paddle/phi/api/yaml/static_ops.yaml
Copy file name to clipboard
+9Lines changed: 9 additions & 0 deletions b/‎paddle/phi/api/yaml/static_ops.yaml
Copy file name to clipboard
+9Lines changed: 9 additions & 0 deletions
diff --git a/‎paddle/phi/infermeta/ternary.cc
Copy file name to clipboard
+112Lines changed: 112 additions & 0 deletions b/‎paddle/phi/infermeta/ternary.cc
Copy file name to clipboard
+112Lines changed: 112 additions & 0 deletions
diff --git a/‎paddle/phi/infermeta/ternary.h
Copy file name to clipboard
+13Lines changed: 13 additions & 0 deletions b/‎paddle/phi/infermeta/ternary.h
Copy file name to clipboard
+13Lines changed: 13 additions & 0 deletions
diff --git a/‎paddle/phi/kernels/funcs/fc_functor.cu
Copy file name to clipboard
+87Lines changed: 87 additions & 0 deletions b/‎paddle/phi/kernels/funcs/fc_functor.cu
Copy file name to clipboard
+87Lines changed: 87 additions & 0 deletions
diff --git a/‎paddle/phi/kernels/funcs/fc_functor.h
Copy file name to clipboard
+21Lines changed: 21 additions & 0 deletions b/‎paddle/phi/kernels/funcs/fc_functor.h
Copy file name to clipboard
+21Lines changed: 21 additions & 0 deletions
@@ -500,6 +500,15 @@
     data_type : x
   backward : prod_grad
 
+- op : quant_linear
+  args: (Tensor x, Tensor w, Tensor bias, int in_num_col_dims = 1, str activation_type = "", bool padding_weights = false, float scale_in = 1.0f, float[] scale_weights = {1.0f}, int quant_round_type = 1, float quant_max_bound = 127.0f, float quant_min_bound = -127.0f)
+  output: Tensor(out)
+  optional: bias
+  infer_meta:
+    func: QuantLinearInferMeta
+  kernel:
+    func: quant_linear
+
 - op : randint
   args : (int low, int high, IntArray shape = {}, DataType dtype = DataType::INT64, int seed = 0)
   output : Tensor(out)
 
@@ -1382,4 +1382,116 @@ void ViterbiDecodeInferMeta(const MetaTensor& input,
   scores->set_dtype(length.dtype());
 }
 
+void QuantLinearInferMeta(const MetaTensor& x,
+                          const MetaTensor& w,
+                          const MetaTensor& bias,
+                          int in_num_col_dims,
+                          const std::string& activation_type,
+                          bool padding_weights,
+                          float scale_in,
+                          const std::vector<float>& scale_weights,
+                          int quant_round_type,
+                          float quant_max_bound,
+                          float quant_min_bound,
+                          MetaTensor* y) {
+  auto w_dims = w.dims();
+  PADDLE_ENFORCE_EQ(
+      w_dims.size(),
+      2,
+      phi::errors::InvalidArgument(
+          "The input Weight of quant_linear is expected to be a 2-D tensor. "
+          "But received the number of Weight's dimensions is %d, "
+          "Weight's shape is %s.",
+          w_dims.size(),
+          w_dims));
+  if (bias) {
+    auto bias_dims = bias.dims();
+    auto w_dims1 = padding_weights ? w_dims[1] - 4 : w_dims[1];
+
+    PADDLE_ENFORCE_LE(bias_dims.size(),
+                      2,
+                      phi::errors::InvalidArgument(
+                          "The input Bias of quant_linear is expected to be a "
+                          "1-D or 2-D tensor. But "
+                          "received the number of Bias's dimensions is %d, "
+                          "Bias's shape is %s.",
+                          bias_dims.size(),
+                          bias_dims));
+
+    PADDLE_ENFORCE_EQ(
+        bias_dims[bias_dims.size() - 1],
+        w_dims1,
+        phi::errors::InvalidArgument(
+            "The last dimension of input Bias is expected be equal "
+            "to the actual width of input Weight. But received the last "
+            "dimension of Bias is %d, Bias's shape is %s; "
+            "the actual width of Weight is %d, Weight's shape is %s.",
+            bias_dims[bias_dims.size() - 1],
+            bias_dims,
+            w_dims1,
+            w_dims));
+
+    if (bias_dims.size() == 2) {
+      PADDLE_ENFORCE_EQ(
+          bias_dims[0],
+          1,
+          phi::errors::InvalidArgument(
+              "The first dimension of input Bias is expected to be 1, "
+              "but received %d, Bias's shape is %s.",
+              bias_dims[0],
+              bias_dims));
+    }
+  }
+
+  auto in_dims = x.dims();
+  PADDLE_ENFORCE_LT(
+      in_num_col_dims,
+      in_dims.size(),
+      phi::errors::InvalidArgument(
+          "The attribute in_num_col_dims used to flatten Input to "
+          "a 2-D tensor, is expected to be less than the number of "
+          "Input's dimensions. But received in_num_col_dims is %d, "
+          "the number of Input's dimensions is %d, Input's shape is %s.",
+          in_num_col_dims,
+          in_dims.size(),
+          in_dims));
+
+  if (!activation_type.empty()) {
+    PADDLE_ENFORCE_EQ(
+        activation_type,
+        "relu",
+        phi::errors::InvalidArgument(
+            "The attribute activation_type of quant_linear is expected "
+            "to be \"relu\", but received %s.",
+            activation_type.c_str()));
+  }
+
+  std::vector<int64_t> output_dims;
+
+  auto in_mat_dims = phi::flatten_to_2d(in_dims, in_num_col_dims);
+  auto w_dims0 = padding_weights ? w_dims[0] - 4 : w_dims[0];
+  auto w_dims1 = padding_weights ? w_dims[1] - 4 : w_dims[1];
+  PADDLE_ENFORCE_EQ(
+      in_mat_dims[1],
+      w_dims0,
+      phi::errors::InvalidArgument(
+          "The input's second dimension and weight's first dimension is "
+          "expected to be the same. But received input's second dimension is "
+          "%d, input's shape is %s; weight's first dimension is %d, weight's "
+          "shape is %s.",
+          in_mat_dims[1],
+          in_mat_dims,
+          w_dims0,
+          phi::make_ddim({w_dims0, w_dims1})));
+  output_dims.reserve(static_cast<size_t>(in_num_col_dims + 1));
+  for (int i = 0; i < in_num_col_dims; ++i) {
+    output_dims.push_back(in_dims[i]);
+  }
+  output_dims.push_back(w_dims1);
+
+  y->set_dims(make_ddim(output_dims));
+  y->share_lod(x);
+  y->set_dtype(x.dtype());
+}
+
 }  // namespace phi
@@ -219,4 +219,17 @@ void ViterbiDecodeInferMeta(const MetaTensor& input,
                             MetaTensor* path,
                             MetaConfig config = MetaConfig());
 
+void QuantLinearInferMeta(const MetaTensor& x,
+                          const MetaTensor& w,
+                          const MetaTensor& bias,
+                          int in_num_col_dims,
+                          const std::string& activation_type,
+                          bool padding_weights,
+                          float scale_in,
+                          const std::vector<float>& scale_weights,
+                          int quant_round_type,
+                          float quant_max_bound,
+                          float quant_min_bound,
+                          MetaTensor* y);
+
 }  // namespace phi
@@ -19,6 +19,12 @@ limitations under the License. */
 #include "paddle/phi/kernels/funcs/blas/blas.h"
 #include "paddle/phi/kernels/funcs/fc_functor.h"
 
+#include "paddle/phi/backends/gpu/gpu_launch_config.h"
+#include "paddle/phi/core/dense_tensor.h"
+#include "paddle/phi/kernels/funcs/blas/blaslt_impl.cu.h"
+#include "paddle/phi/kernels/funcs/quant_dequant.h"
+#include "paddle/phi/kernels/matmul_kernel.h"
+
 namespace phi {
 namespace funcs {
 
@@ -367,5 +373,86 @@ template class FCFunctor<GPUContext, float16>;
 template class FCFunctor<GPUContext, float>;
 template class FCFunctor<GPUContext, double>;
 
+template <typename DeviceContext, typename T>
+void FCInt8Functor<DeviceContext, T>::operator()(
+    const DeviceContext& context,
+    const int M,
+    const int N,
+    const int K,
+    const T* X,
+    const DenseTensor* w_tensor,
+    T* Y,
+    float scale_in,
+    std::vector<float> scale_weights,
+    int quant_round_type,
+    float quant_max_bound,
+    float quant_min_bound,
+    const T* B,
+    bool relu,
+    bool padding_weights) {
+  PADDLE_ENFORCE_EQ(padding_weights,
+                    false,
+                    errors::PermissionDenied(
+                        "Weight padding in fc can not be used in GPU scope."));
+  const int8_t* W = w_tensor->data<int8_t>();
+
+  DenseTensor quant_x_tensor, quant_y_tensor;
+  quant_x_tensor.Resize(phi::make_ddim({M, K}));
+  quant_y_tensor.Resize(phi::make_ddim({M, N}));
+  context.template Alloc<int8_t>(&quant_x_tensor,
+                                 quant_x_tensor.numel() * sizeof(int8_t));
+  context.template Alloc<int32_t>(&quant_y_tensor,
+                                  quant_y_tensor.numel() * sizeof(int32_t));
+  LaunchQuantKernelWithVecSize<T>(X,
+                                  quant_x_tensor.data<int8_t>(),
+                                  scale_in,
+                                  M,
+                                  K,
+                                  quant_round_type,
+                                  quant_max_bound,
+                                  quant_min_bound,
+                                  context.stream());
+
+  MatmulKernel<int8_t, GPUContext>(
+      context, quant_x_tensor, *w_tensor, false, false, &quant_y_tensor);
+
+  DenseTensor scale_weights_dev;
+  scale_weights_dev.Resize(phi::make_ddim({N}));
+  context.template Alloc<float>(&scale_weights_dev,
+                                scale_weights_dev.numel() * sizeof(float));
+  float* scale_weights_dev_ptr = scale_weights_dev.data<float>();
+  cudaMemcpyAsync(scale_weights_dev_ptr,
+                  scale_weights.data(),
+                  N * sizeof(float),
+                  cudaMemcpyHostToDevice);
+
+  phi::backends::gpu::GpuLaunchConfig config;
+  if (N % DequantKernelVecSize == 0) {
+    config = phi::backends::gpu::GetGpuLaunchConfig1D(
+        context, M * N, DequantKernelVecSize);
+  } else {
+    config = phi::backends::gpu::GetGpuLaunchConfig1D(context, M * N, 1);
+  }
+  LaunchDequantKernelWithScaleOfInputAndWeight(quant_y_tensor.data<int32_t>(),
+                                               Y,
+                                               M,
+                                               N,
+                                               context.stream(),
+                                               &config,
+                                               scale_in,
+                                               scale_weights_dev_ptr,
+                                               quant_max_bound);
+
+  if (B == NULL) {
+    return;
+  }
+
+  // M * N
+  AddReluKernel(context.stream(), M, N, Y, B, relu);
+}
+
+template class FCInt8Functor<GPUContext, float16>;
+template class FCInt8Functor<GPUContext, float>;
+template class FCInt8Functor<GPUContext, double>;
 }  // namespace funcs
 }  // namespace phi
@@ -17,6 +17,7 @@ limitations under the License. */
 #include <string>
 
 #include "paddle/phi/backends/all_context.h"
+#include "paddle/phi/core/dense_tensor.h"
 
 namespace phi {
 namespace funcs {
@@ -36,5 +37,25 @@ class FCFunctor {
                   bool weight_pass = false);
 };
 
+template <typename DeviceContext, typename T>
+class FCInt8Functor {
+ public:
+  void operator()(const DeviceContext& context,
+                  const int M,
+                  const int N,
+                  const int K,
+                  const T* X,
+                  const DenseTensor* W,
+                  T* Y,
+                  float scale_in,
+                  std::vector<float> scale_weights,
+                  int quant_round_type,
+                  float quant_max_bound,
+                  float quant_min_bound,
+                  const T* B = nullptr,
+                  bool relu = false,
+                  bool weight_pass = false);
+};
+
 }  // namespace funcs
 }  // namespace phi