diff --git a/paddle/fluid/operators/abs_op.cc b/paddle/fluid/operators/abs_op.cc
index 5a09933e0ee24889f5192a8c84e449e09bdc147e..149a87fe32da16e850d5d64fb519c9bde7afef62 100644
--- a/paddle/fluid/operators/abs_op.cc
+++ b/paddle/fluid/operators/abs_op.cc
@@ -12,12 +12,11 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "paddle/fluid/operators/abs_op.h"
-
 #include <memory>
 #include <string>
 #include <unordered_map>
 #include <vector>
+#include "paddle/fluid/framework/op_registry.h"
 #ifdef PADDLE_WITH_MKLDNN
 #include "paddle/fluid/platform/mkldnn_helper.h"
 #endif
@@ -108,7 +107,7 @@ class AbsDoubleGradMaker : public framework::SingleGradOpMaker<T> {
 
  protected:
   void Apply(GradOpPtr<T> op) const override {
-    op->SetType("abs_grad_grad");
+    op->SetType("abs_double_grad");
     // input1: x
     op->SetInput("X", this->Input("X"));
     // input2: ddx
@@ -159,37 +158,4 @@ REGISTER_OPERATOR(abs_grad, ops::AbsGradOp,
                   ops::AbsDoubleGradMaker<paddle::framework::OpDesc>,
                   ops::AbsDoubleGradMaker<paddle::imperative::OpBase>);
 
-REGISTER_OPERATOR(abs_grad_grad, ops::AbsDoubleGradOp);
-
-REGISTER_OP_CPU_KERNEL(
-    abs, ops::AbsKernel<paddle::platform::CPUDeviceContext, float>,
-    ops::AbsKernel<paddle::platform::CPUDeviceContext, double>,
-    ops::AbsKernel<paddle::platform::CPUDeviceContext, int>,
-    ops::AbsKernel<paddle::platform::CPUDeviceContext, int64_t>,
-    ops::AbsKernel<paddle::platform::CPUDeviceContext,
-                   paddle::platform::complex<float>>,
-    ops::AbsKernel<paddle::platform::CPUDeviceContext,
-                   paddle::platform::complex<double>>);
-
-REGISTER_OP_CPU_KERNEL(
-    abs_grad, ops::AbsGradKernel<paddle::platform::CPUDeviceContext, float>,
-    ops::AbsGradKernel<paddle::platform::CPUDeviceContext, double>,
-    ops::AbsGradKernel<paddle::platform::CPUDeviceContext, int>,
-    ops::AbsGradKernel<paddle::platform::CPUDeviceContext, int64_t>,
-    ops::AbsGradKernel<paddle::platform::CPUDeviceContext,
-                       paddle::platform::complex<float>>,
-    ops::AbsGradKernel<paddle::platform::CPUDeviceContext,
-                       paddle::platform::complex<double>>);
-
-REGISTER_OP_CPU_KERNEL(
-    abs_grad_grad,
-    ops::AbsDoubleGradKernel<paddle::platform::CPUDeviceContext, float>,
-    ops::AbsDoubleGradKernel<paddle::platform::CPUDeviceContext, double>,
-    ops::AbsDoubleGradKernel<paddle::platform::CPUDeviceContext, int>,
-    ops::AbsDoubleGradKernel<paddle::platform::CPUDeviceContext, int64_t>,
-    ops::AbsDoubleGradKernel<paddle::platform::CPUDeviceContext,
-                             paddle::platform::float16>,
-    ops::AbsDoubleGradKernel<paddle::platform::CPUDeviceContext,
-                             paddle::platform::complex<float>>,
-    ops::AbsDoubleGradKernel<paddle::platform::CPUDeviceContext,
-                             paddle::platform::complex<double>>);
+REGISTER_OPERATOR(abs_double_grad, ops::AbsDoubleGradOp);
diff --git a/paddle/fluid/operators/abs_op.cu b/paddle/fluid/operators/abs_op.cu
deleted file mode 100644
index 882c8547a04154778389bd7cd77531b63d19915b..0000000000000000000000000000000000000000
--- a/paddle/fluid/operators/abs_op.cu
+++ /dev/null
@@ -1,89 +0,0 @@
-// Copyright (c) 2021 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
-//
-//     http://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/fluid/operators/abs_op.h"
-#include "paddle/fluid/operators/elementwise/elementwise_op_impl.cu.h"
-#include "paddle/fluid/platform/float16.h"
-
-namespace paddle {
-namespace operators {
-
-template <typename T, typename Enable = void>
-struct CudaAbsFunctor;
-
-template <typename T>
-struct CudaAbsFunctor<T, math::Complex<T, math::Real<T>>> {
-  __device__ __forceinline__ math::Real<T> operator()(const T x) const {
-    return abs(x);
-  }
-};
-
-template <typename T>
-struct CudaAbsFunctor<T, math::NoComplex<T, math::Real<T>>> {
-  __device__ __forceinline__ T operator()(const T x) const {
-    return std::abs(x);
-  }
-};
-
-template <typename T>
-class AbsKernel<platform::CUDADeviceContext, T>
-    : public framework::OpKernel<T> {
- public:
-  void Compute(const framework::ExecutionContext& context) const override {
-    const Tensor* x = context.Input<Tensor>("X");
-    Tensor* out = context.Output<Tensor>("Out");
-    out->mutable_data<math::Real<T>>(context.GetPlace());
-
-    auto& dev_ctx =
-        context.template device_context<platform::CUDADeviceContext>();
-    std::vector<const framework::Tensor*> ins = {x};
-    std::vector<framework::Tensor*> outs = {out};
-    auto functor = CudaAbsFunctor<T>();
-    paddle::operators::LaunchSameDimsElementwiseCudaKernel<math::Real<T>>(
-        dev_ctx, ins, &outs, functor);
-  }
-};
-
-}  // namespace operators
-}  // namespace paddle
-
-namespace ops = paddle::operators;
-namespace plat = paddle::platform;
-
-REGISTER_OP_CUDA_KERNEL(
-    abs, ops::AbsKernel<plat::CUDADeviceContext, float>,
-    ops::AbsKernel<plat::CUDADeviceContext, double>,
-    ops::AbsKernel<plat::CUDADeviceContext, int>,
-    ops::AbsKernel<plat::CUDADeviceContext, int64_t>,
-    ops::AbsKernel<plat::CUDADeviceContext, plat::float16>,
-    ops::AbsKernel<plat::CUDADeviceContext, plat::complex<float>>,
-    ops::AbsKernel<plat::CUDADeviceContext, plat::complex<double>>);
-
-REGISTER_OP_CUDA_KERNEL(
-    abs_grad, ops::AbsGradKernel<plat::CUDADeviceContext, float>,
-    ops::AbsGradKernel<plat::CUDADeviceContext, double>,
-    ops::AbsGradKernel<plat::CUDADeviceContext, int>,
-    ops::AbsGradKernel<plat::CUDADeviceContext, int64_t>,
-    ops::AbsGradKernel<plat::CUDADeviceContext, plat::float16>,
-    ops::AbsGradKernel<plat::CUDADeviceContext, plat::complex<float>>,
-    ops::AbsGradKernel<plat::CUDADeviceContext, plat::complex<double>>);
-
-REGISTER_OP_CUDA_KERNEL(
-    abs_grad_grad, ops::AbsDoubleGradKernel<plat::CUDADeviceContext, float>,
-    ops::AbsDoubleGradKernel<plat::CUDADeviceContext, double>,
-    ops::AbsDoubleGradKernel<plat::CUDADeviceContext, int>,
-    ops::AbsDoubleGradKernel<plat::CUDADeviceContext, int64_t>,
-    ops::AbsDoubleGradKernel<plat::CUDADeviceContext, plat::float16>,
-    ops::AbsDoubleGradKernel<plat::CUDADeviceContext, plat::complex<float>>,
-    ops::AbsDoubleGradKernel<plat::CUDADeviceContext, plat::complex<double>>);
diff --git a/paddle/fluid/operators/abs_op.h b/paddle/fluid/operators/abs_op.h
deleted file mode 100644
index c79e83314f3bd39dcf6736e66c0b12956a2b0e81..0000000000000000000000000000000000000000
--- a/paddle/fluid/operators/abs_op.h
+++ /dev/null
@@ -1,90 +0,0 @@
-// Copyright (c) 2021 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
-//
-//     http://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/fluid/framework/op_registry.h"
-#include "paddle/fluid/framework/operator.h"
-#include "paddle/fluid/operators/math/complex_functors.h"
-#include "paddle/fluid/platform/for_range.h"
-
-namespace paddle {
-namespace operators {
-using Tensor = framework::Tensor;
-
-template <typename DeviceContext, typename T>
-class AbsKernel : public framework::OpKernel<T> {
- public:
-  void Compute(const framework::ExecutionContext& context) const override {
-    const Tensor* x = context.Input<Tensor>("X");
-    Tensor* out = context.Output<Tensor>("Out");
-
-    auto numel = x->numel();
-    auto* x_data = x->data<T>();
-    auto* out_data = out->mutable_data<math::Real<T>>(
-        context.GetPlace(), size_t(x->numel() * sizeof(math::Real<T>)));
-
-    auto& dev_ctx = context.template device_context<DeviceContext>();
-    platform::ForRange<DeviceContext> for_range(dev_ctx, numel);
-    math::AbsFunctor<T> functor(x_data, out_data, numel);
-    for_range(functor);
-  }
-};
-
-template <typename DeviceContext, typename T>
-class AbsGradKernel : public framework::OpKernel<T> {
- public:
-  void Compute(const framework::ExecutionContext& ctx) const {
-    const framework::Tensor* d_out =
-        ctx.Input<framework::Tensor>(framework::GradVarName("Out"));
-    const framework::Tensor* x = ctx.Input<framework::Tensor>("X");
-    framework::Tensor* d_x =
-        ctx.Output<framework::Tensor>(framework::GradVarName("X"));
-
-    auto numel = d_out->numel();
-    auto* dout_data = d_out->data<math::Real<T>>();
-    auto* x_data = x->data<T>();
-    auto* dx_data = d_x->mutable_data<T>(
-        ctx.GetPlace(), static_cast<size_t>(numel * sizeof(T)));
-
-    auto& dev_ctx = ctx.template device_context<DeviceContext>();
-    platform::ForRange<DeviceContext> for_range(dev_ctx, numel);
-    math::AbsGradFunctor<T> functor(dout_data, x_data, dx_data, numel);
-    for_range(functor);
-  }
-};
-
-template <typename DeviceContext, typename T>
-class AbsDoubleGradKernel : public framework::OpKernel<T> {
- public:
-  void Compute(const framework::ExecutionContext& ctx) const {
-    const framework::Tensor* ddx = ctx.Input<framework::Tensor>("DDX");
-    const framework::Tensor* x = ctx.Input<framework::Tensor>("X");
-    framework::Tensor* ddout = ctx.Output<framework::Tensor>("DDOut");
-
-    auto numel = ddx->numel();
-    auto* ddx_data = ddx->data<T>();
-    auto* x_data = x->data<T>();
-    auto* ddout_data = ddout->mutable_data<T>(
-        ctx.GetPlace(), static_cast<size_t>(numel * sizeof(T)));
-
-    auto& dev_ctx = ctx.template device_context<DeviceContext>();
-    platform::ForRange<DeviceContext> for_range(dev_ctx, numel);
-    math::AbsGradGradFunctor<T> functor(ddx_data, x_data, ddout_data, numel);
-    for_range(functor);
-  }
-};
-
-}  // namespace operators
-}  // namespace paddle
diff --git a/paddle/fluid/operators/abs_op_npu.cc b/paddle/fluid/operators/abs_op_npu.cc
index cc2b0925c21e527c6835822161f2dcfd959b1b2d..30ec22cf6d868381a4a78585dc3620a2ea78d466 100644
--- a/paddle/fluid/operators/abs_op_npu.cc
+++ b/paddle/fluid/operators/abs_op_npu.cc
@@ -12,7 +12,7 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 See the License for the specific language governing permissions and
 limitations under the Licnse. */
 
-#include "paddle/fluid/operators/abs_op.h"
+#include "paddle/fluid/framework/op_registry.h"
 #include "paddle/fluid/platform/device/npu/npu_op_runner.h"
 
 namespace paddle {
diff --git a/paddle/pten/kernels/abs_grad_kernel.h b/paddle/pten/kernels/abs_grad_kernel.h
new file mode 100644
index 0000000000000000000000000000000000000000..494f29da783d4d6b6d3f6f940d3591ace578aea1
--- /dev/null
+++ b/paddle/pten/kernels/abs_grad_kernel.h
@@ -0,0 +1,34 @@
+// Copyright (c) 2022 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
+//
+//     http://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/pten/core/dense_tensor.h"
+#include "paddle/pten/core/device_context.h"
+
+namespace pten {
+
+template <typename T, typename Context>
+void AbsGradKernel(const Context& ctx,
+                   const DenseTensor& x,
+                   const DenseTensor& dout,
+                   DenseTensor* dx);
+
+template <typename T, typename Context>
+void AbsDoubleGradKernel(const Context& ctx,
+                         const DenseTensor& x,
+                         const DenseTensor& ddx,
+                         DenseTensor* ddout);
+
+}  // namespace pten
diff --git a/paddle/pten/kernels/abs_kernel.h b/paddle/pten/kernels/abs_kernel.h
new file mode 100644
index 0000000000000000000000000000000000000000..0322afadfab8d2a24a358ef7c747d09174b124f2
--- /dev/null
+++ b/paddle/pten/kernels/abs_kernel.h
@@ -0,0 +1,25 @@
+// Copyright (c) 2022 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
+//
+//     http://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/pten/core/dense_tensor.h"
+#include "paddle/pten/core/device_context.h"
+
+namespace pten {
+
+template <typename T, typename Context>
+void AbsKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out);
+
+}  // namespace pten
diff --git a/paddle/pten/kernels/cpu/abs_grad_kernel.cc b/paddle/pten/kernels/cpu/abs_grad_kernel.cc
new file mode 100644
index 0000000000000000000000000000000000000000..a3f3aabd16c097941ecc615ee51e62b90ef14b34
--- /dev/null
+++ b/paddle/pten/kernels/cpu/abs_grad_kernel.cc
@@ -0,0 +1,41 @@
+// Copyright (c) 2022 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
+//
+//     http://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/fluid/operators/math/complex_functors.h"
+#include "paddle/pten/common/complex.h"
+#include "paddle/pten/core/kernel_registry.h"
+#include "paddle/pten/kernels/impl/abs_grad_kernel_impl.h"
+
+using pten::dtype::complex;
+
+PT_REGISTER_KERNEL(abs_grad,
+                   CPU,
+                   ALL_LAYOUT,
+                   pten::AbsGradKernel,
+                   float,
+                   double,
+                   int,
+                   int64_t,
+                   complex<float>,
+                   complex<double>) {}
+PT_REGISTER_KERNEL(abs_double_grad,
+                   CPU,
+                   ALL_LAYOUT,
+                   pten::AbsDoubleGradKernel,
+                   float,
+                   double,
+                   int,
+                   int64_t,
+                   complex<float>,
+                   complex<double>) {}
diff --git a/paddle/pten/kernels/cpu/abs_kernel.cc b/paddle/pten/kernels/cpu/abs_kernel.cc
new file mode 100644
index 0000000000000000000000000000000000000000..49094f5c64ed3762b8e27d941df2ecc8359f6bf2
--- /dev/null
+++ b/paddle/pten/kernels/cpu/abs_kernel.cc
@@ -0,0 +1,48 @@
+// Copyright (c) 2022 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
+//
+//     http://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/pten/kernels/abs_kernel.h"
+#include "paddle/fluid/operators/math/complex_functors.h"
+#include "paddle/fluid/platform/for_range.h"
+#include "paddle/pten/backends/cpu/cpu_context.h"
+#include "paddle/pten/common/complex.h"
+#include "paddle/pten/core/kernel_registry.h"
+
+namespace pten {
+
+template <typename T, typename Context>
+void AbsKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out) {
+  auto numel = x.numel();
+  auto* x_data = x.data<T>();
+  ctx.template Alloc<paddle::operators::math::Real<T>>(
+      out, size_t(x.numel() * sizeof(paddle::operators::math::Real<T>)));
+  auto* out_data = out->data<paddle::operators::math::Real<T>>();
+
+  paddle::platform::ForRange<Context> for_range(ctx, numel);
+  paddle::operators::math::AbsFunctor<T> functor(x_data, out_data, numel);
+  for_range(functor);
+}
+
+}  // namespace pten
+
+PT_REGISTER_KERNEL(abs,
+                   CPU,
+                   ALL_LAYOUT,
+                   pten::AbsKernel,
+                   float,
+                   double,
+                   int,
+                   int64_t,
+                   pten::dtype::complex<float>,
+                   pten::dtype::complex<double>) {}
diff --git a/paddle/pten/kernels/gpu/abs_grad_kernel.cu b/paddle/pten/kernels/gpu/abs_grad_kernel.cu
new file mode 100644
index 0000000000000000000000000000000000000000..a7257e129ec47ac0b5b33923832855b7907cb719
--- /dev/null
+++ b/paddle/pten/kernels/gpu/abs_grad_kernel.cu
@@ -0,0 +1,44 @@
+// Copyright (c) 2022 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
+//
+//     http://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/pten/common/complex.h"
+#include "paddle/pten/common/float16.h"
+#include "paddle/pten/core/kernel_registry.h"
+#include "paddle/pten/kernels/abs_grad_kernel.h"
+#include "paddle/pten/kernels/impl/abs_grad_kernel_impl.h"
+
+using pten::dtype::complex;
+
+PT_REGISTER_KERNEL(abs_grad,
+                   GPU,
+                   ALL_LAYOUT,
+                   pten::AbsGradKernel,
+                   float,
+                   double,
+                   int,
+                   int64_t,
+                   pten::dtype::float16,
+                   complex<float>,
+                   complex<double>) {}
+PT_REGISTER_KERNEL(abs_double_grad,
+                   GPU,
+                   ALL_LAYOUT,
+                   pten::AbsDoubleGradKernel,
+                   float,
+                   double,
+                   int,
+                   int64_t,
+                   pten::dtype::float16,
+                   complex<float>,
+                   complex<double>) {}
diff --git a/paddle/pten/kernels/gpu/abs_kernel.cu b/paddle/pten/kernels/gpu/abs_kernel.cu
new file mode 100644
index 0000000000000000000000000000000000000000..d97aa79105322e8678fd559f9306a60b15c39080
--- /dev/null
+++ b/paddle/pten/kernels/gpu/abs_kernel.cu
@@ -0,0 +1,71 @@
+// Copyright (c) 2022 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
+//
+//     http://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 <algorithm>
+#include <vector>
+#include "paddle/fluid/operators/math/complex_functors.h"
+#include "paddle/pten/backends/gpu/gpu_context.h"
+#include "paddle/pten/core/dense_tensor.h"
+#include "paddle/pten/core/kernel_registry.h"
+#include "paddle/pten/kernels/abs_kernel.h"
+#include "paddle/pten/kernels/funcs/elementwise_base.h"
+
+namespace pten {
+
+template <typename T, typename Enable = void>
+struct CudaAbsFunctor;
+
+template <typename T>
+struct CudaAbsFunctor<
+    T,
+    paddle::operators::math::Complex<T, paddle::operators::math::Real<T>>> {
+  __device__ __forceinline__ paddle::operators::math::Real<T> operator()(
+      const T x) const {
+    return abs(x);
+  }
+};
+
+template <typename T>
+struct CudaAbsFunctor<
+    T,
+    paddle::operators::math::NoComplex<T, paddle::operators::math::Real<T>>> {
+  __device__ __forceinline__ T operator()(const T x) const {
+    return std::abs(x);
+  }
+};
+
+template <typename T, typename Context>
+void AbsKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out) {
+  ctx.template Alloc<paddle::operators::math::Real<T>>(out);
+  std::vector<const DenseTensor*> ins = {&x};
+  std::vector<DenseTensor*> outs = {out};
+  auto functor = CudaAbsFunctor<T>();
+
+  funcs::LaunchSameDimsElementwiseCudaKernel<paddle::operators::math::Real<T>>(
+      ctx, ins, &outs, functor);
+}
+
+}  // namespace pten
+
+PT_REGISTER_KERNEL(abs,
+                   GPU,
+                   ALL_LAYOUT,
+                   pten::AbsKernel,
+                   float,
+                   double,
+                   int,
+                   int64_t,
+                   pten::dtype::float16,
+                   pten::dtype::complex<float>,
+                   pten::dtype::complex<double>) {}
diff --git a/paddle/pten/kernels/impl/abs_grad_kernel_impl.h b/paddle/pten/kernels/impl/abs_grad_kernel_impl.h
new file mode 100644
index 0000000000000000000000000000000000000000..c702a0042dc4e2c7f75888a76c95b29dae667ec6
--- /dev/null
+++ b/paddle/pten/kernels/impl/abs_grad_kernel_impl.h
@@ -0,0 +1,58 @@
+// Copyright (c) 2022 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
+//
+//     http://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/fluid/operators/math/complex_functors.h"
+#include "paddle/fluid/platform/for_range.h"
+#include "paddle/pten/kernels/abs_grad_kernel.h"
+
+namespace pten {
+
+template <typename T, typename Context>
+void AbsGradKernel(const Context& ctx,
+                   const DenseTensor& x,
+                   const DenseTensor& dout,
+                   DenseTensor* dx) {
+  auto numel = dout.numel();
+  auto* dout_data = dout.data<paddle::operators::math::Real<T>>();
+  auto* x_data = x.data<T>();
+
+  ctx.template Alloc<T>(dx, static_cast<size_t>(numel * sizeof(T)));
+  auto* dx_data = dx->data<T>();
+
+  paddle::platform::ForRange<Context> for_range(ctx, numel);
+  paddle::operators::math::AbsGradFunctor<T> functor(
+      dout_data, x_data, dx_data, numel);
+  for_range(functor);
+}
+
+template <typename T, typename Context>
+void AbsDoubleGradKernel(const Context& ctx,
+                         const DenseTensor& x,
+                         const DenseTensor& ddx,
+                         DenseTensor* ddout) {
+  auto numel = ddx.numel();
+  auto* ddx_data = ddx.data<T>();
+  auto* x_data = x.data<T>();
+  ctx.template Alloc<T>(ddout, static_cast<size_t>(numel * sizeof(T)));
+  auto* ddout_data = ddout->data<T>();
+
+  paddle::platform::ForRange<Context> for_range(ctx, numel);
+  paddle::operators::math::AbsGradGradFunctor<T> functor(
+      ddx_data, x_data, ddout_data, numel);
+  for_range(functor);
+}
+
+}  // namespace pten
diff --git a/paddle/pten/ops/compat/abs_sig.cc b/paddle/pten/ops/compat/abs_sig.cc
new file mode 100644
index 0000000000000000000000000000000000000000..a610db46a16f4da0df1ebbf2cd0d5fda174cde50
--- /dev/null
+++ b/paddle/pten/ops/compat/abs_sig.cc
@@ -0,0 +1,38 @@
+// Copyright (c) 2022 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
+//
+//     http://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/pten/core/compat/op_utils.h"
+
+namespace pten {
+
+KernelSignature AbsOpArgumentMapping(const ArgumentMappingContext& ctx) {
+  return KernelSignature("abs", {"X"}, {}, {"Out"});
+}
+
+KernelSignature AbsGradOpArgumentMapping(const ArgumentMappingContext& ctx) {
+  return KernelSignature(
+      "abs_grad", {"X", GradVarName("Out")}, {}, {GradVarName("X")});
+}
+
+KernelSignature AbsDoubleGradOpArgumentMapping(
+    const ArgumentMappingContext& ctx) {
+  return KernelSignature("abs_double_grad", {"X", "DDX"}, {}, {"DDOut"});
+}
+
+}  // namespace pten
+
+PT_REGISTER_ARG_MAPPING_FN(abs, pten::AbsOpArgumentMapping);
+PT_REGISTER_ARG_MAPPING_FN(abs_grad, pten::AbsGradOpArgumentMapping);
+PT_REGISTER_ARG_MAPPING_FN(abs_double_grad,
+                           pten::AbsDoubleGradOpArgumentMapping);