未验证 提交 d39e7896 编写于 作者: C Chen Weihang 提交者: GitHub

[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
上级 e449f2aa
...@@ -21,10 +21,141 @@ ...@@ -21,10 +21,141 @@
#include "paddle/phi/backends/cpu/cpu_context.h" #include "paddle/phi/backends/cpu/cpu_context.h"
#include "paddle/phi/core/kernel_registry.h" #include "paddle/phi/core/kernel_registry.h"
#include "paddle/fluid/framework/generator.h"
namespace phi { namespace phi {
// reference: https://gist.github.com/lakshayg/d80172fe5ae3c5d2c2aedb53c250320e
template <typename T>
T Erfinv(T x) {
if (x < -1 || x > 1) {
return std::numeric_limits<T>::quiet_NaN();
} else if (x == 1.0) {
return std::numeric_limits<T>::infinity();
} else if (x == -1.0) {
return -std::numeric_limits<T>::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 <typename T>
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 <typename T, typename Context> template <typename T, typename Context>
void TruncatedGaussianRandomKernel(const Context& dev_ctx, void TruncatedGaussianRandomKernel(const Context& dev_ctx,
const std::vector<int>& shape, const std::vector<int>& shape,
...@@ -42,7 +173,13 @@ void TruncatedGaussianRandomKernel(const Context& dev_ctx, ...@@ -42,7 +173,13 @@ void TruncatedGaussianRandomKernel(const Context& dev_ctx,
TruncatedNormal<T> truncated_normal(mean, std); TruncatedNormal<T> truncated_normal(mean, std);
int64_t size = tensor->numel(); int64_t size = tensor->numel();
auto engine = paddle::framework::GetCPURandomEngine(seed); std::shared_ptr<std::mt19937_64> engine;
if (seed) {
engine = std::make_shared<std::mt19937_64>();
engine->seed(seed);
} else {
engine = dev_ctx.GetGenerator()->GetCPUEngine();
}
for (int64_t i = 0; i < size; ++i) { for (int64_t i = 0; i < size; ++i) {
data[i] = truncated_normal(dist(*engine)); data[i] = truncated_normal(dist(*engine));
} }
......
...@@ -24,8 +24,6 @@ ...@@ -24,8 +24,6 @@
#include "paddle/phi/core/dense_tensor.h" #include "paddle/phi/core/dense_tensor.h"
#include "paddle/phi/core/kernel_registry.h" #include "paddle/phi/core/kernel_registry.h"
#include "paddle/fluid/framework/generator.h"
namespace phi { namespace phi {
template <typename T> template <typename T>
...@@ -106,8 +104,7 @@ void TruncatedGaussianRandomKernel(const Context& dev_ctx, ...@@ -106,8 +104,7 @@ void TruncatedGaussianRandomKernel(const Context& dev_ctx,
thrust::counting_iterator<int64_t> index_sequence_begin(0); thrust::counting_iterator<int64_t> index_sequence_begin(0);
int64_t size = tensor->numel(); int64_t size = tensor->numel();
int device_id = dev_ctx.GetPlace().GetDeviceId(); auto gen_cuda = dev_ctx.GetGenerator();
auto gen_cuda = paddle::framework::GetDefaultCUDAGenerator(device_id);
if (gen_cuda->GetIsInitPy() && seed_flag) { if (gen_cuda->GetIsInitPy() && seed_flag) {
auto seed_offset = gen_cuda->IncrementOffset(1); auto seed_offset = gen_cuda->IncrementOffset(1);
......
...@@ -14,149 +14,11 @@ ...@@ -14,149 +14,11 @@
#pragma once #pragma once
#include <limits>
#include <random>
#include "paddle/phi/common/int_array.h" #include "paddle/phi/common/int_array.h"
#include "paddle/phi/core/dense_tensor.h" #include "paddle/phi/core/dense_tensor.h"
#include "paddle/phi/core/device_context.h"
#include "paddle/phi/infermeta/nullary.h"
namespace phi { namespace phi {
// reference: https://gist.github.com/lakshayg/d80172fe5ae3c5d2c2aedb53c250320e
template <typename T>
T Erfinv(T x) {
if (x < -1 || x > 1) {
return std::numeric_limits<T>::quiet_NaN();
} else if (x == 1.0) {
return std::numeric_limits<T>::infinity();
} else if (x == -1.0) {
return -std::numeric_limits<T>::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 <typename T>
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 <typename T, typename Context> template <typename T, typename Context>
void TruncatedGaussianRandomKernel(const Context& dev_ctx, void TruncatedGaussianRandomKernel(const Context& dev_ctx,
const std::vector<int>& shape, const std::vector<int>& shape,
......
...@@ -17,7 +17,7 @@ from __future__ import print_function ...@@ -17,7 +17,7 @@ from __future__ import print_function
import math import math
from . import framework from . import framework
from . import core 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 import numpy as np
from .core import VarDesc from .core import VarDesc
from . import unique_name from . import unique_name
...@@ -417,7 +417,18 @@ class TruncatedNormalInitializer(Initializer): ...@@ -417,7 +417,18 @@ class TruncatedNormalInitializer(Initializer):
out_dtype = var.dtype out_dtype = var.dtype
out_var = var 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( out_var = _C_ops.truncated_gaussian_random(
'shape', var.shape, 'dtype', out_dtype, 'mean', self._mean, 'shape', var.shape, 'dtype', out_dtype, 'mean', self._mean,
'std', self._std_dev, 'seed', self._seed) 'std', self._std_dev, 'seed', self._seed)
......
...@@ -17,10 +17,13 @@ from __future__ import print_function ...@@ -17,10 +17,13 @@ from __future__ import print_function
import unittest import unittest
import numpy import numpy
import paddle
import paddle.fluid as fluid import paddle.fluid as fluid
import paddle.fluid.core as core import paddle.fluid.core as core
from op_test import OpTest
from paddle.fluid.op import Operator from paddle.fluid.op import Operator
from paddle.fluid.executor import Executor from paddle.fluid.executor import Executor
from paddle.fluid.framework import _test_eager_guard
class TestTrunctedGaussianRandomOp(unittest.TestCase): class TestTrunctedGaussianRandomOp(unittest.TestCase):
...@@ -33,15 +36,16 @@ class TestTrunctedGaussianRandomOp(unittest.TestCase): ...@@ -33,15 +36,16 @@ class TestTrunctedGaussianRandomOp(unittest.TestCase):
"std": 1., "std": 1.,
"seed": 10, "seed": 10,
} }
self.outputs = ["Out"] self.outputs = ["Out"]
def test_cpu(self): def test_cpu(self):
self.gaussian_random_test(place=fluid.CPUPlace()) self.gaussian_random_test(place=fluid.CPUPlace())
self.gaussian_random_test_eager(place=fluid.CPUPlace())
def test_gpu(self): def test_gpu(self):
if core.is_compiled_with_cuda(): if core.is_compiled_with_cuda():
self.gaussian_random_test(place=fluid.CUDAPlace(0)) self.gaussian_random_test(place=fluid.CUDAPlace(0))
self.gaussian_random_test_eager(place=fluid.CUDAPlace(0))
def gaussian_random_test(self, place): def gaussian_random_test(self, place):
...@@ -64,6 +68,17 @@ class TestTrunctedGaussianRandomOp(unittest.TestCase): ...@@ -64,6 +68,17 @@ class TestTrunctedGaussianRandomOp(unittest.TestCase):
self.assertAlmostEqual(numpy.mean(tensor), .0, delta=0.1) self.assertAlmostEqual(numpy.mean(tensor), .0, delta=0.1)
self.assertAlmostEqual(numpy.var(tensor), 0.773, 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__": if __name__ == "__main__":
unittest.main() unittest.main()
...@@ -1904,6 +1904,19 @@ ...@@ -1904,6 +1904,19 @@
func : trunc func : trunc
backward : trunc_grad 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 # unfold
- api : unfold - api : unfold
args : (Tensor x, int[] kernel_sizes, int[] strides, int[] paddings, int[] dilations) args : (Tensor x, int[] kernel_sizes, int[] strides, int[] paddings, int[] dilations)
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册