From 00bbb8c59a2150e4cda68e0fae7a362e5cb663f5 Mon Sep 17 00:00:00 2001 From: furnace <34057289+windstamp@users.noreply.github.com> Date: Thu, 3 Mar 2022 10:51:17 +0800 Subject: [PATCH] [Phi] move gaussian_random (#39932) [Phi] move gaussian_random kernel --- paddle/fluid/operators/gaussian_random_op.cc | 23 ---- paddle/fluid/operators/gaussian_random_op.cu | 52 -------- .../phi/kernels/cpu/gaussian_random_kernel.cc | 53 +++++++++ paddle/phi/kernels/gaussian_random_kernel.h | 32 +++++ .../phi/kernels/gpu/gaussian_random_kernel.cu | 111 ++++++++++++++++++ paddle/phi/ops/compat/gaussian_random_sig.cc | 45 +++++++ 6 files changed, 241 insertions(+), 75 deletions(-) create mode 100644 paddle/phi/kernels/cpu/gaussian_random_kernel.cc create mode 100644 paddle/phi/kernels/gaussian_random_kernel.h create mode 100644 paddle/phi/kernels/gpu/gaussian_random_kernel.cu create mode 100644 paddle/phi/ops/compat/gaussian_random_sig.cc diff --git a/paddle/fluid/operators/gaussian_random_op.cc b/paddle/fluid/operators/gaussian_random_op.cc index 774ff0bd065..6b559885c56 100644 --- a/paddle/fluid/operators/gaussian_random_op.cc +++ b/paddle/fluid/operators/gaussian_random_op.cc @@ -26,27 +26,6 @@ namespace paddle { namespace operators { using Tensor = framework::Tensor; -template -class CPUGaussianRandomKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - float mean = context.Attr("mean"); - float std = context.Attr("std"); - auto* tensor = context.Output("Out"); - - std::normal_distribution dist(mean, std); - auto shape = GetShape(context); - tensor->Resize(shape); - int64_t size = tensor->numel(); - T* data = tensor->mutable_data(context.GetPlace()); - unsigned int seed = static_cast(context.Attr("seed")); - auto engine = framework::GetCPURandomEngine(seed); - - for (int64_t i = 0; i < size; ++i) { - data[i] = dist(*engine); - } - } -}; // namespace operators template class CPUGaussianRandomBatchSizeLikeKernel : public framework::OpKernel { @@ -194,8 +173,6 @@ Used to initialize tensors with gaussian random generator. namespace ops = paddle::operators; REGISTER_OP_WITHOUT_GRADIENT(gaussian_random, ops::GaussianRandomOp, ops::GaussianRandomOpMaker); -REGISTER_OP_CPU_KERNEL(gaussian_random, ops::CPUGaussianRandomKernel, - ops::CPUGaussianRandomKernel); REGISTER_OP_CPU_KERNEL(gaussian_random_batch_size_like, ops::CPUGaussianRandomBatchSizeLikeKernel, ops::CPUGaussianRandomBatchSizeLikeKernel); diff --git a/paddle/fluid/operators/gaussian_random_op.cu b/paddle/fluid/operators/gaussian_random_op.cu index 21d827c7920..d419bd70e67 100644 --- a/paddle/fluid/operators/gaussian_random_op.cu +++ b/paddle/fluid/operators/gaussian_random_op.cu @@ -52,53 +52,6 @@ struct GaussianGenerator { } }; -template -class GPUGaussianRandomKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - auto* tensor = context.Output("Out"); - unsigned int seed = static_cast(context.Attr("seed")); - bool seed_flag = false; - if (seed == 0) { - std::random_device rd; - seed = rd(); - seed_flag = true; - } - T mean = static_cast(context.Attr("mean")); - T std = static_cast(context.Attr("std")); - auto shape = GetShape(context); - tensor->Resize(shape); - - auto& dev_cxt = - context.template device_context(); - T* data = tensor->mutable_data(dev_cxt.GetPlace()); - - int64_t size = tensor->numel(); - - int device_id = context.GetPlace().GetDeviceId(); - auto gen_cuda = framework::GetDefaultCUDAGenerator(device_id); - - if (gen_cuda->GetIsInitPy() && seed_flag) { - if (FLAGS_use_curand) { - using MT = typename details::MPTypeTrait::Type; - distribution::normal_distribution dist; - distribution::normal_transform trans(mean, std); - distribution::distribution_and_transform(dev_cxt, tensor, dist, - trans); - } else { - auto seed_offset = gen_cuda->IncrementOffset(1); - int64_t gen_offset = size * seed_offset.second; - auto func = - GaussianGenerator(mean, std, seed_offset.first, gen_offset); - IndexKernel>(dev_cxt, tensor, func); - } - } else { - auto func = GaussianGenerator(mean, std, seed); - IndexKernel>(dev_cxt, tensor, func); - } - } -}; - template class GPUGaussianRandomBatchSizeLikeKernel : public framework::OpKernel { public: @@ -136,11 +89,6 @@ class GPUGaussianRandomBatchSizeLikeKernel : public framework::OpKernel { } // namespace operators } // namespace paddle -REGISTER_OP_CUDA_KERNEL( - gaussian_random, - paddle::operators::GPUGaussianRandomKernel, - paddle::operators::GPUGaussianRandomKernel, - paddle::operators::GPUGaussianRandomKernel); REGISTER_OP_CUDA_KERNEL( gaussian_random_batch_size_like, paddle::operators::GPUGaussianRandomBatchSizeLikeKernel< diff --git a/paddle/phi/kernels/cpu/gaussian_random_kernel.cc b/paddle/phi/kernels/cpu/gaussian_random_kernel.cc new file mode 100644 index 00000000000..7e336f18bf8 --- /dev/null +++ b/paddle/phi/kernels/cpu/gaussian_random_kernel.cc @@ -0,0 +1,53 @@ +// 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/kernels/gaussian_random_kernel.h" + +#include "paddle/phi/backends/cpu/cpu_context.h" +#include "paddle/phi/core/kernel_registry.h" + +#include "paddle/fluid/framework/generator.h" + +namespace phi { + +template +void GaussianRandomKernel(const Context& dev_ctx, + const ScalarArray& shape, + float mean, + float std, + int seed, + DataType dtype, + DenseTensor* out) { + auto tensor = out; + + std::normal_distribution dist(mean, std); + + tensor->Resize(phi::make_ddim(shape.GetData())); + int64_t size = tensor->numel(); + T* data = dev_ctx.template Alloc(tensor); + auto engine = paddle::framework::GetCPURandomEngine(seed); + + for (int64_t i = 0; i < size; ++i) { + data[i] = dist(*engine); + } +} + +} // namespace phi + +PD_REGISTER_KERNEL(gaussian_random, + CPU, + ALL_LAYOUT, + phi::GaussianRandomKernel, + float, + double) {} diff --git a/paddle/phi/kernels/gaussian_random_kernel.h b/paddle/phi/kernels/gaussian_random_kernel.h new file mode 100644 index 00000000000..2903d80d22d --- /dev/null +++ b/paddle/phi/kernels/gaussian_random_kernel.h @@ -0,0 +1,32 @@ +// 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/common/scalar_array.h" +#include "paddle/phi/core/dense_tensor.h" +#include "paddle/phi/core/device_context.h" + +namespace phi { + +template +void GaussianRandomKernel(const Context& ctx, + const ScalarArray& shape, + float mean, + float std, + int seed, + DataType dtype, + DenseTensor* out); + +} // namespace phi diff --git a/paddle/phi/kernels/gpu/gaussian_random_kernel.cu b/paddle/phi/kernels/gpu/gaussian_random_kernel.cu new file mode 100644 index 00000000000..d5acc60a360 --- /dev/null +++ b/paddle/phi/kernels/gpu/gaussian_random_kernel.cu @@ -0,0 +1,111 @@ +// 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/kernels/gaussian_random_kernel.h" + +#include +#include +#include +#include + +#include "paddle/phi/backends/gpu/gpu_context.h" +#include "paddle/phi/core/dense_tensor.h" +#include "paddle/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/funcs/distribution_helper.h" +#include "paddle/phi/kernels/funcs/index_impl.cu.h" + +#include "paddle/fluid/framework/generator.h" + +DECLARE_bool(use_curand); + +namespace phi { + +template +struct GaussianGenerator { + T mean_, std_; + unsigned int seed_; + unsigned int offset_ = 0; + + __host__ __device__ GaussianGenerator(T mean, T std, int seed) + : mean_(mean), std_(std), seed_(seed) {} + + __host__ __device__ GaussianGenerator(T mean, T std, int seed, int offset) + : mean_(mean), std_(std), seed_(seed), offset_(offset) {} + + __host__ __device__ T operator()(const unsigned int n) const { + thrust::minstd_rand rng; + rng.seed(seed_); + using MT = typename phi::kps::details::MPTypeTrait::Type; + thrust::normal_distribution dist(mean_, std_); + unsigned int new_n = n + offset_; + rng.discard(new_n); + MT out = dist(rng); + return static_cast(out); + } +}; + +template +void GaussianRandomKernel(const Context& dev_ctx, + const ScalarArray& shape, + float mean, + float std, + int seed, + DataType dtype, + DenseTensor* out) { + auto tensor = out; + + bool seed_flag = false; + if (seed == 0) { + std::random_device rd; + seed = rd(); + seed_flag = true; + } + + tensor->Resize(phi::make_ddim(shape.GetData())); + + T* data = dev_ctx.template Alloc(tensor); + + int64_t size = tensor->numel(); + + int device_id = dev_ctx.GetPlace().GetDeviceId(); + auto gen_cuda = paddle::framework::GetDefaultCUDAGenerator(device_id); + + using MT = typename phi::kps::details::MPTypeTrait::Type; + if (gen_cuda->GetIsInitPy() && seed_flag) { + if (FLAGS_use_curand) { + funcs::normal_distribution dist; + funcs::normal_transform trans(mean, std); + funcs::distribution_and_transform(dev_ctx, tensor, dist, trans); + } else { + auto seed_offset = gen_cuda->IncrementOffset(1); + int64_t gen_offset = size * seed_offset.second; + auto func = + GaussianGenerator(mean, std, seed_offset.first, gen_offset); + IndexKernel>(dev_ctx, tensor, func); + } + } else { + auto func = GaussianGenerator(mean, std, seed); + IndexKernel>(dev_ctx, tensor, func); + } +} + +} // namespace phi + +PD_REGISTER_KERNEL(gaussian_random, + GPU, + ALL_LAYOUT, + phi::GaussianRandomKernel, + phi::dtype::float16, + float, + double) {} diff --git a/paddle/phi/ops/compat/gaussian_random_sig.cc b/paddle/phi/ops/compat/gaussian_random_sig.cc new file mode 100644 index 00000000000..cddcb80ebea --- /dev/null +++ b/paddle/phi/ops/compat/gaussian_random_sig.cc @@ -0,0 +1,45 @@ +// 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 GaussianRandomOpArgumentMapping( + const ArgumentMappingContext& ctx) { + if (ctx.InputSize("ShapeTensorList") > 0) { + return KernelSignature("gaussian_random", + {}, + {"ShapeTensorList", "mean", "std", "seed", "dtype"}, + {"Out"}); + } + + const auto& shape = paddle::any_cast>(ctx.Attr("shape")); + if (ctx.HasInput("ShapeTensor") && shape.empty()) { + return KernelSignature("gaussian_random", + {}, + {"ShapeTensor", "mean", "std", "seed", "dtype"}, + {"Out"}); + } + + return KernelSignature("gaussian_random", + {}, + {"shape", "mean", "std", "seed", "dtype"}, + {"Out"}); +} + +} // namespace phi + +PD_REGISTER_ARG_MAPPING_FN(gaussian_random, + phi::GaussianRandomOpArgumentMapping); -- GitLab