From d39e7896c730011a56de4c65ae9387ceed3534d1 Mon Sep 17 00:00:00 2001 From: Chen Weihang Date: Thu, 7 Apr 2022 10:36:20 +0800 Subject: [PATCH] [Phi] Polish truncated normal kernel and add yaml (#41280) * polish truncated normal kernel * add yaml * add truncated normal kernel and add yaml * polish unittests and yaml * import dygraph mehtod --- .../cpu/truncated_gaussian_random_kernel.cc | 143 +++++++++++++++++- .../gpu/truncated_gaussian_random_kernel.cu | 5 +- .../truncated_gaussian_random_kernel.h | 138 ----------------- python/paddle/fluid/initializer.py | 15 +- .../test_truncated_gaussian_random_op.py | 17 ++- python/paddle/utils/code_gen/api.yaml | 13 ++ 6 files changed, 183 insertions(+), 148 deletions(-) diff --git a/paddle/phi/kernels/cpu/truncated_gaussian_random_kernel.cc b/paddle/phi/kernels/cpu/truncated_gaussian_random_kernel.cc index 4247e597ace..10280082619 100644 --- a/paddle/phi/kernels/cpu/truncated_gaussian_random_kernel.cc +++ b/paddle/phi/kernels/cpu/truncated_gaussian_random_kernel.cc @@ -21,10 +21,141 @@ #include "paddle/phi/backends/cpu/cpu_context.h" #include "paddle/phi/core/kernel_registry.h" -#include "paddle/fluid/framework/generator.h" - namespace phi { +// reference: https://gist.github.com/lakshayg/d80172fe5ae3c5d2c2aedb53c250320e +template +T Erfinv(T x) { + if (x < -1 || x > 1) { + return std::numeric_limits::quiet_NaN(); + } else if (x == 1.0) { + return std::numeric_limits::infinity(); + } else if (x == -1.0) { + return -std::numeric_limits::infinity(); + } + + const T LN2 = 6.931471805599453094172321214581e-1; + + const T A0 = 1.1975323115670912564578e0; + const T A1 = 4.7072688112383978012285e1; + const T A2 = 6.9706266534389598238465e2; + const T A3 = 4.8548868893843886794648e3; + const T A4 = 1.6235862515167575384252e4; + const T A5 = 2.3782041382114385731252e4; + const T A6 = 1.1819493347062294404278e4; + const T A7 = 8.8709406962545514830200e2; + + const T B0 = 1.0000000000000000000e0; + const T B1 = 4.2313330701600911252e1; + const T B2 = 6.8718700749205790830e2; + const T B3 = 5.3941960214247511077e3; + const T B4 = 2.1213794301586595867e4; + const T B5 = 3.9307895800092710610e4; + const T B6 = 2.8729085735721942674e4; + const T B7 = 5.2264952788528545610e3; + + const T C0 = 1.42343711074968357734e0; + const T C1 = 4.63033784615654529590e0; + const T C2 = 5.76949722146069140550e0; + const T C3 = 3.64784832476320460504e0; + const T C4 = 1.27045825245236838258e0; + const T C5 = 2.41780725177450611770e-1; + const T C6 = 2.27238449892691845833e-2; + const T C7 = 7.74545014278341407640e-4; + + const T D0 = 1.4142135623730950488016887e0; + const T D1 = 2.9036514445419946173133295e0; + const T D2 = 2.3707661626024532365971225e0; + const T D3 = 9.7547832001787427186894837e-1; + const T D4 = 2.0945065210512749128288442e-1; + const T D5 = 2.1494160384252876777097297e-2; + const T D6 = 7.7441459065157709165577218e-4; + const T D7 = 1.4859850019840355905497876e-9; + + const T E0 = 6.65790464350110377720e0; + const T E1 = 5.46378491116411436990e0; + const T E2 = 1.78482653991729133580e0; + const T E3 = 2.96560571828504891230e-1; + const T E4 = 2.65321895265761230930e-2; + const T E5 = 1.24266094738807843860e-3; + const T E6 = 2.71155556874348757815e-5; + const T E7 = 2.01033439929228813265e-7; + + const T F0 = 1.414213562373095048801689e0; + const T F1 = 8.482908416595164588112026e-1; + const T F2 = 1.936480946950659106176712e-1; + const T F3 = 2.103693768272068968719679e-2; + const T F4 = 1.112800997078859844711555e-3; + const T F5 = 2.611088405080593625138020e-5; + const T F6 = 2.010321207683943062279931e-7; + const T F7 = 2.891024605872965461538222e-15; + + T abs_x = abs(x); + + if (abs_x <= 0.85) { + T r = 0.180625 - 0.25 * x * x; + T num = + (((((((A7 * r + A6) * r + A5) * r + A4) * r + A3) * r + A2) * r + A1) * + r + + A0); + T den = + (((((((B7 * r + B6) * r + B5) * r + B4) * r + B3) * r + B2) * r + B1) * + r + + B0); + return x * num / den; + } + + T r = sqrt(LN2 - log(1.0 - abs_x)); + + T num, den; + if (r <= 5.0) { + r = r - 1.6; + num = + (((((((C7 * r + C6) * r + C5) * r + C4) * r + C3) * r + C2) * r + C1) * + r + + C0); + den = + (((((((D7 * r + D6) * r + D5) * r + D4) * r + D3) * r + D2) * r + D1) * + r + + D0); + } else { + r = r - 5.0; + num = + (((((((E7 * r + E6) * r + E5) * r + E4) * r + E3) * r + E2) * r + E1) * + r + + E0); + den = + (((((((F7 * r + F6) * r + F5) * r + F4) * r + F3) * r + F2) * r + F1) * + r + + F0); + } + + if (x < 0) { + return -num / den; + } else { + return num / den; + } +} + +template +struct TruncatedNormal { + T mean, std; + T a_normal_cdf; + T b_normal_cdf; + TruncatedNormal(T mean, T std) : mean(mean), std(std) { + auto normal_cdf = [](T x) { + return (1.0 + std::erf(x / std::sqrt(2.0))) / 2.0; + }; + a_normal_cdf = normal_cdf(-2.0); + b_normal_cdf = normal_cdf(2.0); + } + + T operator()(T value) const { + auto p = a_normal_cdf + (b_normal_cdf - a_normal_cdf) * value; + return std::sqrt(2.0) * Erfinv(2 * p - 1) * std + mean; + } +}; + template void TruncatedGaussianRandomKernel(const Context& dev_ctx, const std::vector& shape, @@ -42,7 +173,13 @@ void TruncatedGaussianRandomKernel(const Context& dev_ctx, TruncatedNormal truncated_normal(mean, std); int64_t size = tensor->numel(); - auto engine = paddle::framework::GetCPURandomEngine(seed); + std::shared_ptr engine; + if (seed) { + engine = std::make_shared(); + engine->seed(seed); + } else { + engine = dev_ctx.GetGenerator()->GetCPUEngine(); + } for (int64_t i = 0; i < size; ++i) { data[i] = truncated_normal(dist(*engine)); } diff --git a/paddle/phi/kernels/gpu/truncated_gaussian_random_kernel.cu b/paddle/phi/kernels/gpu/truncated_gaussian_random_kernel.cu index f27b32ca7b8..5b6ae9d09bf 100644 --- a/paddle/phi/kernels/gpu/truncated_gaussian_random_kernel.cu +++ b/paddle/phi/kernels/gpu/truncated_gaussian_random_kernel.cu @@ -24,8 +24,6 @@ #include "paddle/phi/core/dense_tensor.h" #include "paddle/phi/core/kernel_registry.h" -#include "paddle/fluid/framework/generator.h" - namespace phi { template @@ -106,8 +104,7 @@ void TruncatedGaussianRandomKernel(const Context& dev_ctx, thrust::counting_iterator index_sequence_begin(0); int64_t size = tensor->numel(); - int device_id = dev_ctx.GetPlace().GetDeviceId(); - auto gen_cuda = paddle::framework::GetDefaultCUDAGenerator(device_id); + auto gen_cuda = dev_ctx.GetGenerator(); if (gen_cuda->GetIsInitPy() && seed_flag) { auto seed_offset = gen_cuda->IncrementOffset(1); diff --git a/paddle/phi/kernels/truncated_gaussian_random_kernel.h b/paddle/phi/kernels/truncated_gaussian_random_kernel.h index 2781b79520a..773bfc8c71e 100644 --- a/paddle/phi/kernels/truncated_gaussian_random_kernel.h +++ b/paddle/phi/kernels/truncated_gaussian_random_kernel.h @@ -14,149 +14,11 @@ #pragma once -#include -#include - #include "paddle/phi/common/int_array.h" #include "paddle/phi/core/dense_tensor.h" -#include "paddle/phi/core/device_context.h" -#include "paddle/phi/infermeta/nullary.h" namespace phi { -// reference: https://gist.github.com/lakshayg/d80172fe5ae3c5d2c2aedb53c250320e -template -T Erfinv(T x) { - if (x < -1 || x > 1) { - return std::numeric_limits::quiet_NaN(); - } else if (x == 1.0) { - return std::numeric_limits::infinity(); - } else if (x == -1.0) { - return -std::numeric_limits::infinity(); - } - - const T LN2 = 6.931471805599453094172321214581e-1; - - const T A0 = 1.1975323115670912564578e0; - const T A1 = 4.7072688112383978012285e1; - const T A2 = 6.9706266534389598238465e2; - const T A3 = 4.8548868893843886794648e3; - const T A4 = 1.6235862515167575384252e4; - const T A5 = 2.3782041382114385731252e4; - const T A6 = 1.1819493347062294404278e4; - const T A7 = 8.8709406962545514830200e2; - - const T B0 = 1.0000000000000000000e0; - const T B1 = 4.2313330701600911252e1; - const T B2 = 6.8718700749205790830e2; - const T B3 = 5.3941960214247511077e3; - const T B4 = 2.1213794301586595867e4; - const T B5 = 3.9307895800092710610e4; - const T B6 = 2.8729085735721942674e4; - const T B7 = 5.2264952788528545610e3; - - const T C0 = 1.42343711074968357734e0; - const T C1 = 4.63033784615654529590e0; - const T C2 = 5.76949722146069140550e0; - const T C3 = 3.64784832476320460504e0; - const T C4 = 1.27045825245236838258e0; - const T C5 = 2.41780725177450611770e-1; - const T C6 = 2.27238449892691845833e-2; - const T C7 = 7.74545014278341407640e-4; - - const T D0 = 1.4142135623730950488016887e0; - const T D1 = 2.9036514445419946173133295e0; - const T D2 = 2.3707661626024532365971225e0; - const T D3 = 9.7547832001787427186894837e-1; - const T D4 = 2.0945065210512749128288442e-1; - const T D5 = 2.1494160384252876777097297e-2; - const T D6 = 7.7441459065157709165577218e-4; - const T D7 = 1.4859850019840355905497876e-9; - - const T E0 = 6.65790464350110377720e0; - const T E1 = 5.46378491116411436990e0; - const T E2 = 1.78482653991729133580e0; - const T E3 = 2.96560571828504891230e-1; - const T E4 = 2.65321895265761230930e-2; - const T E5 = 1.24266094738807843860e-3; - const T E6 = 2.71155556874348757815e-5; - const T E7 = 2.01033439929228813265e-7; - - const T F0 = 1.414213562373095048801689e0; - const T F1 = 8.482908416595164588112026e-1; - const T F2 = 1.936480946950659106176712e-1; - const T F3 = 2.103693768272068968719679e-2; - const T F4 = 1.112800997078859844711555e-3; - const T F5 = 2.611088405080593625138020e-5; - const T F6 = 2.010321207683943062279931e-7; - const T F7 = 2.891024605872965461538222e-15; - - T abs_x = abs(x); - - if (abs_x <= 0.85) { - T r = 0.180625 - 0.25 * x * x; - T num = - (((((((A7 * r + A6) * r + A5) * r + A4) * r + A3) * r + A2) * r + A1) * - r + - A0); - T den = - (((((((B7 * r + B6) * r + B5) * r + B4) * r + B3) * r + B2) * r + B1) * - r + - B0); - return x * num / den; - } - - T r = sqrt(LN2 - log(1.0 - abs_x)); - - T num, den; - if (r <= 5.0) { - r = r - 1.6; - num = - (((((((C7 * r + C6) * r + C5) * r + C4) * r + C3) * r + C2) * r + C1) * - r + - C0); - den = - (((((((D7 * r + D6) * r + D5) * r + D4) * r + D3) * r + D2) * r + D1) * - r + - D0); - } else { - r = r - 5.0; - num = - (((((((E7 * r + E6) * r + E5) * r + E4) * r + E3) * r + E2) * r + E1) * - r + - E0); - den = - (((((((F7 * r + F6) * r + F5) * r + F4) * r + F3) * r + F2) * r + F1) * - r + - F0); - } - - if (x < 0) { - return -num / den; - } else { - return num / den; - } -} - -template -struct TruncatedNormal { - T mean, std; - T a_normal_cdf; - T b_normal_cdf; - TruncatedNormal(T mean, T std) : mean(mean), std(std) { - auto normal_cdf = [](T x) { - return (1.0 + std::erf(x / std::sqrt(2.0))) / 2.0; - }; - a_normal_cdf = normal_cdf(-2.0); - b_normal_cdf = normal_cdf(2.0); - } - - T operator()(T value) const { - auto p = a_normal_cdf + (b_normal_cdf - a_normal_cdf) * value; - return std::sqrt(2.0) * Erfinv(2 * p - 1) * std + mean; - } -}; - template void TruncatedGaussianRandomKernel(const Context& dev_ctx, const std::vector& shape, diff --git a/python/paddle/fluid/initializer.py b/python/paddle/fluid/initializer.py index bdc97eca0d8..37eff6d132d 100644 --- a/python/paddle/fluid/initializer.py +++ b/python/paddle/fluid/initializer.py @@ -17,7 +17,7 @@ from __future__ import print_function import math from . import framework from . import core -from .framework import _non_static_mode, default_main_program +from .framework import _non_static_mode, in_dygraph_mode, _in_legacy_dygraph, default_main_program, _current_expected_place import numpy as np from .core import VarDesc from . import unique_name @@ -417,7 +417,18 @@ class TruncatedNormalInitializer(Initializer): out_dtype = var.dtype out_var = var - if framework._non_static_mode(): + if in_dygraph_mode(): + out_var = _C_ops.final_state_truncated_gaussian_random( + var.shape, self._mean, self._std_dev, self._seed, out_dtype, + _current_expected_place()) + if var.dtype in [VarDesc.VarType.FP16, VarDesc.VarType.BF16]: + var_tmp = _C_ops.final_state_cast(out_var, var.dtype) + var_tmp._share_underline_tensor_to(var) + else: + out_var._share_underline_tensor_to(var) + return None + + if _in_legacy_dygraph(): out_var = _C_ops.truncated_gaussian_random( 'shape', var.shape, 'dtype', out_dtype, 'mean', self._mean, 'std', self._std_dev, 'seed', self._seed) diff --git a/python/paddle/fluid/tests/unittests/test_truncated_gaussian_random_op.py b/python/paddle/fluid/tests/unittests/test_truncated_gaussian_random_op.py index 4abeae77d26..fe28e0c9638 100644 --- a/python/paddle/fluid/tests/unittests/test_truncated_gaussian_random_op.py +++ b/python/paddle/fluid/tests/unittests/test_truncated_gaussian_random_op.py @@ -17,10 +17,13 @@ from __future__ import print_function import unittest import numpy +import paddle import paddle.fluid as fluid import paddle.fluid.core as core +from op_test import OpTest from paddle.fluid.op import Operator from paddle.fluid.executor import Executor +from paddle.fluid.framework import _test_eager_guard class TestTrunctedGaussianRandomOp(unittest.TestCase): @@ -33,15 +36,16 @@ class TestTrunctedGaussianRandomOp(unittest.TestCase): "std": 1., "seed": 10, } - self.outputs = ["Out"] def test_cpu(self): self.gaussian_random_test(place=fluid.CPUPlace()) + self.gaussian_random_test_eager(place=fluid.CPUPlace()) def test_gpu(self): if core.is_compiled_with_cuda(): self.gaussian_random_test(place=fluid.CUDAPlace(0)) + self.gaussian_random_test_eager(place=fluid.CUDAPlace(0)) def gaussian_random_test(self, place): @@ -64,6 +68,17 @@ class TestTrunctedGaussianRandomOp(unittest.TestCase): self.assertAlmostEqual(numpy.mean(tensor), .0, delta=0.1) self.assertAlmostEqual(numpy.var(tensor), 0.773, delta=0.1) + # TruncatedNormal.__call__ has no return value, so here call _C_ops api + # directly + def gaussian_random_test_eager(self, place): + with fluid.dygraph.guard(place): + with _test_eager_guard(): + out = paddle._C_ops.final_state_truncated_gaussian_random( + self.attrs["shape"], self.attrs["mean"], self.attrs["std"], + self.attrs["seed"], core.VarDesc.VarType.FP32, place) + self.assertAlmostEqual(numpy.mean(out.numpy()), .0, delta=0.1) + self.assertAlmostEqual(numpy.var(out.numpy()), 0.773, delta=0.1) + if __name__ == "__main__": unittest.main() diff --git a/python/paddle/utils/code_gen/api.yaml b/python/paddle/utils/code_gen/api.yaml index cb26fecb8e5..3266a43bd1d 100644 --- a/python/paddle/utils/code_gen/api.yaml +++ b/python/paddle/utils/code_gen/api.yaml @@ -1904,6 +1904,19 @@ func : trunc backward : trunc_grad +# python API: paddle.nn.initializer.TruncatedNormal +- api : truncated_gaussian_random + args : (int[] shape, float mean, float std, int seed, DataType dtype=DataType::FLOAT32, Place place={}) + output : Tensor + infer_meta : + func : TruncatedGaussianRandomInferMeta + param : [shape, mean, std, seed, dtype] + kernel : + func : truncated_gaussian_random + param : [shape, mean, std, seed, dtype] + backend : place + data_type : dtype + # unfold - api : unfold args : (Tensor x, int[] kernel_sizes, int[] strides, int[] paddings, int[] dilations) -- GitLab