PaddlePaddle
diff --git a/‎paddle/fluid/operators/fake_dequantize_op.cc
Copy file name to clipboard
-71Lines changed: 0 additions & 71 deletions b/‎paddle/fluid/operators/fake_dequantize_op.cc
Copy file name to clipboard
-71Lines changed: 0 additions & 71 deletions
diff --git a/‎paddle/fluid/operators/fake_dequantize_op.cu
Copy file name to clipboard
-7Lines changed: 0 additions & 7 deletions b/‎paddle/fluid/operators/fake_dequantize_op.cu
Copy file name to clipboard
-7Lines changed: 0 additions & 7 deletions
diff --git a/‎paddle/fluid/operators/fake_dequantize_op.cu.h
Copy file name to clipboard
-36Lines changed: 0 additions & 36 deletions b/‎paddle/fluid/operators/fake_dequantize_op.cu.h
Copy file name to clipboard
-36Lines changed: 0 additions & 36 deletions
diff --git a/‎paddle/fluid/operators/fake_dequantize_op.h
Copy file name to clipboard
-27Lines changed: 0 additions & 27 deletions b/‎paddle/fluid/operators/fake_dequantize_op.h
Copy file name to clipboard
-27Lines changed: 0 additions & 27 deletions
diff --git a/‎paddle/phi/infermeta/binary.cc
Copy file name to clipboard
+9Lines changed: 9 additions & 0 deletions b/‎paddle/phi/infermeta/binary.cc
Copy file name to clipboard
+9Lines changed: 9 additions & 0 deletions
diff --git a/‎paddle/phi/infermeta/binary.h
Copy file name to clipboard
+5Lines changed: 5 additions & 0 deletions b/‎paddle/phi/infermeta/binary.h
Copy file name to clipboard
+5Lines changed: 5 additions & 0 deletions
diff --git a/‎paddle/phi/kernels/cpu/fake_dequantize_kernel.cc
Copy file name to clipboard
+23Lines changed: 23 additions & 0 deletions b/‎paddle/phi/kernels/cpu/fake_dequantize_kernel.cc
Copy file name to clipboard
+23Lines changed: 23 additions & 0 deletions
diff --git a/‎paddle/phi/kernels/fake_dequantize_kernel.h
Copy file name to clipboard
+28Lines changed: 28 additions & 0 deletions b/‎paddle/phi/kernels/fake_dequantize_kernel.h
Copy file name to clipboard
+28Lines changed: 28 additions & 0 deletions
diff --git a/‎paddle/phi/kernels/funcs/fake_dequantize_functor.cc
Copy file name to clipboard
+38Lines changed: 38 additions & 0 deletions b/‎paddle/phi/kernels/funcs/fake_dequantize_functor.cc
Copy file name to clipboard
+38Lines changed: 38 additions & 0 deletions
diff --git a/‎paddle/phi/kernels/funcs/fake_dequantize_functor.cu
Copy file name to clipboard
+57Lines changed: 57 additions & 0 deletions b/‎paddle/phi/kernels/funcs/fake_dequantize_functor.cu
Copy file name to clipboard
+57Lines changed: 57 additions & 0 deletions
@@ -22,22 +22,6 @@ limitations under the License. */
 namespace paddle {
 namespace operators {
 
-template <typename T>
-struct DequantizeFunctor<phi::CPUContext, T> {
-  void operator()(const phi::CPUContext& dev_ctx,
-                  const phi::DenseTensor* in,
-                  const phi::DenseTensor* scale,
-                  T max_range,
-                  phi::DenseTensor* out) {
-    auto in_e = phi::EigenVector<T>::Flatten(*in);
-    const T* scale_factor = scale->data<T>();
-    auto out_e = phi::EigenVector<T>::Flatten(*out);
-
-    auto& dev = *dev_ctx.eigen_device();
-    out_e.device(dev) = in_e * scale_factor[0] / max_range;
-  }
-};
-
 template <typename T>
 struct ChannelDequantizeFunctor<phi::CPUContext, T> {
   void operator()(const phi::CPUContext& dev_ctx,
@@ -139,51 +123,9 @@ struct ChannelDequantizeFunctor<phi::CPUContext, T> {
   }
 };
 
-template struct DequantizeFunctor<phi::CPUContext, float>;
-template struct DequantizeFunctor<phi::CPUContext, double>;
 template struct ChannelDequantizeFunctor<phi::CPUContext, float>;
 template struct ChannelDequantizeFunctor<phi::CPUContext, double>;
 
-class FakeDequantizeMaxAbsOp : public framework::OperatorWithKernel {
- public:
-  FakeDequantizeMaxAbsOp(const std::string& type,
-                         const framework::VariableNameMap& inputs,
-                         const framework::VariableNameMap& outputs,
-                         const framework::AttributeMap& attrs)
-      : OperatorWithKernel(type, inputs, outputs, attrs) {}
-
-  void InferShape(framework::InferShapeContext* ctx) const override {
-    OP_INOUT_CHECK(ctx->HasInput("X"), "Input", "X", "FakeDequantizeMaxAbs");
-    OP_INOUT_CHECK(
-        ctx->HasOutput("Out"), "Output", "Out", "FakeDequantizeMaxAbs");
-
-    ctx->ShareDim("X", /*->*/ "Out");
-    ctx->ShareLoD("X", /*->*/ "Out");
-  }
-};
-
-class FakeDequantizeMaxAbsOpMaker : public framework::OpProtoAndCheckerMaker {
- public:
-  void Make() override {
-    AddInput("X",
-             "(Tensor) The input with float-32/64 type is the "
-             "low precision tensor.");
-    AddInput("Scale", "(float) The scale in quantization stage.");
-    AddOutput("Out",
-              "(Tensor) The output is the dequantized high "
-              "precision tensor.");
-    AddAttr<float>("max_range", "(float) The max range in quantization stage.");
-    AddComment(R"DOC(
-FakeDequantizeMaxAbsOp operator.
-
-This calculation is an opposite operation of FakeQuantizeMaxAbsOp:
-
-$$Out = \frac{scale*X}{ max_range }$$
-
-)DOC");
-  }
-};
-
 class FakeChannelWiseDequantizeMaxAbsOp : public framework::OperatorWithKernel {
  public:
   using framework::OperatorWithKernel::OperatorWithKernel;
@@ -271,19 +213,6 @@ Notes: In general, the per-channel quantization is only applied to weights and t
 namespace ops = paddle::operators;
 using CPU = phi::CPUContext;
 
-REGISTER_OPERATOR(
-    fake_dequantize_max_abs,
-    ops::FakeDequantizeMaxAbsOp,
-    ops::FakeDequantizeMaxAbsOpMaker,
-    paddle::framework::EmptyGradOpMaker<paddle::framework::OpDesc>,
-    paddle::framework::EmptyGradOpMaker<paddle::imperative::OpBase>);
-PD_REGISTER_STRUCT_KERNEL(fake_dequantize_max_abs,
-                          CPU,
-                          ALL_LAYOUT,
-                          ops::FakeDequantizeMaxAbsKernel,
-                          float,
-                          double) {}
-
 REGISTER_OPERATOR(
     fake_channel_wise_dequantize_max_abs,
     ops::FakeChannelWiseDequantizeMaxAbsOp,
 
@@ -18,13 +18,6 @@ limitations under the License. */
 namespace ops = paddle::operators;
 using float16 = phi::dtype::float16;
 
-PD_REGISTER_STRUCT_KERNEL(fake_dequantize_max_abs,
-                          GPU,
-                          ALL_LAYOUT,
-                          ops::FakeDequantizeMaxAbsKernel,
-                          float,
-                          double,
-                          float16) {}
 PD_REGISTER_STRUCT_KERNEL(fake_channel_wise_dequantize_max_abs,
                           GPU,
                           ALL_LAYOUT,
 
@@ -21,40 +21,6 @@ limitations under the License. */
 namespace paddle {
 namespace operators {
 
-template <typename T>
-__global__ void KeDequantize(
-    const T* in, const T* scale, T max_range, int64_t num, T* out) {
-  int64_t idx = threadIdx.x + blockIdx.x * blockDim.x;
-  for (int64_t i = idx; i < num; i += blockDim.x * gridDim.x) {
-    out[i] = in[i] * scale[0] / max_range;
-  }
-}
-
-template <typename T>
-struct DequantizeFunctor<phi::GPUContext, T> {
-  void operator()(const phi::GPUContext& dev_ctx,
-                  const phi::DenseTensor* in,
-                  const phi::DenseTensor* scale,
-                  T max_range,
-                  phi::DenseTensor* out) {
-    const T* in_data = in->data<T>();
-    const T* scale_factor = scale->data<T>();
-    T* out_data = out->mutable_data<T>(dev_ctx.GetPlace());
-
-    int64_t num = in->numel();
-    int64_t block_size = std::min(
-        num, static_cast<int64_t>(dev_ctx.GetMaxThreadsPerBlock() / 4));
-    int64_t max_threads =
-        dev_ctx.GetMaxPhysicalThreadCount();  // SM * block_per_SM
-    const int64_t max_blocks =
-        std::max(((max_threads - 1) / block_size + 1), static_cast<int64_t>(1));
-    const int64_t grid_size =
-        std::min(max_blocks, (num + block_size - 1) / block_size);
-    KeDequantize<T><<<grid_size, block_size, 0, dev_ctx.stream()>>>(
-        in_data, scale_factor, max_range, num, out_data);
-  }
-};
-
 template <typename T>
 __global__ void DequantizeOneScaleQuantAxis0(
     const T* in, const T* scale, T max_range, int num, int channel, T* out) {
@@ -177,8 +143,6 @@ struct ChannelDequantizeFunctor<phi::GPUContext, T> {
   }
 };
 
-template struct DequantizeFunctor<phi::GPUContext, float>;
-template struct DequantizeFunctor<phi::GPUContext, double>;
 template struct ChannelDequantizeFunctor<phi::GPUContext, float>;
 template struct ChannelDequantizeFunctor<phi::GPUContext, double>;
 
 
@@ -23,15 +23,6 @@ limitations under the License. */
 namespace paddle {
 namespace operators {
 
-template <typename DeviceContext, typename T>
-struct DequantizeFunctor {
-  void operator()(const DeviceContext& dev_ctx,
-                  const phi::DenseTensor* in,
-                  const phi::DenseTensor* scale,
-                  T max_range,
-                  phi::DenseTensor* out);
-};
-
 template <typename DeviceContext, typename T>
 struct ChannelDequantizeFunctor {
   void operator()(const DeviceContext& dev_ctx,
@@ -44,24 +35,6 @@ struct ChannelDequantizeFunctor {
                   phi::DenseTensor* out);
 };
 
-template <typename T, typename DeviceContext>
-class FakeDequantizeMaxAbsKernel : public framework::OpKernel<T> {
- public:
-  virtual void Compute(const framework::ExecutionContext& ctx) const {
-    auto* in = ctx.Input<phi::DenseTensor>("X");
-    auto* scale = ctx.Input<phi::DenseTensor>("Scale");
-    auto* out = ctx.Output<phi::DenseTensor>("Out");
-
-    float max_range = ctx.Attr<float>("max_range");
-
-    auto& dev_ctx = ctx.template device_context<DeviceContext>();
-    out->mutable_data<T>(dev_ctx.GetPlace());
-
-    DequantizeFunctor<DeviceContext, T>()(
-        dev_ctx, in, scale, static_cast<T>(max_range), out);
-  }
-};
-
 template <typename T, typename DeviceContext>
 class FakeChannelWiseDequantizeMaxAbsKernel : public framework::OpKernel<T> {
  public:
 
@@ -1674,6 +1674,15 @@ void ExpandAsInferMeta(const MetaTensor& x,
 #undef MAX_RANK_SUPPORTED
 }
 
+void FakeDequantizeMaxAbsInferMeta(const MetaTensor& x,
+                                   const MetaTensor& scale,
+                                   float max_range,
+                                   MetaTensor* out) {
+  out->set_dtype(x.dtype());
+  out->share_dims(x);
+  out->share_lod(x);
+}
+
 void FillDiagonalTensorInferMeta(const MetaTensor& x,
                                  const MetaTensor& y,
                                  int64_t offset,
 
@@ -324,6 +324,11 @@ void ExpandAsInferMeta(const MetaTensor& x,
                        const std::vector<int>& target_shape,
                        MetaTensor* out);
 
+void FakeDequantizeMaxAbsInferMeta(const MetaTensor& x,
+                                   const MetaTensor& scale,
+                                   float max_range,
+                                   MetaTensor* out);
+
 void FillDiagonalTensorInferMeta(const MetaTensor& x,
                                  const MetaTensor& y,
                                  int64_t offset,
 
@@ -0,0 +1,23 @@
+// Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "paddle/phi/core/kernel_registry.h"
+#include "paddle/phi/kernels/impl/fake_dequantize_kernel_impl.h"
+
+PD_REGISTER_KERNEL(fake_dequantize_max_abs,
+                   CPU,
+                   ALL_LAYOUT,
+                   phi::FakeDequantizeMaxAbsKernel,
+                   float,
+                   double) {}
@@ -0,0 +1,28 @@
+// Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#pragma once
+
+#include "paddle/phi/core/dense_tensor.h"
+
+namespace phi {
+
+template <typename T, typename Context>
+void FakeDequantizeMaxAbsKernel(const Context& dev_ctx,
+                                const DenseTensor& x,
+                                const DenseTensor& scale,
+                                float max_range,
+                                DenseTensor* out);
+
+}  // namespace phi
@@ -0,0 +1,38 @@
+/* Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    https://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/phi/kernels/funcs/fake_dequantize_functor.h"
+
+namespace phi {
+namespace funcs {
+
+template <typename Context, typename T>
+void DequantizeFunctor<Context, T>::operator()(const Context& dev_ctx,
+                                               const DenseTensor* in,
+                                               const DenseTensor* scale,
+                                               T max_range,
+                                               DenseTensor* out) {
+  auto in_e = phi::EigenVector<T>::Flatten(*in);
+  const T* scale_factor = scale->data<T>();
+  auto out_e = phi::EigenVector<T>::Flatten(*out);
+
+  auto& dev = *dev_ctx.eigen_device();
+  out_e.device(dev) = in_e * scale_factor[0] / max_range;
+}
+
+template class DequantizeFunctor<CPUContext, float>;
+template class DequantizeFunctor<CPUContext, double>;
+
+}  // namespace funcs
+}  // namespace phi
@@ -0,0 +1,57 @@
+/* Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    https://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/phi/kernels/funcs/fake_dequantize_functor.h"
+
+namespace phi {
+namespace funcs {
+
+template <typename T>
+__global__ void KeDequantize(
+    const T* in, const T* scale, T max_range, int64_t num, T* out) {
+  int64_t idx = threadIdx.x + blockIdx.x * blockDim.x;
+  for (int64_t i = idx; i < num; i += blockDim.x * gridDim.x) {
+    out[i] = in[i] * scale[0] / max_range;
+  }
+}
+
+template <typename Context, typename T>
+void DequantizeFunctor<Context, T>::operator()(const Context& dev_ctx,
+                                               const DenseTensor* in,
+                                               const DenseTensor* scale,
+                                               T max_range,
+                                               DenseTensor* out) {
+  const T* in_data = in->data<T>();
+  const T* scale_factor = scale->data<T>();
+  T* out_data = dev_ctx.template Alloc<T>(out);
+
+  int64_t num = in->numel();
+  int64_t block_size =
+      std::min(num, static_cast<int64_t>(dev_ctx.GetMaxThreadsPerBlock() / 4));
+  int64_t max_threads =
+      dev_ctx.GetMaxPhysicalThreadCount();  // SM * block_per_SM
+  const int64_t max_blocks =
+      std::max(((max_threads - 1) / block_size + 1), static_cast<int64_t>(1));
+  const int64_t grid_size =
+      std::min(max_blocks, (num + block_size - 1) / block_size);
+  KeDequantize<T><<<grid_size, block_size, 0, dev_ctx.stream()>>>(
+      in_data, scale_factor, max_range, num, out_data);
+}
+
+template class DequantizeFunctor<GPUContext, float>;
+template class DequantizeFunctor<GPUContext, double>;
+template class DequantizeFunctor<GPUContext, float16>;
+
+}  // namespace funcs
+}  // namespace phi