add mv op(c++, python, unit test) (#27024)

paddle/fluid/operators/mv_op.cc

+/* 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 <typename T>
+class MVOpGradMaker : public framework::SingleGradOpMaker<T> {
+ public:
+  using framework::SingleGradOpMaker<T>::SingleGradOpMaker;
+
+ protected:
+  void Apply(GradOpPtr<T> 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<paddle::framework::OpDesc>,
+                  ops::MVOpGradMaker<paddle::imperative::OpBase>);
+REGISTER_OPERATOR(mv_grad, ops::MVOpGrad);
+
+REGISTER_OP_CPU_KERNEL(
+    mv, ops::MVKernel<paddle::platform::CPUDeviceContext, float>,
+    ops::MVKernel<paddle::platform::CPUDeviceContext, double>);
+REGISTER_OP_CPU_KERNEL(
+    mv_grad, ops::MVGradKernel<paddle::platform::CPUDeviceContext, float>,
+    ops::MVGradKernel<paddle::platform::CPUDeviceContext, double>);

paddle/fluid/operators/mv_op.cu

+/* 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 <typename T>
+__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 <typename T>
+class MVGradKernel<platform::CUDADeviceContext, T>
+    : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext &context) const override {
+    auto *x = context.Input<framework::Tensor>("X");
+    auto *vec = context.Input<framework::Tensor>("Vec");
+    auto *dout =
+        context.Input<framework::Tensor>(framework::GradVarName("Out"));
+    auto *dx = context.Output<framework::Tensor>(framework::GradVarName("X"));
+    auto *dvec =
+        context.Output<framework::Tensor>(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<T>();
+    const T *vec_data = vec->data<T>();
+    const T *dout_data = dout->data<T>();
+
+    T *dx_data = dx->mutable_data<T>(context.GetPlace());
+    T *dvec_data = dvec->mutable_data<T>(context.GetPlace());
+
+    auto &dev_ctx =
+        context.template device_context<platform::CUDADeviceContext>();
+    auto blas = math::GetBlas<platform::CUDADeviceContext, T>(dev_ctx);
+
+    // calculate dx
+    auto stream = context.cuda_device_context().stream();
+    auto config = GetGpuLaunchConfig1D(dev_ctx, m * n);
+    MVGradCUDAKernel<
+        T><<<config.block_per_grid.x, config.thread_per_block.x, 0, stream>>>(
+        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<T>(1), x_data, dout_data,
+              static_cast<T>(0), dvec_data);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+namespace plat = paddle::platform;
+
+REGISTER_OP_CUDA_KERNEL(
+    mv, ops::MVKernel<paddle::platform::CUDADeviceContext, float>,
+    ops::MVKernel<paddle::platform::CUDADeviceContext, double>);
+REGISTER_OP_CUDA_KERNEL(
+    mv_grad, ops::MVGradKernel<paddle::platform::CUDADeviceContext, float>,
+    ops::MVGradKernel<paddle::platform::CUDADeviceContext, double>);

paddle/fluid/operators/mv_op.h

+/* 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 <algorithm>
+#include <utility>
+#include <vector>
+#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 <typename DeviceContext, typename T>
+class MVKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext &context) const override {
+    auto *x = context.Input<framework::Tensor>("X");
+    auto *vec = context.Input<framework::Tensor>("Vec");
+
+    auto *out = context.Output<framework::Tensor>("Out");
+
+    auto dim_x = x->dims();
+
+    // get data ptr
+    const T *x_data = x->data<T>();
+    const T *vec_data = vec->data<T>();
+    T *out_data = out->mutable_data<T>(context.GetPlace());
+
+    auto &dev_ctx = context.template device_context<DeviceContext>();
+    auto blas = math::GetBlas<DeviceContext, T>(dev_ctx);
+
+    blas.GEMV(false, dim_x[0], dim_x[1], static_cast<T>(1), x_data, vec_data,
+              static_cast<T>(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 <typename DeviceContext, typename T>
+class MVGradKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext &context) const override {
+    auto *x = context.Input<framework::Tensor>("X");
+    auto *vec = context.Input<framework::Tensor>("Vec");
+    auto *dout =
+        context.Input<framework::Tensor>(framework::GradVarName("Out"));
+    auto *dx = context.Output<framework::Tensor>(framework::GradVarName("X"));
+    auto *dvec =
+        context.Output<framework::Tensor>(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<T>();
+    const T *vec_data = vec->data<T>();
+    const T *dout_data = dout->data<T>();
+
+    T *dx_data = dx->mutable_data<T>(context.GetPlace());
+    T *dvec_data = dvec->mutable_data<T>(context.GetPlace());
+
+    auto &dev_ctx = context.template device_context<DeviceContext>();
+    auto blas = math::GetBlas<DeviceContext, T>(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<T>(1), x_data, dout_data,
+              static_cast<T>(0), dvec_data);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

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

python/paddle/fluid/tests/unittests/test_mv_op.py

+#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()

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

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