From fb4730672dc0a0f9b1c93a347f0fc421587eeb4e Mon Sep 17 00:00:00 2001 From: From00 Date: Tue, 15 Feb 2022 11:12:20 +0800 Subject: [PATCH] Move Abs OP to pten (#39492) * Move Abs op to pten * Fix NPU compilation error * Fix CI error * Use LaunchSameDimsElementwiseCudaKernel in pten --- paddle/fluid/operators/abs_op.cc | 40 +-------- paddle/fluid/operators/abs_op.cu | 89 ------------------ paddle/fluid/operators/abs_op.h | 90 ------------------- paddle/fluid/operators/abs_op_npu.cc | 2 +- paddle/pten/kernels/abs_grad_kernel.h | 34 +++++++ paddle/pten/kernels/abs_kernel.h | 25 ++++++ paddle/pten/kernels/cpu/abs_grad_kernel.cc | 41 +++++++++ paddle/pten/kernels/cpu/abs_kernel.cc | 48 ++++++++++ paddle/pten/kernels/gpu/abs_grad_kernel.cu | 44 +++++++++ paddle/pten/kernels/gpu/abs_kernel.cu | 71 +++++++++++++++ .../pten/kernels/impl/abs_grad_kernel_impl.h | 58 ++++++++++++ paddle/pten/ops/compat/abs_sig.cc | 38 ++++++++ 12 files changed, 363 insertions(+), 217 deletions(-) delete mode 100644 paddle/fluid/operators/abs_op.cu delete mode 100644 paddle/fluid/operators/abs_op.h create mode 100644 paddle/pten/kernels/abs_grad_kernel.h create mode 100644 paddle/pten/kernels/abs_kernel.h create mode 100644 paddle/pten/kernels/cpu/abs_grad_kernel.cc create mode 100644 paddle/pten/kernels/cpu/abs_kernel.cc create mode 100644 paddle/pten/kernels/gpu/abs_grad_kernel.cu create mode 100644 paddle/pten/kernels/gpu/abs_kernel.cu create mode 100644 paddle/pten/kernels/impl/abs_grad_kernel_impl.h create mode 100644 paddle/pten/ops/compat/abs_sig.cc diff --git a/paddle/fluid/operators/abs_op.cc b/paddle/fluid/operators/abs_op.cc index 5a09933e0ee..149a87fe32d 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 #include #include #include +#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 { protected: void Apply(GradOpPtr 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, ops::AbsDoubleGradMaker); -REGISTER_OPERATOR(abs_grad_grad, ops::AbsDoubleGradOp); - -REGISTER_OP_CPU_KERNEL( - abs, ops::AbsKernel, - ops::AbsKernel, - ops::AbsKernel, - ops::AbsKernel, - ops::AbsKernel>, - ops::AbsKernel>); - -REGISTER_OP_CPU_KERNEL( - abs_grad, ops::AbsGradKernel, - ops::AbsGradKernel, - ops::AbsGradKernel, - ops::AbsGradKernel, - ops::AbsGradKernel>, - ops::AbsGradKernel>); - -REGISTER_OP_CPU_KERNEL( - abs_grad_grad, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel>, - ops::AbsDoubleGradKernel>); +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 882c8547a04..00000000000 --- 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 -struct CudaAbsFunctor; - -template -struct CudaAbsFunctor>> { - __device__ __forceinline__ math::Real operator()(const T x) const { - return abs(x); - } -}; - -template -struct CudaAbsFunctor>> { - __device__ __forceinline__ T operator()(const T x) const { - return std::abs(x); - } -}; - -template -class AbsKernel - : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - const Tensor* x = context.Input("X"); - Tensor* out = context.Output("Out"); - out->mutable_data>(context.GetPlace()); - - auto& dev_ctx = - context.template device_context(); - std::vector ins = {x}; - std::vector outs = {out}; - auto functor = CudaAbsFunctor(); - paddle::operators::LaunchSameDimsElementwiseCudaKernel>( - dev_ctx, ins, &outs, functor); - } -}; - -} // namespace operators -} // namespace paddle - -namespace ops = paddle::operators; -namespace plat = paddle::platform; - -REGISTER_OP_CUDA_KERNEL( - abs, ops::AbsKernel, - ops::AbsKernel, - ops::AbsKernel, - ops::AbsKernel, - ops::AbsKernel, - ops::AbsKernel>, - ops::AbsKernel>); - -REGISTER_OP_CUDA_KERNEL( - abs_grad, ops::AbsGradKernel, - ops::AbsGradKernel, - ops::AbsGradKernel, - ops::AbsGradKernel, - ops::AbsGradKernel, - ops::AbsGradKernel>, - ops::AbsGradKernel>); - -REGISTER_OP_CUDA_KERNEL( - abs_grad_grad, ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel, - ops::AbsDoubleGradKernel>, - ops::AbsDoubleGradKernel>); diff --git a/paddle/fluid/operators/abs_op.h b/paddle/fluid/operators/abs_op.h deleted file mode 100644 index c79e83314f3..00000000000 --- 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 -class AbsKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - const Tensor* x = context.Input("X"); - Tensor* out = context.Output("Out"); - - auto numel = x->numel(); - auto* x_data = x->data(); - auto* out_data = out->mutable_data>( - context.GetPlace(), size_t(x->numel() * sizeof(math::Real))); - - auto& dev_ctx = context.template device_context(); - platform::ForRange for_range(dev_ctx, numel); - math::AbsFunctor functor(x_data, out_data, numel); - for_range(functor); - } -}; - -template -class AbsGradKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& ctx) const { - const framework::Tensor* d_out = - ctx.Input(framework::GradVarName("Out")); - const framework::Tensor* x = ctx.Input("X"); - framework::Tensor* d_x = - ctx.Output(framework::GradVarName("X")); - - auto numel = d_out->numel(); - auto* dout_data = d_out->data>(); - auto* x_data = x->data(); - auto* dx_data = d_x->mutable_data( - ctx.GetPlace(), static_cast(numel * sizeof(T))); - - auto& dev_ctx = ctx.template device_context(); - platform::ForRange for_range(dev_ctx, numel); - math::AbsGradFunctor functor(dout_data, x_data, dx_data, numel); - for_range(functor); - } -}; - -template -class AbsDoubleGradKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& ctx) const { - const framework::Tensor* ddx = ctx.Input("DDX"); - const framework::Tensor* x = ctx.Input("X"); - framework::Tensor* ddout = ctx.Output("DDOut"); - - auto numel = ddx->numel(); - auto* ddx_data = ddx->data(); - auto* x_data = x->data(); - auto* ddout_data = ddout->mutable_data( - ctx.GetPlace(), static_cast(numel * sizeof(T))); - - auto& dev_ctx = ctx.template device_context(); - platform::ForRange for_range(dev_ctx, numel); - math::AbsGradGradFunctor 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 cc2b0925c21..30ec22cf6d8 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 00000000000..494f29da783 --- /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 +void AbsGradKernel(const Context& ctx, + const DenseTensor& x, + const DenseTensor& dout, + DenseTensor* dx); + +template +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 00000000000..0322afadfab --- /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 +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 00000000000..a3f3aabd16c --- /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, + complex) {} +PT_REGISTER_KERNEL(abs_double_grad, + CPU, + ALL_LAYOUT, + pten::AbsDoubleGradKernel, + float, + double, + int, + int64_t, + complex, + complex) {} diff --git a/paddle/pten/kernels/cpu/abs_kernel.cc b/paddle/pten/kernels/cpu/abs_kernel.cc new file mode 100644 index 00000000000..49094f5c64e --- /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 +void AbsKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out) { + auto numel = x.numel(); + auto* x_data = x.data(); + ctx.template Alloc>( + out, size_t(x.numel() * sizeof(paddle::operators::math::Real))); + auto* out_data = out->data>(); + + paddle::platform::ForRange for_range(ctx, numel); + paddle::operators::math::AbsFunctor 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, + pten::dtype::complex) {} 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 00000000000..a7257e129ec --- /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, + complex) {} +PT_REGISTER_KERNEL(abs_double_grad, + GPU, + ALL_LAYOUT, + pten::AbsDoubleGradKernel, + float, + double, + int, + int64_t, + pten::dtype::float16, + complex, + complex) {} diff --git a/paddle/pten/kernels/gpu/abs_kernel.cu b/paddle/pten/kernels/gpu/abs_kernel.cu new file mode 100644 index 00000000000..d97aa791053 --- /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 +#include +#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 +struct CudaAbsFunctor; + +template +struct CudaAbsFunctor< + T, + paddle::operators::math::Complex>> { + __device__ __forceinline__ paddle::operators::math::Real operator()( + const T x) const { + return abs(x); + } +}; + +template +struct CudaAbsFunctor< + T, + paddle::operators::math::NoComplex>> { + __device__ __forceinline__ T operator()(const T x) const { + return std::abs(x); + } +}; + +template +void AbsKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out) { + ctx.template Alloc>(out); + std::vector ins = {&x}; + std::vector outs = {out}; + auto functor = CudaAbsFunctor(); + + funcs::LaunchSameDimsElementwiseCudaKernel>( + 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, + pten::dtype::complex) {} 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 00000000000..c702a0042dc --- /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 +void AbsGradKernel(const Context& ctx, + const DenseTensor& x, + const DenseTensor& dout, + DenseTensor* dx) { + auto numel = dout.numel(); + auto* dout_data = dout.data>(); + auto* x_data = x.data(); + + ctx.template Alloc(dx, static_cast(numel * sizeof(T))); + auto* dx_data = dx->data(); + + paddle::platform::ForRange for_range(ctx, numel); + paddle::operators::math::AbsGradFunctor functor( + dout_data, x_data, dx_data, numel); + for_range(functor); +} + +template +void AbsDoubleGradKernel(const Context& ctx, + const DenseTensor& x, + const DenseTensor& ddx, + DenseTensor* ddout) { + auto numel = ddx.numel(); + auto* ddx_data = ddx.data(); + auto* x_data = x.data(); + ctx.template Alloc(ddout, static_cast(numel * sizeof(T))); + auto* ddout_data = ddout->data(); + + paddle::platform::ForRange for_range(ctx, numel); + paddle::operators::math::AbsGradGradFunctor 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 00000000000..a610db46a16 --- /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); -- GitLab