diff --git a/paddle/fluid/operators/trunc_op.cc b/paddle/fluid/operators/trunc_op.cc index 2b79e2152b2f3414c3e3b7794e8c07c00a2aee00..bd3dc002990a7cf3af738eb2d914b3fc3dd9e79a 100644 --- a/paddle/fluid/operators/trunc_op.cc +++ b/paddle/fluid/operators/trunc_op.cc @@ -12,7 +12,8 @@ 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/trunc_op.h" +#include "paddle/fluid/framework/op_registry.h" +#include "paddle/fluid/framework/operator.h" namespace paddle { namespace operators { @@ -80,10 +81,3 @@ REGISTER_OPERATOR(trunc, ops::TruncOp, ops::TruncOpMaker, ops::TruncGradOpMaker); REGISTER_OPERATOR(trunc_grad, ops::TruncGradOp); - -REGISTER_OP_CPU_KERNEL(trunc, ops::TruncKernel, ops::TruncKernel, - ops::TruncKernel, ops::TruncKernel); - -REGISTER_OP_CPU_KERNEL(trunc_grad, ops::TruncGradKernel, - ops::TruncGradKernel, ops::TruncGradKernel, - ops::TruncGradKernel); diff --git a/paddle/fluid/operators/trunc_op.cu b/paddle/fluid/operators/trunc_op.cu deleted file mode 100644 index 68d8c608f6338802067fbba3141a58087dcafe62..0000000000000000000000000000000000000000 --- a/paddle/fluid/operators/trunc_op.cu +++ /dev/null @@ -1,115 +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/trunc_op.h" -#include "paddle/fluid/platform/device/gpu/gpu_info.h" -#include "paddle/fluid/platform/device/gpu/gpu_primitives.h" - -namespace paddle { -namespace operators { - -using platform::PADDLE_CUDA_NUM_THREADS; - -template -class TruncFunctor { - public: - __device__ TruncFunctor(const T x) : x_(x) {} - __device__ T operator()() { return trunc(x_); } - - public: - const T x_; -}; - -template <> -class TruncFunctor { - public: - __device__ TruncFunctor(const int x) : x_(x) {} - __device__ int operator()() { return x_; } - - public: - const int x_; -}; - -template <> -class TruncFunctor { - public: - __device__ TruncFunctor(const int64_t x) : x_(x) {} - __device__ int64_t operator()() { return x_; } - - public: - const int64_t x_; -}; - -template -__global__ void Trunc(const T* x, T* out, int64_t N) { - CUDA_KERNEL_LOOP(index, N) { - TruncFunctor functor(x[index]); - out[index] = functor(); - } -} - -template -__global__ void TruncGrad(T* dx, int64_t N) { - CUDA_KERNEL_LOOP(index, N) { dx[index] = static_cast(0.0); } -} - -template -class TruncCUDAKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - auto* x = context.Input("X"); - auto* out = context.Output("Out"); - - const auto* x_data = x->data(); - auto* out_data = out->mutable_data(context.GetPlace()); - - int64_t numel = x->numel(); - - int theads = PADDLE_CUDA_NUM_THREADS; - int blocks = (numel + theads - 1) / theads; - - Trunc<<>>(x_data, out_data, numel); - } -}; - -template -class TruncCUDAGradKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - auto* dout = context.Input(framework::GradVarName("Out")); - auto* dx = context.Output(framework::GradVarName("X")); - - const auto* dout_data = dout->data(); - auto* dx_data = dx->mutable_data(context.GetPlace()); - - int64_t numel = dout->numel(); - - int theads = PADDLE_CUDA_NUM_THREADS; - int blocks = (numel + theads - 1) / theads; - - TruncGrad<<>>(dx_data, numel); - } -}; - -} // namespace operators -} // namespace paddle - -namespace ops = paddle::operators; -REGISTER_OP_CUDA_KERNEL(trunc, ops::TruncCUDAKernel, - ops::TruncCUDAKernel, ops::TruncCUDAKernel, - ops::TruncCUDAKernel); - -REGISTER_OP_CUDA_KERNEL(trunc_grad, ops::TruncCUDAGradKernel, - ops::TruncCUDAGradKernel, - ops::TruncCUDAGradKernel, - ops::TruncCUDAGradKernel); diff --git a/paddle/fluid/operators/trunc_op.h b/paddle/fluid/operators/trunc_op.h deleted file mode 100644 index 0f788eae5249c57b92c7558451eca641a6840a41..0000000000000000000000000000000000000000 --- a/paddle/fluid/operators/trunc_op.h +++ /dev/null @@ -1,55 +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 -#include "paddle/fluid/framework/op_registry.h" -#include "paddle/fluid/framework/operator.h" - -namespace paddle { -namespace operators { - -using Tensor = framework::Tensor; - -template -class TruncKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - const Tensor* x = context.Input("X"); - Tensor* out = context.Output("Out"); - - size_t numel = x->numel(); - const T* x_data = x->data(); - T* out_data = out->mutable_data(context.GetPlace()); - - for (size_t i = 0; i < numel; i++) { - out_data[i] = trunc(x_data[i]); - } - } -}; - -template -class TruncGradKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - auto* dx = context.Output(framework::GradVarName("X")); - T* dx_data = dx->mutable_data(context.GetPlace()); - - int numel = dx->numel(); - memset(dx_data, 0.0, numel * sizeof(T)); - } -}; - -} // namespace operators -} // namespace paddle diff --git a/paddle/pten/kernels/cpu/trunc_grad_kernel.cc b/paddle/pten/kernels/cpu/trunc_grad_kernel.cc new file mode 100644 index 0000000000000000000000000000000000000000..2fa06321c558255c0e4190f0c64999217d4bec06 --- /dev/null +++ b/paddle/pten/kernels/cpu/trunc_grad_kernel.cc @@ -0,0 +1,40 @@ +// 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/trunc_grad_kernel.h" +#include "paddle/pten/backends/cpu/cpu_context.h" +#include "paddle/pten/core/kernel_registry.h" + +namespace pten { + +template +void TruncGradKernel(const Context& dev_ctx, + const DenseTensor& out_grad, + DenseTensor* in_grad) { + T* dx_data = dev_ctx.template Alloc(in_grad); + + int numel = in_grad->numel(); + memset(dx_data, 0.0, numel * sizeof(T)); +} + +} // namespace pten + +PT_REGISTER_KERNEL(trunc_grad, + CPU, + ALL_LAYOUT, + pten::TruncGradKernel, + float, + double, + int, + int64_t) {} diff --git a/paddle/pten/kernels/cpu/trunc_kernel.cc b/paddle/pten/kernels/cpu/trunc_kernel.cc new file mode 100644 index 0000000000000000000000000000000000000000..8bcef2351b80830699f8da53aba8c5a4c9f25975 --- /dev/null +++ b/paddle/pten/kernels/cpu/trunc_kernel.cc @@ -0,0 +1,39 @@ +// 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 "paddle/pten/backends/cpu/cpu_context.h" +#include "paddle/pten/core/kernel_registry.h" +#include "paddle/pten/kernels/trunc_kernel.h" + +namespace pten { + +template +void TruncKernel(const Context& dev_ctx, + const DenseTensor& x, + DenseTensor* out) { + size_t numel = x.numel(); + const T* x_data = x.data(); + T* out_data = dev_ctx.template Alloc(out); + + for (size_t i = 0; i < numel; i++) { + out_data[i] = trunc(x_data[i]); + } +} + +} // namespace pten + +PT_REGISTER_KERNEL( + trunc, CPU, ALL_LAYOUT, pten::TruncKernel, float, double, int, int64_t) {} diff --git a/paddle/pten/kernels/gpu/trunc_grad_kernel.cu b/paddle/pten/kernels/gpu/trunc_grad_kernel.cu new file mode 100644 index 0000000000000000000000000000000000000000..abb756351f765fabd2f4ee19d7172cc23a82d911 --- /dev/null +++ b/paddle/pten/kernels/gpu/trunc_grad_kernel.cu @@ -0,0 +1,54 @@ +// 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/platform/device/gpu/gpu_primitives.h" +#include "paddle/pten/backends/gpu/gpu_context.h" +#include "paddle/pten/backends/gpu/gpu_info.h" +#include "paddle/pten/core/kernel_registry.h" +#include "paddle/pten/kernels/trunc_grad_kernel.h" + +namespace pten { + +using paddle::platform::PADDLE_CUDA_NUM_THREADS; + +template +__global__ void TruncGrad(T* dx, int64_t N) { + CUDA_KERNEL_LOOP(index, N) { dx[index] = static_cast(0.0); } +} + +template +void TruncGradKernel(const Context& dev_ctx, + const DenseTensor& out_grad, + DenseTensor* in_grad) { + const auto* out_grad_data = out_grad.data(); + T* in_grad_data = dev_ctx.template Alloc(in_grad); + + int64_t numel = out_grad.numel(); + + int theads = PADDLE_CUDA_NUM_THREADS; + int blocks = (numel + theads - 1) / theads; + + TruncGrad<<>>(in_grad_data, numel); +} + +} // namespace pten + +PT_REGISTER_KERNEL(trunc_grad, + GPU, + ALL_LAYOUT, + pten::TruncGradKernel, + float, + double, + int, + int64_t) {} diff --git a/paddle/pten/kernels/gpu/trunc_kernel.cu b/paddle/pten/kernels/gpu/trunc_kernel.cu new file mode 100644 index 0000000000000000000000000000000000000000..3c9f491066cd2bdacbe7e661c3c8e1f82be7fc71 --- /dev/null +++ b/paddle/pten/kernels/gpu/trunc_kernel.cu @@ -0,0 +1,81 @@ +// 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/platform/device/gpu/gpu_primitives.h" +#include "paddle/pten/backends/gpu/gpu_context.h" +#include "paddle/pten/backends/gpu/gpu_info.h" +#include "paddle/pten/core/kernel_registry.h" +#include "paddle/pten/kernels/trunc_kernel.h" + +namespace pten { + +using paddle::platform::PADDLE_CUDA_NUM_THREADS; + +template +class TruncFunctor { + public: + __device__ TruncFunctor(const T x) : x_(x) {} + __device__ T operator()() { return trunc(x_); } + + public: + const T x_; +}; + +template <> +class TruncFunctor { + public: + __device__ TruncFunctor(const int x) : x_(x) {} + __device__ int operator()() { return x_; } + + public: + const int x_; +}; + +template <> +class TruncFunctor { + public: + __device__ TruncFunctor(const int64_t x) : x_(x) {} + __device__ int64_t operator()() { return x_; } + + public: + const int64_t x_; +}; + +template +__global__ void Trunc(const T* x, T* out, int64_t N) { + CUDA_KERNEL_LOOP(index, N) { + TruncFunctor functor(x[index]); + out[index] = functor(); + } +} + +template +void TruncKernel(const Context& dev_ctx, + const DenseTensor& x, + DenseTensor* out) { + const auto* x_data = x.data(); + auto* out_data = dev_ctx.template Alloc(out); + + int64_t numel = x.numel(); + + int theads = PADDLE_CUDA_NUM_THREADS; + int blocks = (numel + theads - 1) / theads; + + Trunc<<>>(x_data, out_data, numel); +} + +} // namespace pten + +PT_REGISTER_KERNEL( + trunc, GPU, ALL_LAYOUT, pten::TruncKernel, float, double, int, int64_t) {} diff --git a/paddle/pten/kernels/trunc_grad_kernel.h b/paddle/pten/kernels/trunc_grad_kernel.h new file mode 100644 index 0000000000000000000000000000000000000000..d567b84a1ebaea90182f085a9cefccf792ee0ff7 --- /dev/null +++ b/paddle/pten/kernels/trunc_grad_kernel.h @@ -0,0 +1,26 @@ +// 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" + +namespace pten { + +template +void TruncGradKernel(const Context& dev_ctx, + const DenseTensor& out_grad, + DenseTensor* in_grad); + +} // namespace pten diff --git a/paddle/pten/kernels/trunc_kernel.h b/paddle/pten/kernels/trunc_kernel.h new file mode 100644 index 0000000000000000000000000000000000000000..e8046d8b85d725dcf172e80253e762b8b65f27cc --- /dev/null +++ b/paddle/pten/kernels/trunc_kernel.h @@ -0,0 +1,26 @@ +// 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" + +namespace pten { + +template +void TruncKernel(const Context& dev_ctx, + const DenseTensor& x, + DenseTensor* out); + +} // namespace pten diff --git a/paddle/pten/ops/compat/trunc_sig.cc b/paddle/pten/ops/compat/trunc_sig.cc new file mode 100644 index 0000000000000000000000000000000000000000..1434418200cad662126b547dbbb5ea10e192ed16 --- /dev/null +++ b/paddle/pten/ops/compat/trunc_sig.cc @@ -0,0 +1,31 @@ +// 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 TruncOpArgumentMapping(const ArgumentMappingContext& ctx) { + return KernelSignature("trunc", {"X"}, {}, {"Out"}); +} + +KernelSignature TruncGradOpArgumentMapping(const ArgumentMappingContext& ctx) { + return KernelSignature( + "trunc_grad", {GradVarName("Out")}, {}, {GradVarName("X")}); +} + +} // namespace pten + +PT_REGISTER_ARG_MAPPING_FN(trunc, pten::TruncOpArgumentMapping); +PT_REGISTER_ARG_MAPPING_FN(trunc_grad, pten::TruncGradOpArgumentMapping);