未验证 提交 33185000 编写于 作者: zhouweiwei2014's avatar zhouweiwei2014 提交者: GitHub

add new API/OP:paddle.Tensor.exponential_ (#38256)

* add new API/OP:paddle.Tensor.exponential_

* fix CI
上级 c396ee65
/* 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
#ifdef __NVCC__
#include <curand_kernel.h>
#endif
#ifdef __HIPCC__
#include <hiprand_kernel.h>
#endif
#include "paddle/fluid/framework/tensor.h"
#include "paddle/fluid/platform/device/gpu/gpu_info.h"
#include "paddle/fluid/platform/device_context.h"
#include "paddle/fluid/platform/for_range.h"
#include "paddle/fluid/platform/hostdevice.h"
namespace paddle {
namespace distribution {
using Tensor = framework::Tensor;
template <typename T>
struct exponential_transform {
explicit exponential_transform(T lambda) : lambda_(lambda) {}
HOSTDEVICE inline T operator()(T val) const {
#if defined(__NVCC__) || defined(__HIPCC__)
if (std::is_same<T, double>::value) {
return static_cast<T>(-1.0) / lambda_ * log(val);
} else {
return static_cast<T>(-1.0) / lambda_ * __logf(val);
}
#else
return static_cast<T>(-1.0) / lambda_ * std::log(static_cast<T>(1.0) - val);
#endif
}
private:
T lambda_;
};
#if defined(__NVCC__) || defined(__HIPCC__)
template <typename T>
struct uniform_distribution;
template <typename T>
struct normal_distribution;
#if defined(__NVCC__)
template <>
struct uniform_distribution<float> {
__device__ inline float4 operator()(curandStatePhilox4_32_10_t *state) const {
return curand_uniform4(state);
}
static constexpr int kReturnsCount = 4;
};
template <>
struct uniform_distribution<double> {
__device__ inline double2 operator()(
curandStatePhilox4_32_10_t *state) const {
return curand_uniform2_double(state);
}
static constexpr int kReturnsCount = 2;
};
template <>
struct normal_distribution<float> {
__device__ inline float4 operator()(curandStatePhilox4_32_10_t *state) const {
return curand_normal4(state);
}
static constexpr int kReturnsCount = 4;
};
template <>
struct normal_distribution<double> {
__device__ inline double2 operator()(
curandStatePhilox4_32_10_t *state) const {
return curand_normal2_double(state);
}
static constexpr int kReturnsCount = 2;
};
#else
template <>
struct uniform_distribution<float> {
__device__ inline float4 operator()(
hiprandStatePhilox4_32_10_t *state) const {
return hiprand_uniform4(state);
}
static constexpr int kReturnsCount = 4;
};
template <>
struct uniform_distribution<double> {
__device__ inline double2 operator()(
hiprandStatePhilox4_32_10_t *state) const {
return hiprand_uniform2_double(state);
}
static constexpr int kReturnsCount = 2;
};
template <>
struct normal_distribution<float> {
__device__ inline float4 operator()(
hiprandStatePhilox4_32_10_t *state) const {
return hiprand_normal4(state);
}
static constexpr int kReturnsCount = 4;
};
template <>
struct normal_distribution<double> {
__device__ inline double2 operator()(
hiprandStatePhilox4_32_10_t *state) const {
return hiprand_normal2_double(state);
}
static constexpr int kReturnsCount = 2;
};
#endif
template <typename T, typename DistOp, typename TransformOp>
__global__ void DistributionKernel(size_t size, uint64_t seed, uint64_t offset,
DistOp dist, TransformOp trans,
T *out_data) {
size_t idx = static_cast<size_t>(blockIdx.x * blockDim.x + threadIdx.x);
int32_t returns_count = DistOp::kReturnsCount;
#if defined(__NVCC__)
curandStatePhilox4_32_10_t state;
curand_init(seed, idx, offset, &state);
#else
hiprandStatePhilox4_32_10_t state;
hiprand_init(seed, idx, offset, &state);
#endif
size_t total_thread = gridDim.x * blockDim.x;
for (size_t i = idx; i < size; i += total_thread * returns_count) {
auto random_tuple = dist(&state);
for (size_t j = 0; j < returns_count; j++) {
size_t index = i + j * total_thread;
if (index < size) {
auto random = static_cast<T>((&random_tuple.x)[j]);
out_data[index] = trans(random);
}
}
}
}
template <typename T, typename DistOp, typename TransformOp>
void distribution_and_transform(const platform::CUDADeviceContext &dev_ctx,
Tensor *out, DistOp dist, TransformOp trans) {
T *out_data = out->mutable_data<T>(dev_ctx.GetPlace());
auto size = out->numel();
int64_t device_id =
BOOST_GET_CONST(platform::CUDAPlace, dev_ctx.GetPlace()).GetDeviceId();
auto gen_cuda = framework::GetDefaultCUDAGenerator(device_id);
size_t block_size = 256;
size_t expect_grid_size = (size + block_size - 1) / block_size;
const auto &prop = platform::GetDeviceProperties(device_id);
size_t max_grid_size = (prop.maxThreadsPerMultiProcessor / block_size) *
prop.multiProcessorCount;
size_t grid_size =
expect_grid_size > max_grid_size ? max_grid_size : expect_grid_size;
size_t total_thread = block_size * grid_size;
size_t curand4_loop_times =
(size + 4 * total_thread - 1) / (4 * total_thread);
// 'increment' shoulde be multiple of 4
uint64_t increment = curand4_loop_times * 4;
auto seed_offset = gen_cuda->IncrementOffset(increment);
uint64_t seed = seed_offset.first;
uint64_t offset = seed_offset.second;
DistributionKernel<
T, DistOp, TransformOp><<<grid_size, block_size, 0, dev_ctx.stream()>>>(
size, seed, offset, dist, trans, out_data);
}
#endif
} // namespace distribution
} // namespace paddle
/* 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/exponential_op.h"
namespace paddle {
namespace operators {
class ExponentialOp : public framework::OperatorWithKernel {
public:
using framework::OperatorWithKernel::OperatorWithKernel;
void InferShape(framework::InferShapeContext *ctx) const override {
OP_INOUT_CHECK(ctx->HasInput("X"), "Input", "X", "ExponentialOp");
OP_INOUT_CHECK(ctx->HasOutput("Out"), "Output", "Out", "ExponentialOp");
auto dim = ctx->GetInputDim("X");
ctx->SetOutputDim("Out", dim);
}
protected:
framework::OpKernelType GetExpectedKernelType(
const framework::ExecutionContext &ctx) const override {
return framework::OpKernelType(
OperatorWithKernel::IndicateVarDataType(ctx, "X"), ctx.GetPlace());
}
};
class ExponentialOpMaker : public framework::OpProtoAndCheckerMaker {
public:
void Make() override {
AddComment(R"DOC(
This operator fills the input tensor with random values sampled from a
exponential distribution.
)DOC");
AddInput("X", "The input tensor.");
AddOutput("Out", "The output tensor of exponential OP.");
AddAttr<float>(
"lambda", "lambd parameter of exponential distribution. [default 1.0].")
.SetDefault(1.0f);
}
};
class ExponentialOpInferVarType
: public framework::PassInDtypeAndVarTypeToOutput {
protected:
std::unordered_map<std::string, std::string> &GetInputOutputWithSameType()
const override {
static std::unordered_map<std::string, std::string> m{{"X", /*->*/ "Out"}};
return m;
}
};
template <typename T>
class ExponentialKernel<platform::CPUDeviceContext, T>
: public framework::OpKernel<T> {
public:
void Compute(const framework::ExecutionContext &ctx) const override {
auto *out = ctx.Output<framework::Tensor>("Out");
T *out_data = out->mutable_data<T>(ctx.GetPlace());
T lambda = static_cast<T>(ctx.Attr<float>("lambda"));
int64_t size = out->numel();
auto gen = framework::DefaultCPUGenerator();
auto engine = gen->GetCPUEngine();
std::uniform_real_distribution<T> uniform(0.0, 1.0);
distribution::exponential_transform<T> trans(lambda);
for (int64_t i = 0; i < size; ++i) {
out_data[i] = trans(uniform(*engine));
}
}
};
class ExponentialGradOp : public framework::OperatorWithKernel {
public:
using framework::OperatorWithKernel::OperatorWithKernel;
void InferShape(framework::InferShapeContext *ctx) const override {
OP_INOUT_CHECK(ctx->HasInput(framework::GradVarName("Out")), "Input",
"Out_Grad", "ExponentialGradOp");
auto dout_dim = ctx->GetInputDim(framework::GradVarName("Out"));
ctx->SetOutputDim(framework::GradVarName("X"), dout_dim);
}
};
template <typename T>
class ExponentialGradOpMaker : public framework::SingleGradOpMaker<T> {
public:
using framework::SingleGradOpMaker<T>::SingleGradOpMaker;
protected:
void Apply(GradOpPtr<T> retv) const override {
retv->SetType("exponential_grad");
retv->SetInput(framework::GradVarName("Out"), this->OutputGrad("Out"));
retv->SetOutput(framework::GradVarName("X"), this->InputGrad("X"));
retv->SetAttrMap(this->Attrs());
}
};
} // namespace operators
} // namespace paddle
namespace ops = paddle::operators;
namespace plat = paddle::platform;
DECLARE_INPLACE_OP_INFERER(ExponentialInferer, {"X", "Out"});
DECLARE_INPLACE_OP_INFERER(ExponentialGradInferer,
{paddle::framework::GradVarName("Out"),
paddle::framework::GradVarName("X")});
REGISTER_OPERATOR(exponential, ops::ExponentialOp, ops::ExponentialOpMaker,
ops::ExponentialOpInferVarType,
ops::ExponentialGradOpMaker<paddle::framework::OpDesc>,
ops::ExponentialGradOpMaker<paddle::imperative::OpBase>,
ExponentialInferer);
REGISTER_OPERATOR(exponential_grad, ops::ExponentialGradOp,
ExponentialGradInferer);
REGISTER_OP_CPU_KERNEL(exponential,
ops::ExponentialKernel<plat::CPUDeviceContext, float>,
ops::ExponentialKernel<plat::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL(
exponential_grad, ops::ExponentialGradKernel<plat::CPUDeviceContext, float>,
ops::ExponentialGradKernel<plat::CPUDeviceContext, double>);
/* 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/exponential_op.h"
namespace paddle {
namespace operators {
template <typename T>
class ExponentialKernel<platform::CUDADeviceContext, T>
: public framework::OpKernel<T> {
public:
void Compute(const framework::ExecutionContext& ctx) const override {
framework::Tensor* out = ctx.Output<framework::Tensor>("Out");
auto& dev_cxt = ctx.template device_context<platform::CUDADeviceContext>();
T lambda = static_cast<T>(ctx.Attr<float>("lambda"));
distribution::uniform_distribution<T> dist;
distribution::exponential_transform<T> trans(lambda);
distribution::distribution_and_transform<T>(dev_cxt, out, dist, trans);
}
};
} // namespace operators
} // namespace paddle
namespace ops = paddle::operators;
namespace plat = paddle::platform;
REGISTER_OP_CUDA_KERNEL(
exponential, ops::ExponentialKernel<plat::CUDADeviceContext, float>,
ops::ExponentialKernel<plat::CUDADeviceContext, double>);
REGISTER_OP_CUDA_KERNEL(
exponential_grad,
ops::ExponentialGradKernel<plat::CUDADeviceContext, float>,
ops::ExponentialGradKernel<plat::CUDADeviceContext, double>);
// 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/fluid/operators/distribution_helper.h"
#include "paddle/fluid/operators/math/math_function.h"
namespace paddle {
namespace operators {
template <typename DeviceContext, typename T>
class ExponentialKernel;
template <typename DeviceContext, typename T>
class ExponentialGradKernel : public framework::OpKernel<T> {
public:
void Compute(const framework::ExecutionContext& ctx) const override {
auto* dx = ctx.Output<framework::Tensor>(framework::GradVarName("X"));
dx->mutable_data<T>(ctx.GetPlace());
math::SetConstant<DeviceContext, T> functor;
auto& dev_ctx = ctx.template device_context<DeviceContext>();
functor(dev_ctx, dx, static_cast<T>(0));
}
};
} // namespace operators
} // namespace paddle
# 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.
import unittest
import paddle
import numpy as np
from op_test import OpTest
paddle.enable_static()
class TestExponentialOp1(OpTest):
def setUp(self):
self.op_type = "exponential"
self.config()
self.attrs = {"lambda": self.lam}
self.inputs = {'X': np.empty([1024, 1024], dtype=self.dtype)}
self.outputs = {'Out': np.ones([1024, 1024], dtype=self.dtype)}
def config(self):
self.lam = 0.5
self.dtype = "float64"
def test_check_output(self):
self.check_output_customized(self.verify_output)
def verify_output(self, outs):
hist1, _ = np.histogram(outs[0], range=(0, 5))
hist1 = hist1.astype("float32")
hist1 = hist1 / float(outs[0].size)
data_np = np.random.exponential(1. / self.lam, [1024, 1024])
hist2, _ = np.histogram(data_np, range=(0, 5))
hist2 = hist2.astype("float32")
hist2 = hist2 / float(data_np.size)
self.assertTrue(
np.allclose(
hist1, hist2, rtol=0.02),
"actual: {}, expected: {}".format(hist1, hist2))
def test_check_grad_normal(self):
self.check_grad(
['X'],
'Out',
user_defined_grads=[np.zeros(
[1024, 1024], dtype=self.dtype)],
user_defined_grad_outputs=[
np.random.rand(1024, 1024).astype(self.dtype)
])
class TestExponentialOp2(TestExponentialOp1):
def config(self):
self.lam = 0.25
self.dtype = "float32"
class TestExponentialAPI(unittest.TestCase):
def test_static(self):
with paddle.static.program_guard(paddle.static.Program(),
paddle.static.Program()):
x_np = np.full([10, 10], -1.)
x = paddle.static.data(name="X", shape=[10, 10], dtype='float64')
x.exponential_(1.0)
exe = paddle.static.Executor()
out = exe.run(paddle.static.default_main_program(),
feed={"X": x_np},
fetch_list=[x])
self.assertTrue(np.min(out) >= 0)
def test_dygraph(self):
paddle.disable_static()
x = paddle.full([10, 10], -1., dtype='float32')
x.exponential_(0.5)
self.assertTrue(np.min(x.numpy()) >= 0)
paddle.enable_static()
# Test GPU Fixed random number, which is generated by 'curandStatePhilox4_32_10_t'
def test_fixed_random_number(self):
if not paddle.is_compiled_with_cuda():
return
# Note(zhouwei): The Number of threads is determined by
# 'multiProcessorCount * maxThreadsPerMultiProcessor'. So, different
# GPU have different number of threads, which result in different
# random value. Only test on V100 GPU here.
if not "V100" in paddle.device.cuda.get_device_name():
return
print("Test Fixed Random number on V100 GPU------>")
paddle.disable_static()
paddle.set_device('gpu')
paddle.seed(2021)
x = paddle.empty([64, 3, 1024, 1024], dtype="float32")
x.exponential_(1.0)
x_np = x.numpy()
expect = [
0.80073667, 0.2249291, 0.07734892, 1.25392, 0.14013891, 0.45736602,
1.9735607, 0.30490234, 0.57100505, 0.8115938
]
self.assertTrue(np.allclose(x_np[0, 0, 0, 0:10], expect))
expect = [
1.4296371e+00, 9.5411777e-01, 5.2575850e-01, 2.4805880e-01,
1.2322118e-04, 8.4604341e-01, 2.1111444e-01, 1.4143821e+00,
2.8194717e-01, 1.1360573e+00
]
self.assertTrue(np.allclose(x_np[16, 1, 300, 200:210], expect))
expect = [
1.3448033, 0.35146526, 1.7380928, 0.32012638, 0.10396296,
0.51344526, 0.15308502, 0.18712929, 0.03888268, 0.20771872
]
self.assertTrue(np.allclose(x_np[32, 1, 600, 500:510], expect))
expect = [
0.5107464, 0.20970327, 2.1986802, 1.580056, 0.31036147, 0.43966478,
0.9056133, 0.30119267, 1.4797124, 1.4319834
]
self.assertTrue(np.allclose(x_np[48, 2, 900, 800:810], expect))
expect = [
3.4640615, 1.1019983, 0.41195083, 0.22681557, 0.291846, 0.53617656,
1.5791925, 2.4645927, 0.04094889, 0.9057725
]
self.assertTrue(np.allclose(x_np[63, 2, 1023, 1000:1010], expect))
x = paddle.empty([10, 10], dtype="float32")
x.exponential_(3.0)
x_np = x.numpy()
expect = [
0.02831675, 0.1691551, 0.6798956, 0.69347525, 0.0243443, 0.22180498,
0.30574575, 0.9839696, 0.2834912, 0.59420055
]
self.assertTrue(np.allclose(x_np[5, 0:10], expect))
x = paddle.empty([16, 2, 1024, 768], dtype="float64")
x.exponential_(0.25)
x_np = x.numpy()
expect = [
10.0541229, 12.67860643, 1.09850734, 7.35289643, 2.65471225,
3.86217432, 2.97902086, 2.92744479, 2.67927152, 0.19667352
]
self.assertTrue(np.allclose(x_np[0, 0, 0, 100:110], expect))
expect = [
0.68328125, 3.1454553, 0.92158376, 1.95842188, 1.05296941,
12.93242051, 5.20255978, 3.3588624, 1.57377174, 5.73194183
]
self.assertTrue(np.allclose(x_np[4, 0, 300, 190:200], expect))
expect = [
1.37973974, 3.45036798, 7.94625406, 1.62610973, 0.31032122,
4.13596493, 1.98494535, 1.13207041, 8.30592769, 2.81460147
]
self.assertTrue(np.allclose(x_np[8, 1, 600, 300:310], expect))
expect = [
2.27710811, 12.25003028, 2.96409124, 4.72405788, 0.67917249,
4.35856718, 0.46870976, 2.31120149, 9.61595826, 4.64446271
]
self.assertTrue(np.allclose(x_np[12, 1, 900, 500:510], expect))
expect = [
0.95883744, 1.57316361, 15.22524512, 20.49559882, 13.70008548,
3.29430143, 3.90390424, 0.9146657, 0.80972249, 0.33376219
]
self.assertTrue(np.allclose(x_np[15, 1, 1023, 750:760], expect))
x = paddle.empty([512, 768], dtype="float64")
x.exponential_(0.3)
x_np = x.numpy()
expect = [
8.79266704, 4.79596009, 2.75480243, 6.04670011, 0.35379556,
0.76864868, 3.17428251, 0.26556859, 12.22485885, 10.51690383
]
self.assertTrue(np.allclose(x_np[0, 200:210], expect))
expect = [
5.6341126, 0.52243418, 5.36410796, 6.83672002, 11.9243311,
5.85985566, 5.75169548, 0.13877972, 6.1348385, 3.82436519
]
self.assertTrue(np.allclose(x_np[300, 400:410], expect))
expect = [
4.94883581, 0.56345306, 0.85841585, 1.92287801, 6.10036656,
1.19524847, 3.64735434, 5.19618716, 2.57467974, 3.49152791
]
self.assertTrue(np.allclose(x_np[500, 700:710], expect))
x = paddle.empty([10, 10], dtype="float64")
x.exponential_(4.0)
x_np = x.numpy()
expect = [
0.15713826, 0.56395964, 0.0680941, 0.00316643, 0.27046853,
0.19852724, 0.12776634, 0.09642974, 0.51977551, 1.33739699
]
self.assertTrue(np.allclose(x_np[5, 0:10], expect))
paddle.enable_static()
if __name__ == "__main__":
unittest.main()
......@@ -227,6 +227,7 @@ from .random import randint # noqa: F401
from .random import randint_like # noqa: F401
from .random import randperm # noqa: F401
from .random import poisson # noqa: F401
from .random import exponential_ # noqa: F401
from .search import argmax # noqa: F401
from .search import argmin # noqa: F401
from .search import argsort # noqa: F401
......@@ -453,6 +454,7 @@ tensor_method_func = [ #noqa
'angle',
'moveaxis',
'repeat_interleave',
'exponential_',
]
#this list used in math_op_patch.py for magic_method bind
......
......@@ -100,13 +100,13 @@ def poisson(x, name=None):
.. code-block:: python
import paddle
paddle.set_device('gpu')
paddle.set_device('cpu')
paddle.seed(2021)
x = paddle.uniform([2,3], min=1.0, max=5.0)
out = paddle.poisson(x)
# [[0., 5., 1.],
# [4., 3., 0.]])
# [[2., 1., 4.],
# [4., 5., 1.]]
"""
......@@ -980,3 +980,49 @@ def rand(shape, dtype=None, name=None):
"""
return uniform(shape, dtype, min=0.0, max=1.0, name=name)
def exponential_(x, lam=1.0, name=None):
"""
This inplace OP fill input Tensor ``x`` with random number from a Exponential Distribution.
``lam`` is :math:`\lambda` parameter of Exponential Distribution.
.. math::
f(x) = \lambda e^{-\lambda x}
Args:
x(Tensor): Input tensor. The data type should be float32, float64.
lam(float): :math:`\lambda` parameter of Exponential Distribution.
name(str, optional): The default value is None. Normally there is no
need for user to set this property. For more information, please
refer to :ref:`api_guide_Name`.
Returns:
Tensor: Input Tensor ``x``.
Examples:
.. code-block:: python
import paddle
paddle.set_device('cpu')
paddle.seed(100)
x = paddle.empty([2,3])
x.exponential_()
# [[0.80643415, 0.23211166, 0.01169797],
# [0.72520673, 0.45208144, 0.30234432]]
"""
if in_dygraph_mode():
return _C_ops.exponential_(x, "lambda", lam)
check_variable_and_dtype(x, "x", ["float32", "float64"], "exponential")
helper = LayerHelper("exponential", **locals())
helper.append_op(
type='exponential',
inputs={"X": x},
outputs={'Out': x},
attrs={"lambda": lam})
return x
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册