From bbe441fc74a203e464c0b73c47693a65818837ee Mon Sep 17 00:00:00 2001 From: Zhou Wei <1183042833@qq.com> Date: Thu, 24 Feb 2022 13:28:38 +0800 Subject: [PATCH] =?UTF-8?q?=E3=80=90Phi=E3=80=91Migrate=20poisson=20op=20i?= =?UTF-8?q?nto=20phi=20(#39814)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Migrate poisson op into phi * fix CI * fix comment --- paddle/fluid/operators/poisson_op.cc | 51 ++--------- paddle/fluid/operators/poisson_op.cu | 91 ------------------- paddle/fluid/operators/poisson_op.h | 41 --------- paddle/phi/infermeta/unary.h | 1 + paddle/phi/kernels/cpu/poisson_grad_kernel.cc | 19 ++++ paddle/phi/kernels/cpu/poisson_kernel.cc | 41 +++++++++ paddle/phi/kernels/gpu/poisson_grad_kernel.cu | 19 ++++ paddle/phi/kernels/gpu/poisson_kernel.cu | 77 ++++++++++++++++ .../kernels/impl/poisson_grad_kernel_impl.h | 29 ++++++ paddle/phi/kernels/poisson_grad_kernel.h | 25 +++++ paddle/phi/kernels/poisson_kernel.h | 24 +++++ paddle/phi/ops/compat/poisson_sig.cc | 26 ++++++ 12 files changed, 270 insertions(+), 174 deletions(-) delete mode 100644 paddle/fluid/operators/poisson_op.cu delete mode 100644 paddle/fluid/operators/poisson_op.h create mode 100644 paddle/phi/kernels/cpu/poisson_grad_kernel.cc create mode 100644 paddle/phi/kernels/cpu/poisson_kernel.cc create mode 100644 paddle/phi/kernels/gpu/poisson_grad_kernel.cu create mode 100644 paddle/phi/kernels/gpu/poisson_kernel.cu create mode 100644 paddle/phi/kernels/impl/poisson_grad_kernel_impl.h create mode 100644 paddle/phi/kernels/poisson_grad_kernel.h create mode 100644 paddle/phi/kernels/poisson_kernel.h create mode 100644 paddle/phi/ops/compat/poisson_sig.cc diff --git a/paddle/fluid/operators/poisson_op.cc b/paddle/fluid/operators/poisson_op.cc index cc4b6e5e07..0cecbf0b9c 100644 --- a/paddle/fluid/operators/poisson_op.cc +++ b/paddle/fluid/operators/poisson_op.cc @@ -13,8 +13,10 @@ See the License for the specific language governing permissions and limitations under the License. */ #include - -#include "paddle/fluid/operators/poisson_op.h" +#include "paddle/fluid/framework/infershape_utils.h" +#include "paddle/fluid/framework/op_registry.h" +#include "paddle/phi/core/infermeta_utils.h" +#include "paddle/phi/infermeta/unary.h" namespace paddle { namespace operators { @@ -23,14 +25,6 @@ class PoissonOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - void InferShape(framework::InferShapeContext *ctx) const override { - OP_INOUT_CHECK(ctx->HasInput("X"), "Input", "X", "PoissonOp"); - OP_INOUT_CHECK(ctx->HasOutput("Out"), "Output", "Out", "PoissonOp"); - - auto dim = ctx->GetInputDim("X"); - ctx->SetOutputDim("Out", dim); - } - protected: framework::OpKernelType GetExpectedKernelType( const framework::ExecutionContext &ctx) const override { @@ -61,29 +55,6 @@ class PoissonOpInferVarType : public framework::PassInDtypeAndVarTypeToOutput { } }; -template -class PoissonKernel - : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext &ctx) const override { - const auto *x = ctx.Input("X"); - auto *out = ctx.Output("Out"); - - const T *x_data = x->data(); - T *out_data = out->mutable_data(ctx.GetPlace()); - - int64_t size = x->numel(); - - auto gen = framework::DefaultCPUGenerator(); - auto engine = gen->GetCPUEngine(); - - for (int64_t i = 0; i < size; ++i) { - std::poisson_distribution<> dist(x_data[i]); - out_data[i] = static_cast(dist(*engine)); - } - } -}; - class PoissonGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; @@ -116,17 +87,13 @@ class PoissonGradOpMaker : public framework::SingleGradOpMaker { namespace ops = paddle::operators; namespace plat = paddle::platform; +DELCARE_INFER_SHAPE_FUNCTOR(poisson, PoissonInferShapeFunctor, + PT_INFER_META(phi::UnchangedInferMeta)); + REGISTER_OPERATOR(poisson, ops::PoissonOp, ops::PoissonOpMaker, ops::PoissonOpInferVarType, ops::PoissonGradOpMaker, - ops::PoissonGradOpMaker); + ops::PoissonGradOpMaker, + PoissonInferShapeFunctor); REGISTER_OPERATOR(poisson_grad, ops::PoissonGradOp); - -REGISTER_OP_CPU_KERNEL(poisson, - ops::PoissonKernel, - ops::PoissonKernel); - -REGISTER_OP_CPU_KERNEL(poisson_grad, - ops::PoissonGradKernel, - ops::PoissonGradKernel); diff --git a/paddle/fluid/operators/poisson_op.cu b/paddle/fluid/operators/poisson_op.cu deleted file mode 100644 index ef2f6d4665..0000000000 --- a/paddle/fluid/operators/poisson_op.cu +++ /dev/null @@ -1,91 +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. */ - -#ifdef __NVCC__ -#include -#endif -#ifdef __HIPCC__ -#include -#endif -#include "paddle/fluid/operators/poisson_op.h" -#include "paddle/fluid/platform/for_range.h" - -namespace paddle { -namespace operators { - -template -struct PoissonCudaFunctor { - public: - PoissonCudaFunctor(const T* in, T* out, unsigned int seed, - unsigned int offset) - : in_(in), out_(out), seed_(seed), offset_(offset) {} - - __device__ void operator()(int64_t idx) { -#ifdef __NVCC__ - curandStatePhilox4_32_10_t state; - curand_init(seed_, idx, offset_, &state); - out_[idx] = static_cast(curand_poisson(&state, in_[idx])); -#elif __HIPCC__ - hiprandStatePhilox4_32_10_t state; - hiprand_init(seed_, idx, offset_, &state); - out_[idx] = static_cast(hiprand_poisson(&state, in_[idx])); -#endif - } - - private: - const T* in_; - T* out_; - const unsigned int seed_; - const unsigned int offset_; -}; - -template -class PoissonKernel - : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& ctx) const override { - const auto* x = ctx.Input("X"); - auto* out = ctx.Output("Out"); - - const T* x_data = x->data(); - T* out_data = out->mutable_data(ctx.GetPlace()); - auto size = x->numel(); - int64_t device_id = ctx.GetPlace().GetDeviceId(); - - auto gen_cuda = framework::GetDefaultCUDAGenerator(device_id); - auto seed_offset = gen_cuda->IncrementOffset(20); - uint64_t seed = seed_offset.first; - uint64_t offset = seed_offset.second; - - auto& dev_ctx = ctx.template device_context(); - platform::ForRange for_range(dev_ctx, size); - - PoissonCudaFunctor functor(x_data, out_data, seed, offset); - for_range(functor); - } -}; - -} // namespace operators -} // namespace paddle - -namespace ops = paddle::operators; -namespace plat = paddle::platform; - -REGISTER_OP_CUDA_KERNEL(poisson, - ops::PoissonKernel, - ops::PoissonKernel); - -REGISTER_OP_CUDA_KERNEL( - poisson_grad, ops::PoissonGradKernel, - ops::PoissonGradKernel); diff --git a/paddle/fluid/operators/poisson_op.h b/paddle/fluid/operators/poisson_op.h deleted file mode 100644 index 2bcb524401..0000000000 --- a/paddle/fluid/operators/poisson_op.h +++ /dev/null @@ -1,41 +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/generator.h" -#include "paddle/fluid/framework/op_registry.h" -#include "paddle/fluid/framework/operator.h" -#include "paddle/phi/kernels/funcs/math_function.h" - -namespace paddle { -namespace operators { - -template -class PoissonKernel; - -template -class PoissonGradKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& ctx) const override { - auto* dx = ctx.Output(framework::GradVarName("X")); - dx->mutable_data(ctx.GetPlace()); - phi::funcs::SetConstant functor; - auto& dev_ctx = ctx.template device_context(); - functor(dev_ctx, dx, static_cast(0)); - } -}; - -} // namespace operators -} // namespace paddle diff --git a/paddle/phi/infermeta/unary.h b/paddle/phi/infermeta/unary.h index 7d15f497ea..21cbe76bb1 100644 --- a/paddle/phi/infermeta/unary.h +++ b/paddle/phi/infermeta/unary.h @@ -103,4 +103,5 @@ void UnfoldInferMeta(const MetaTensor& x, const std::vector& dilations, MetaTensor* out, MetaConfig config = MetaConfig()); + } // namespace phi diff --git a/paddle/phi/kernels/cpu/poisson_grad_kernel.cc b/paddle/phi/kernels/cpu/poisson_grad_kernel.cc new file mode 100644 index 0000000000..4e274a7af9 --- /dev/null +++ b/paddle/phi/kernels/cpu/poisson_grad_kernel.cc @@ -0,0 +1,19 @@ +// 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/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/impl/poisson_grad_kernel_impl.h" + +PD_REGISTER_KERNEL( + poisson_grad, CPU, ALL_LAYOUT, phi::PoissonGradKernel, float, double) {} diff --git a/paddle/phi/kernels/cpu/poisson_kernel.cc b/paddle/phi/kernels/cpu/poisson_kernel.cc new file mode 100644 index 0000000000..6a3e32c2f0 --- /dev/null +++ b/paddle/phi/kernels/cpu/poisson_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 + +#include "paddle/phi/backends/cpu/cpu_context.h" +#include "paddle/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/poisson_kernel.h" + +namespace phi { + +template +void PoissonKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out) { + const T* x_data = x.data(); + T* out_data = ctx.template Alloc(out); + int64_t size = x.numel(); + + auto gen = ctx.GetGenerator(); + auto engine = gen->GetCPUEngine(); + + for (int64_t i = 0; i < size; ++i) { + std::poisson_distribution<> dist(x_data[i]); + out_data[i] = static_cast(dist(*engine)); + } +} + +} // namespace phi + +PD_REGISTER_KERNEL( + poisson, CPU, ALL_LAYOUT, phi::PoissonKernel, float, double) {} diff --git a/paddle/phi/kernels/gpu/poisson_grad_kernel.cu b/paddle/phi/kernels/gpu/poisson_grad_kernel.cu new file mode 100644 index 0000000000..8c16bc51ff --- /dev/null +++ b/paddle/phi/kernels/gpu/poisson_grad_kernel.cu @@ -0,0 +1,19 @@ +// 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/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/impl/poisson_grad_kernel_impl.h" + +PD_REGISTER_KERNEL( + poisson_grad, GPU, ALL_LAYOUT, phi::PoissonGradKernel, float, double) {} diff --git a/paddle/phi/kernels/gpu/poisson_kernel.cu b/paddle/phi/kernels/gpu/poisson_kernel.cu new file mode 100644 index 0000000000..ae97f2fca6 --- /dev/null +++ b/paddle/phi/kernels/gpu/poisson_kernel.cu @@ -0,0 +1,77 @@ +/* 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. */ + +#ifdef __NVCC__ +#include +#endif +#ifdef __HIPCC__ +#include +#endif + +#include "paddle/fluid/platform/for_range.h" +#include "paddle/phi/backends/gpu/gpu_context.h" +#include "paddle/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/poisson_kernel.h" + +namespace phi { + +template +struct PoissonCudaFunctor { + public: + PoissonCudaFunctor(const T* in, + T* out, + unsigned int seed, + unsigned int offset) + : in_(in), out_(out), seed_(seed), offset_(offset) {} + + __device__ void operator()(int64_t idx) { +#ifdef __NVCC__ + curandStatePhilox4_32_10_t state; + curand_init(seed_, idx, offset_, &state); + out_[idx] = static_cast(curand_poisson(&state, in_[idx])); +#elif __HIPCC__ + hiprandStatePhilox4_32_10_t state; + hiprand_init(seed_, idx, offset_, &state); + out_[idx] = static_cast(hiprand_poisson(&state, in_[idx])); +#endif + } + + private: + const T* in_; + T* out_; + const unsigned int seed_; + const unsigned int offset_; +}; + +template +void PoissonKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out) { + const T* x_data = x.data(); + T* out_data = ctx.template Alloc(out); + auto size = x.numel(); + + auto gen_cuda = ctx.GetGenerator(); + auto seed_offset = gen_cuda->IncrementOffset(20); + uint64_t seed = seed_offset.first; + uint64_t offset = seed_offset.second; + + paddle::platform::ForRange for_range(ctx, size); + + PoissonCudaFunctor functor(x_data, out_data, seed, offset); + for_range(functor); +} + +} // namespace phi + +PD_REGISTER_KERNEL( + poisson, GPU, ALL_LAYOUT, phi::PoissonKernel, float, double) {} diff --git a/paddle/phi/kernels/impl/poisson_grad_kernel_impl.h b/paddle/phi/kernels/impl/poisson_grad_kernel_impl.h new file mode 100644 index 0000000000..4e82cccac3 --- /dev/null +++ b/paddle/phi/kernels/impl/poisson_grad_kernel_impl.h @@ -0,0 +1,29 @@ +// 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/phi/kernels/funcs/math_function.h" +#include "paddle/phi/kernels/poisson_grad_kernel.h" + +namespace phi { + +template +void PoissonGradKernel(const Context& ctx, DenseTensor* x_grad) { + ctx.template Alloc(x_grad); + phi::funcs::SetConstant functor; + functor(ctx, x_grad, static_cast(0)); +} + +} // namespace phi diff --git a/paddle/phi/kernels/poisson_grad_kernel.h b/paddle/phi/kernels/poisson_grad_kernel.h new file mode 100644 index 0000000000..21720474f4 --- /dev/null +++ b/paddle/phi/kernels/poisson_grad_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/phi/core/dense_tensor.h" +#include "paddle/phi/core/device_context.h" + +namespace phi { + +template +void PoissonGradKernel(const Context& ctx, DenseTensor* x_grad); + +} // namespace phi diff --git a/paddle/phi/kernels/poisson_kernel.h b/paddle/phi/kernels/poisson_kernel.h new file mode 100644 index 0000000000..f67c9c4631 --- /dev/null +++ b/paddle/phi/kernels/poisson_kernel.h @@ -0,0 +1,24 @@ +// 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/phi/core/dense_tensor.h" + +namespace phi { + +template +void PoissonKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out); + +} // namespace phi diff --git a/paddle/phi/ops/compat/poisson_sig.cc b/paddle/phi/ops/compat/poisson_sig.cc new file mode 100644 index 0000000000..cb6ae28804 --- /dev/null +++ b/paddle/phi/ops/compat/poisson_sig.cc @@ -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. + +#include "paddle/phi/core/compat/op_utils.h" + +namespace phi { + +KernelSignature PoissonGradOpArgumentMapping( + const ArgumentMappingContext& ctx) { + return KernelSignature("poisson_grad", {}, {}, {GradVarName("X")}); +} + +} // namespace phi + +PD_REGISTER_ARG_MAPPING_FN(poisson_grad, phi::PoissonGradOpArgumentMapping); -- GitLab