diff --git a/paddle/fluid/operators/mv_op.cc b/paddle/fluid/operators/mv_op.cc new file mode 100644 index 0000000000000000000000000000000000000000..1339982adaab162056bdefd3eecb405e95188a0d --- /dev/null +++ b/paddle/fluid/operators/mv_op.cc @@ -0,0 +1,125 @@ +/* Copyright (c) 2020 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/mv_op.h" +namespace paddle { +namespace operators { + +class MVOpMaker : public framework::OpProtoAndCheckerMaker { + public: + void Make() override { + AddInput("X", "The matrix input of mv op"); + AddInput("Vec", "The vector input of mv op"); + AddOutput("Out", "The output of mv op"); + AddComment(R"DOC( +MV Operator. + +This operator is used to perform matrix vector multiplication +of the input tensors `X` and `Vec`. +)DOC"); + } +}; + +class MVOp : public framework::OperatorWithKernel { + public: + using framework::OperatorWithKernel::OperatorWithKernel; + + protected: + void InferShape(framework::InferShapeContext *context) const override { + OP_INOUT_CHECK(context->HasInput("X"), "Input", "X", "mv"); + OP_INOUT_CHECK(context->HasInput("Vec"), "Input", "Vec", "mv"); + OP_INOUT_CHECK(context->HasOutput("Out"), "Output", "Out", "mv"); + + auto dim_x = context->GetInputDim("X"); + auto dim_y = context->GetInputDim("Vec"); + PADDLE_ENFORCE_EQ( + dim_x.size(), 2, + platform::errors::InvalidArgument( + "The rank of input X should be 2, but is %d", dim_x.size())); + PADDLE_ENFORCE_EQ( + dim_y.size(), 1, + platform::errors::InvalidArgument( + "The rank of input Vec should be 1, but is %d", dim_y.size())); + PADDLE_ENFORCE_EQ(dim_x[1] == dim_y[0], true, + platform::errors::InvalidArgument( + "The length of input X' second dim should equal the " + "length of input Vec," + " but X[%d, %d], Vec[%d]", + dim_x[0], dim_x[1], dim_y[0])); + + framework::DDim dim_out = framework::make_ddim({dim_x[0]}); + + context->SetOutputDim("Out", dim_out); + context->ShareLoD("X", /*->*/ "Out"); + } +}; + +template +class MVOpGradMaker : public framework::SingleGradOpMaker { + public: + using framework::SingleGradOpMaker::SingleGradOpMaker; + + protected: + void Apply(GradOpPtr retv) const override { + retv->SetType("mv_grad"); + retv->SetInput("X", this->Input("X")); + retv->SetInput("Vec", this->Input("Vec")); + retv->SetInput(framework::GradVarName("Out"), this->OutputGrad("Out")); + retv->SetOutput(framework::GradVarName("X"), this->InputGrad("X")); + retv->SetOutput(framework::GradVarName("Vec"), this->InputGrad("Vec")); + } +}; + +class MVOpGrad : public framework::OperatorWithKernel { + public: + using framework::OperatorWithKernel::OperatorWithKernel; + + protected: + void InferShape(framework::InferShapeContext *context) const override { + OP_INOUT_CHECK(context->HasInput("X"), "Input", "X", "mv"); + OP_INOUT_CHECK(context->HasInput("Vec"), "Input", "Vec", "mv"); + OP_INOUT_CHECK(context->HasInput(framework::GradVarName("Out")), "Input", + "Out@GRAD", "mv"); + auto x_dims = context->GetInputDim("X"); + auto vec_dims = context->GetInputDim("Vec"); + + auto x_grad_name = framework::GradVarName("X"); + auto vec_grad_name = framework::GradVarName("Vec"); + + if (context->HasOutput(x_grad_name)) { + context->SetOutputDim(x_grad_name, x_dims); + } + if (context->HasOutput(vec_grad_name)) { + context->SetOutputDim(vec_grad_name, vec_dims); + } + } +}; + +} // namespace operators +} // namespace paddle + +namespace ops = paddle::operators; +namespace plat = paddle::platform; + +REGISTER_OPERATOR(mv, ops::MVOp, ops::MVOpMaker, + ops::MVOpGradMaker, + ops::MVOpGradMaker); +REGISTER_OPERATOR(mv_grad, ops::MVOpGrad); + +REGISTER_OP_CPU_KERNEL( + mv, ops::MVKernel, + ops::MVKernel); +REGISTER_OP_CPU_KERNEL( + mv_grad, ops::MVGradKernel, + ops::MVGradKernel); diff --git a/paddle/fluid/operators/mv_op.cu b/paddle/fluid/operators/mv_op.cu new file mode 100644 index 0000000000000000000000000000000000000000..9a16fe025cd71457faade38f92f56e56c26b3b32 --- /dev/null +++ b/paddle/fluid/operators/mv_op.cu @@ -0,0 +1,95 @@ +/* Copyright (c) 2020 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/mv_op.h" +#include "paddle/fluid/platform/gpu_launch_param_config.h" + +namespace paddle { +namespace operators { + +template +__global__ void MVGradCUDAKernel(const int m, const int n, const T *dout, + const T *vec, T *dx) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + for (; idx < m * n; idx += blockDim.x * gridDim.x) { + int i = idx / n; + int j = idx % n; + dx[idx] = dout[i] * vec[j]; + } +} + +// Using dimensional constraints on matrix multiplication, it is +// straight-forward to check the following table for when X and Y +// are both matrices. +// +// dX = | dOut Vec^T +// dVec = | X^T dOut +template +class MVGradKernel + : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext &context) const override { + auto *x = context.Input("X"); + auto *vec = context.Input("Vec"); + auto *dout = + context.Input(framework::GradVarName("Out")); + auto *dx = context.Output(framework::GradVarName("X")); + auto *dvec = + context.Output(framework::GradVarName("Vec")); + + auto dim_x = x->dims(); + int m = dim_x[0]; + int n = dim_x[1]; + + dx->Resize(framework::make_ddim({m * n})); + + // get data ptr + const T *x_data = x->data(); + const T *vec_data = vec->data(); + const T *dout_data = dout->data(); + + T *dx_data = dx->mutable_data(context.GetPlace()); + T *dvec_data = dvec->mutable_data(context.GetPlace()); + + auto &dev_ctx = + context.template device_context(); + auto blas = math::GetBlas(dev_ctx); + + // calculate dx + auto stream = context.cuda_device_context().stream(); + auto config = GetGpuLaunchConfig1D(dev_ctx, m * n); + MVGradCUDAKernel< + T><<>>( + m, n, dout_data, vec_data, dx_data); + + dx->Resize(framework::make_ddim({m, n})); + + // calculate dvec + blas.GEMV(true, dim_x[0], dim_x[1], static_cast(1), x_data, dout_data, + static_cast(0), dvec_data); + } +}; + +} // namespace operators +} // namespace paddle + +namespace ops = paddle::operators; +namespace plat = paddle::platform; + +REGISTER_OP_CUDA_KERNEL( + mv, ops::MVKernel, + ops::MVKernel); +REGISTER_OP_CUDA_KERNEL( + mv_grad, ops::MVGradKernel, + ops::MVGradKernel); diff --git a/paddle/fluid/operators/mv_op.h b/paddle/fluid/operators/mv_op.h new file mode 100644 index 0000000000000000000000000000000000000000..3c63f3640ff46f5592a244a930a191a23959baf7 --- /dev/null +++ b/paddle/fluid/operators/mv_op.h @@ -0,0 +1,105 @@ +/* Copyright (c) 2020 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 +#include +#include +#include "paddle/fluid/framework/op_registry.h" +#include "paddle/fluid/operators/math/blas.h" +#ifdef PADDLE_WITH_MKLDNN +#include "paddle/fluid/platform/mkldnn_helper.h" +#endif + +namespace paddle { +namespace operators { + +using Tensor = framework::Tensor; + +template +class MVKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext &context) const override { + auto *x = context.Input("X"); + auto *vec = context.Input("Vec"); + + auto *out = context.Output("Out"); + + auto dim_x = x->dims(); + + // get data ptr + const T *x_data = x->data(); + const T *vec_data = vec->data(); + T *out_data = out->mutable_data(context.GetPlace()); + + auto &dev_ctx = context.template device_context(); + auto blas = math::GetBlas(dev_ctx); + + blas.GEMV(false, dim_x[0], dim_x[1], static_cast(1), x_data, vec_data, + static_cast(0), out_data); + } +}; + +// Using dimensional constraints on matrix multiplication, it is +// straight-forward to check the following table for when X and Y +// are both matrices. +// +// dX = | dOut vec^T +// dVec = | X^T dOut +template +class MVGradKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext &context) const override { + auto *x = context.Input("X"); + auto *vec = context.Input("Vec"); + auto *dout = + context.Input(framework::GradVarName("Out")); + auto *dx = context.Output(framework::GradVarName("X")); + auto *dvec = + context.Output(framework::GradVarName("Vec")); + + auto dim_x = x->dims(); + int m = dim_x[0]; + int n = dim_x[1]; + + dx->Resize(framework::make_ddim({m * n})); + + // get data ptr + const T *x_data = x->data(); + const T *vec_data = vec->data(); + const T *dout_data = dout->data(); + + T *dx_data = dx->mutable_data(context.GetPlace()); + T *dvec_data = dvec->mutable_data(context.GetPlace()); + + auto &dev_ctx = context.template device_context(); + auto blas = math::GetBlas(dev_ctx); + + // calculate dx + for (int i = 0; i < m; ++i) { + for (int j = 0; j < n; ++j) + dx_data[i * n + j] = dout_data[i] * vec_data[j]; + } + + dx->Resize(framework::make_ddim({m, n})); + + // calculate dvec + blas.GEMV(true, dim_x[0], dim_x[1], static_cast(1), x_data, dout_data, + static_cast(0), dvec_data); + } +}; + +} // namespace operators +} // namespace paddle diff --git a/python/paddle/__init__.py b/python/paddle/__init__.py index 661471599cb080da7a65c11fecc339830f2c00ee..1e0dc0e07b464e90262e1773cba844e9e58033c5 100755 --- a/python/paddle/__init__.py +++ b/python/paddle/__init__.py @@ -90,6 +90,7 @@ from .tensor.linalg import cholesky #DEFINE_ALIAS # from .tensor.linalg import tensordot #DEFINE_ALIAS from .tensor.linalg import bmm #DEFINE_ALIAS from .tensor.linalg import histogram #DEFINE_ALIAS +from .tensor.linalg import mv #DEFINE_ALIAS from .tensor.logic import equal #DEFINE_ALIAS from .tensor.logic import greater_equal #DEFINE_ALIAS from .tensor.logic import greater_than #DEFINE_ALIAS diff --git a/python/paddle/fluid/tests/unittests/test_mv_op.py b/python/paddle/fluid/tests/unittests/test_mv_op.py new file mode 100644 index 0000000000000000000000000000000000000000..6b930e59aa554c57ba1ecae2c01aaefabbe578e9 --- /dev/null +++ b/python/paddle/fluid/tests/unittests/test_mv_op.py @@ -0,0 +1,94 @@ +#Copyright (c) 2020 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. + +from __future__ import print_function + +import unittest +import numpy as np +import paddle +import paddle.fluid as fluid +import paddle.fluid.layers as layers +import paddle.fluid.core as core +from op_test import OpTest + + +class TestMVOp(OpTest): + def setUp(self): + self.op_type = "mv" + self.init_config() + self.inputs = {'X': self.x, 'Vec': self.vec} + self.outputs = {'Out': np.dot(self.x, self.vec)} + + def test_check_output(self): + self.check_output() + + def test_check_grad(self): + self.check_grad(['X', 'Vec'], 'Out') + + def init_config(self): + self.x = np.random.random((5, 100)).astype("float64") + self.vec = np.random.random((100)).astype("float64") + + +class TestMVAPI(unittest.TestCase): + def test_dygraph_api_out(self): + paddle.disable_static() + + self.x_data = np.random.random((5, 100)).astype("float64") + self.x = paddle.to_tensor(self.x_data) + self.vec_data = np.random.random((100)).astype("float64") + self.vec = paddle.to_tensor(self.vec_data) + z = paddle.mv(self.x, self.vec) + np_z = z.numpy() + z_expected = np.array(np.dot(self.x_data, self.vec_data)) + self.assertTrue(np.allclose(np_z, z_expected)) + + paddle.enable_static() + + def test_static_graph(self): + paddle.enable_static() + + self.input_x = np.random.rand(5, 100).astype("float64") + self.input_vec = np.random.rand(100).astype("float64") + + data_x = paddle.static.data("x", shape=[5, 100], dtype="float64") + data_vec = paddle.static.data("vec", shape=[100], dtype="float64") + result_vec = paddle.mv(data_x, data_vec) + self.place = paddle.CPUPlace() + exe = paddle.static.Executor(self.place) + res, = exe.run(feed={"x": self.input_x, + "vec": self.input_vec}, + fetch_list=[result_vec]) + z_expected = np.array(np.dot(self.input_x, self.input_vec)) + self.assertTrue(np.allclose(res, z_expected)) + + +class TestMVError(unittest.TestCase): + def test_input(self): + def test_shape(): + paddle.enable_static() + + self.input_x = np.random.rand(5, 100).astype("float64") + self.input_vec = np.random.rand(100).astype("float64") + + data_x = paddle.static.data("x", shape=[5, 100], dtype="float64") + data_vec = paddle.static.data( + "vec", shape=[100, 2], dtype="float64") + result_vec = paddle.mv(data_x, data_vec) + + self.assertRaises(ValueError, test_shape) + + +if __name__ == '__main__': + unittest.main() diff --git a/python/paddle/tensor/__init__.py b/python/paddle/tensor/__init__.py index a713663e1822d4af2d09efb2986aeb513930bbc0..2df9473c4b2dabee1039d15ba42f4295a553c331 100755 --- a/python/paddle/tensor/__init__.py +++ b/python/paddle/tensor/__init__.py @@ -56,6 +56,7 @@ from .linalg import cholesky #DEFINE_ALIAS # from .linalg import tensordot #DEFINE_ALIAS from .linalg import bmm #DEFINE_ALIAS from .linalg import histogram #DEFINE_ALIAS +from .linalg import mv #DEFINE_ALIAS from .logic import equal #DEFINE_ALIAS from .logic import greater_equal #DEFINE_ALIAS from .logic import greater_than #DEFINE_ALIAS diff --git a/python/paddle/tensor/linalg.py b/python/paddle/tensor/linalg.py index 67e3ce21ffba0c312eb01163cdf32f87c6433ee1..f27cfba487d78f284408815eaba933b18f303df9 100644 --- a/python/paddle/tensor/linalg.py +++ b/python/paddle/tensor/linalg.py @@ -32,7 +32,8 @@ __all__ = [ 'cholesky', # 'tensordot', 'bmm', - 'histogram' + 'histogram', + 'mv' ] @@ -920,3 +921,64 @@ def histogram(input, bins=100, min=0, max=0): 'min': min, 'max': max}) return out + + +def mv(x, vec, name=None): + """ + Performs a matrix-vector product of the matrix x and the vector vec. + + Args: + x (Variable): A tensor with shape :math:`[M, N]` , The data type of the input Tensor x + should be one of float32, float64. + vec (Variable): A tensor with shape :math:`[N]` , The data type of the input Tensor x + should be one of float32, float64. + 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: The tensor which is producted by x and vec. + + Examples: + .. code-block:: python + + # x: [M, N], vec: [N] + # paddle.mv(x, vec) # out: [M] + + import numpy as np + import paddle + + paddle.disable_static() + x_data = np.array([[2, 1, 3], [3, 0, 1]]).astype("float64") + x = paddle.to_tensor(x_data) + vec_data = np.array([3, 5, 1]) + vec = paddle.to_tensor(vec_data).astype("float64") + out = paddle.mv(x, vec) + paddle.enable_static() + """ + if in_dygraph_mode(): + out = core.ops.mv(x, vec) + return out + + def __check_input(x, vec): + var_names = {'x': x, 'vec': vec} + for name, val in var_names.items(): + check_variable_and_dtype(val, name, ['float32', 'float64'], 'mv') + x_shape = list(x.shape) + vec_shape = list(vec.shape) + if len(x_shape) != 2: + raise ValueError( + "x should be 2-dimensional. But received x's dimention: {}". + format(x_shape)) + if len(vec_shape) != 1: + raise ValueError( + "vec should be 1-dimensional. But received vec's dimention: {}". + format(vec_shape)) + + __check_input(x, vec) + + helper = LayerHelper('mv', **locals()) + out = helper.create_variable_for_type_inference(dtype=x.dtype) + helper.append_op( + type='mv', inputs={'X': x, + 'Vec': vec}, outputs={'Out': out}) + return out