Add topk op #3760

typhoonzero · 2017-08-30T09:00:11Z

Still got problems when using a 3d tensor input.

… top_k_op

jacquesqiao · 2017-08-30T16:40:07Z

python/paddle/v2/framework/tests/CMakeLists.txt

@@ -16,6 +16,7 @@ py_test(test_cross_entropy_op SRCS test_cross_entropy_op.py)
 py_test(test_gather_op SRCS test_gather_op.py)
 py_test(test_scatter_op SRCS test_scatter_op.py)
 py_test(test_fill_zeros_like_op SRCS test_fill_zeros_like_op.py)
+py_test(test_top_k_op SRCS test_top_k_op.py)


do not find file test_top_k_op.py

Thanks. Added.

typhoonzero · 2017-09-01T04:54:44Z

I'm looking into http://www.math.grin.edu/%7Eblanchaj/Research/ABGS_KSelection.pdf for some details of the performance of topk kernel.

qingqing01 · 2017-09-01T06:27:52Z

paddle/operators/top_k_op.cc

+limitations under the License. */
+
+#include "paddle/operators/top_k_op.h"
+#include <iostream>


remove this line.

qingqing01 · 2017-09-01T06:31:04Z

paddle/operators/top_k_op.cc

+        R"DOC(If the input is a vector (1d tensor), finds the k largest entries in the vector and outputs their values and indices as vectors. Thus values[j] is the j-th largest entry in input, and its index is indices[j].
+
+    For matrices, computes the top k entries in each row. )DOC");
+    AddAttr<AttrType>("k",


AddAttr<int>( "k",

The template <typename AttrType> is not needed here.

qingqing01 · 2017-09-01T06:35:23Z

paddle/operators/top_k_op.cu

+  }
+
+  T v_;
+  int id_;


The member variables in structs and classes have different naming rules. There is no _

https://google.github.io/styleguide/cppguide.html#Variable_Names

qingqing01 · 2017-09-01T06:39:09Z

paddle/operators/top_k_op.cu

+    auto* input = ctx.Input<Tensor>("X");
+    auto* output = ctx.Output<Tensor>("Out");
+    auto* indices = ctx.Output<Tensor>("Indices");
+    size_t k = static_cast<AttrType>(ctx.op_.GetAttr<AttrType>("k"));


Same as above comments, the typename AttrType = int is not necessary.

qingqing01 · 2017-09-01T06:39:19Z

paddle/operators/top_k_op.cu

+  void Compute(const framework::ExecutionContext& ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "It must use GPUPlace.");
+    std::cout << "in TopkOpCUDAKernel" << std::endl;


remove this line.

qingqing01 · 2017-09-01T06:40:37Z

paddle/operators/top_k_op.cu

+    output->mutable_data<T>(ctx.GetPlace());
+    indices->mutable_data<int>(ctx.GetPlace());
+    T* output_data = output->data<T>();
+    int* indices_data = indices->data<int>();


line 296 to line 299 can be write as follows:

T* output_data = output->mutable_data<T>(ctx.GetPlace()); int* indices_data =indices->mutable_data<int>(ctx.GetPlace());

qingqing01 · 2017-09-01T06:45:22Z

paddle/operators/top_k_op.cu

+    dim3 threads(256, 1);
+    dim3 grid(input_height, 1);
+
+    KeMatrixTopK<T, 5, 256><<<grid, threads>>>(


Now the cuda stream can be got by reinterpret_cast<platform::CUDADeviceContext*>(ctx.device_context())->stream(), so

auto stream = reinterpret_cast<platform::CUDADeviceContext*>(ctx.device_context())->stream(); KeMatrixTopK<T, 5, 256><<<grid, threads, 0, stream>>>( ...

ctx.device_context() returns const type currently, can not be reinterpret_casted currently?

qingqing01 · 2017-09-01T06:59:44Z

paddle/operators/top_k_op.h

+    auto* output = ctx.Output<Tensor>("Out");
+    auto* indices = ctx.Output<Tensor>("Indices");
+    // k is determined by Attr
+    const size_t k = static_cast<AttrType>(ctx.op_.GetAttr<AttrType>("k"));


For the typename AttrType = int, same problem with above comments.

qingqing01 · 2017-09-01T07:03:46Z

paddle/operators/top_k_op.h

+
+    auto X = EigenMatrix<T>::From(*input);
+    auto Out = EigenMatrix<T>::From(*output);
+    auto Indices = EigenMatrix<T>::From(*indices);


If using the std::partial_sort and not using the function in Eigen, there is no need to convert to Eigen tensor.

qingqing01 · 2017-09-01T07:05:33Z

paddle/operators/top_k_op.h

+      // TODO(typhoonzero): make this more efficient
+      std::vector<std::pair<T, size_t>> vec;
+      for (size_t j = 0; j < col; j++) {
+        vec.push_back(std::pair<T, size_t>(X(i, j), j));


Maybe we can use the plain pointer to get data and there is no need to use Eigen.

Follow comments, all done. Please review again, Thanks very much!

… top_k_op

qingqing01

some minor problems

qingqing01 · 2017-09-06T13:24:37Z

paddle/operators/top_k_op.cc

+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input of TopkOP must be initialized.");
+    auto *input = ctx.Input<framework::Tensor>("X");
+    const int k = static_cast<int>(ctx.op().GetAttr<int>("k"));


ctx.op().GetAttr ->ctx.GetAttr.

if merged after this PR: #3903
please use Attr.

qingqing01 · 2017-09-06T13:25:29Z

paddle/operators/top_k_op.cc

+    // input must have >= 1d shape.
+    PADDLE_ENFORCE_GE(input->dims().size(), 1);
+    // input must have >= k columns.
+    PADDLE_ENFORCE_GE(input->dims()[input->dims().size() - 1], k);


need comments in PADDLE_ENFORCE_GE

qingqing01 · 2017-09-06T13:27:52Z

paddle/operators/top_k_op.cu

+    dim3 grid(input_height, 1);
+
+    // auto stream = reinterpret_cast<platform::CUDADeviceContext*>(
+    // ctx.device_context())->stream();


please refer this PR: https://github.com/PaddlePaddle/Paddle/pull/3753/files#diff-be198775c4ad57a522949cb9ff47c6b2R89

qingqing01 · 2017-09-06T13:29:39Z

paddle/operators/top_k_op.h

+        framework::slice_ddim(inputdims, 0, inputdims.size() - 1));
+    const size_t col = inputdims[inputdims.size() - 1];
+    Eigen::DSizes<int, 2> flat2dims(row, col);
+    X.reshape(flat2dims);


X -> x. all variable names should be lowercase.

qingqing01 · 2017-09-06T13:31:59Z

paddle/operators/top_k_op.h

+    const size_t row = framework::product(
+        framework::slice_ddim(inputdims, 0, inputdims.size() - 1));
+    const size_t col = inputdims[inputdims.size() - 1];
+    Eigen::DSizes<int, 2> flat2dims(row, col);


我觉得通常是： row = inputdims[0], col = product(inputdims[1: ...])

这应该是两种不同的flattern的方法，我提过一个issue：https://github.com/PaddlePaddle/Paddle/issues/3771，参考TF的实现是这样的。

qingqing01 · 2017-09-06T13:33:44Z

paddle/operators/top_k_op.h

+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    // Get the top k elements of each row of input tensor
+    // FIXME: only deal with matrix(2d tensor).


实际上我觉得当前可以只支持rank=2的，用户如果需要可以TopK之前加reshap op.

嗯嗯，同感可以先merge后续增加。

… top_k_op

qingqing01

need to add cuda stream in next PR.

typhoonzero added 4 commits August 28, 2017 21:12

init add

a1348f2

Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into…

5fb4271

… top_k_op

add topk op

5f53184

someupdate

b933d69

qingqing01 requested review from hedaoyuan, QiJune, qingqing01 and luotao1 August 30, 2017 10:20

jacquesqiao reviewed Aug 30, 2017

View reviewed changes

typhoonzero added 3 commits August 31, 2017 10:43

fix style check

95da792

add test py file

0504b5b

update top k cuda kernel

a975b2f

qingqing01 reviewed Sep 1, 2017

View reviewed changes

typhoonzero added 4 commits September 5, 2017 12:12

Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into…

861d43e

… top_k_op

Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into…

3aafa66

… top_k_op

follow comments

e3b78dc

remove debug print

cb99e4d

qingqing01 added the OpPorting label Sep 5, 2017

typhoonzero added 4 commits September 6, 2017 15:11

fix casting error

68d2c5a

fix casting error

5add8bd

fix casting error

fc53ed0

Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into…

d1b39ee

… top_k_op

qingqing01 reviewed Sep 6, 2017

View reviewed changes

typhoonzero added 4 commits September 7, 2017 13:13

fix rename bug...

d24885f

follow comments

5e342a5

Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into…

a586303

… top_k_op

fix travis

aa6876d

typhoonzero changed the title ~~[WIP]Add topk op~~ Add topk op Sep 7, 2017

qingqing01 approved these changes Sep 7, 2017

View reviewed changes

typhoonzero merged commit 3fbb692 into PaddlePaddle:develop Sep 7, 2017

typhoonzero deleted the top_k_op branch December 22, 2017 05:46

Add topk op #3760

Add topk op #3760

Uh oh!

Conversation

typhoonzero commented Aug 30, 2017

Uh oh!

jacquesqiao Aug 30, 2017 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

typhoonzero commented Sep 1, 2017

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

qingqing01 left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

qingqing01 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

jacquesqiao Aug 30, 2017 •

edited

Loading