Add clip op #3937

wanghaoshuang · 2017-09-06T22:50:57Z

… clip_op

reyoung · 2017-09-07T02:43:00Z

paddle/operators/clip_op.cu

+template <typename T>
+__global__ void ClipGradientKernel(const int N, const T min, const T max,
+                                   const T* Y, const T* dY, T* dX) {
+  CUDA_1D_KERNEL_LOOP(i, N) { dX[i] = dY[i] * (Y[i] > min && Y[i] < max); }


This implementation depends on true is 1 in C++ and CUDA. It is not good.

reyoung

Well done. It seems that Relu is an alias of Clipping Operator with min=0.0, max=+inf. To alias an operator, please reference identity_op

qingqing01 · 2017-09-14T13:13:24Z

paddle/operators/clip_op.cc

+    AddInput("X", "The input of clip op");
+    AddOutput("Out", "The output of clip op");
+    AddComment(R"DOC(
+Clip Operator.


An better comments example:

https://github.com/PaddlePaddle/Paddle/pull/3927/files#diff-3f391a144c02ad510b3e6fbb17413cf7R105

qingqing01 · 2017-09-14T13:14:52Z

paddle/operators/clip_op.cc

+Clip Operator.
+)DOC");
+    AddAttr<float>("min", "min value to be clipped.");
+    AddAttr<float>("max", "max value to be clipped.");


AddAttr放在DOC前面吧，min、max也需要类型模板，参考： https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/scale_op.cc#L36

一个op的多个attrs的类型必须一样么？

qingqing01 · 2017-09-14T13:15:32Z

paddle/operators/clip_op.cc

+  void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"), "Input(X) should not be null");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar(framework::GradVarName("Out")),
+                            "Input(Out@GRAD) should not be null");


The check is also needed in ClipOp.

qingqing01 · 2017-09-14T13:15:49Z

paddle/operators/clip_op.cc

+    auto max = Attr<float>("max");
+    auto min = Attr<float>("min");
+    PADDLE_ENFORCE_LT(min, max, "max should be greater than min.");
+    ctx.Output<Tensor>("Out")->Resize(x_dims);


Output< framework::LoDTensor>

qingqing01 · 2017-09-14T13:16:02Z

paddle/operators/clip_op.cc

+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar(framework::GradVarName("Out")),
+                            "Input(Out@GRAD) should not be null");
+    auto x_dims = ctx.Input<Tensor>("X")->dims();
+    auto *x_grad = ctx.Output<Tensor>(framework::GradVarName("X"));


Output< framework::LoDTensor>

qingqing01 · 2017-09-14T13:17:34Z

paddle/operators/clip_op.cu

+ public:
+  void Compute(const framework::ExecutionContext& context) const override {
+    auto max = context.op().Attr<float>("max");
+    auto min = context.op().Attr<float>("min");


context.op().Attr<> -> context.Attr<>

qingqing01 · 2017-09-14T13:18:33Z

paddle/operators/clip_op.cu

+    size_t count = 1;
+    for (int i = 0; i < dims.size(); ++i) {
+      count *= dims[i];
+    }


int64_t count = dims->numel()

qingqing01 · 2017-09-14T13:19:59Z

paddle/operators/clip_op.cu

+    int block = 512;
+    int grid = (N * D + block - 1) / block;
+
+    ClipGradientKernel<T><<<grid, block>>>(count, min, max, x_data, d_out_data,


Need to use CUDA stream when launching kernel. Please refer to im2col.cu in paddle/operator/math.

qingqing01 · 2017-09-14T13:21:27Z

paddle/operators/clip_op.h

+    size_t count = 1;
+    for (int i = 0; i < dims.size(); ++i) {
+      count *= dims[i];
+    }


same as above.

qingqing01 · 2017-09-14T13:22:39Z

python/paddle/v2/framework/tests/test_clip_op.py

+import numpy as np
+from paddle.v2.framework.op import Operator
+from gradient_checker import GradientChecker
+from op_test_util import OpTestMeta


Use the new unit testing framework.

2. Fix unitest. 3. Fix comments and some issues. Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into clip_op

… clip_op

JiayiFeng · 2017-09-19T03:55:06Z

paddle/operators/clip_op.cc

+    auto x_dims = ctx.Input<LoDTensor>("X")->dims();
+    auto *x_grad = ctx.Output<LoDTensor>(framework::GradVarName("X"));
+
+    x_grad->Resize(x_dims);


Check whether x_grad is nullptr before it is resized.

JiayiFeng · 2017-09-19T04:03:07Z

paddle/operators/clip_op.cu

+    if (Y[i] > min && Y[i] < max) {
+      dX[i] = dY[i];
+    } else {
+      dX[i] = 0;


clip to 0 or min / max ?

It's the gradient computation. Please refer ClipGradientOp of caffe2

JiayiFeng · 2017-09-19T04:06:59Z

paddle/operators/clip_op.cu

+    auto* x = context.Input<LoDTensor>("X");
+    auto dims = d_x->dims();
+    int64_t count = d_out->numel();
+    auto d_x_data = d_x->mutable_data<T>(context.GetPlace());


Why not just d_x->mutable_data<T>(context.GetPlace()) and keep using d_x?

d_x_data is needed by ClipGradientKernel

I mean d_x_data and d_x they are totally same. They point to the same memory.

… clip_op

JiayiFeng · 2017-09-19T20:25:41Z

paddle/operators/clip_op.cu

+      int D = d_x->dims()[1];
+      int block = 512;
+      int grid = (N * D + block - 1) / block;
+      ClipGradientKernel<T><<<


This implement is too complex. @reyoung provids a global function transform to apply a functor to every elements in an iterable object.
See https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/platform/transform.h
It works on both CPU and GPU.

Here is a demo of how to use it: https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/prelu_op.h#L57

JiayiFeng · 2017-09-19T20:37:01Z

paddle/operators/clip_op.h

+  auto x_tensor = EigenTensor<T, D>::From(*x);
+  auto out_tensor = EigenTensor<T, D>::From(*out);
+  auto place = context.GetEigenDevice<Place>();
+  out_tensor.device(place) = x_tensor.cwiseMin(max).cwiseMax(min);


transform() can also be used here. With it, there is no necessity to use Eigen and we can get rid of template parameter size_t D.

… clip_op

JiayiFeng · 2017-09-21T01:33:56Z

paddle/operators/clip_op.h

+    if (y > min_ && y < max_)
+      return x;
+    else
+      return 0;


return y > min_ && y < max_ ? x : 0;

JiayiFeng

LGTM

JiayiFeng · 2017-09-21T01:43:06Z

paddle/operators/clip_op.h

+    auto* out = context.Output<Tensor>("Out");
+    T* out_data = out->mutable_data<T>(context.GetPlace());
+    const T* x_data = x->data<T>();
+    int numel = x->numel();


… clip_op

JiayiFeng · 2017-09-21T02:14:34Z

paddle/operators/clip_op.h

+    if (d_x != nullptr) {
+      auto* x = context.Input<Tensor>("X");
+      int64_t numel = d_out->numel();
+      auto d_x_data = d_x->mutable_data<T>(context.GetPlace());


d_x_data is a pointer, so please use auto*

wanghaoshuang added 3 commits September 7, 2017 06:49

Add clip op

987cdf1

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

320df7a

… clip_op

Fix issues

6e964ad

qingqing01 added the OpPorting label Sep 7, 2017

reyoung reviewed Sep 7, 2017

View reviewed changes

qingqing01 requested a review from Superjomn September 9, 2017 02:22

qingqing01 reviewed Sep 14, 2017

View reviewed changes

wanghaoshuang added 4 commits September 18, 2017 10:48

fix issues

2321a37

1. Add CUDA stream when launching kernel.

a345b71

2. Fix unitest. 3. Fix comments and some issues. Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into clip_op

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

3102a52

… clip_op

remove gradient_checker.py

44224f4

JiayiFeng reviewed Sep 19, 2017

View reviewed changes

wanghaoshuang added 3 commits September 19, 2017 23:27

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

a3c3b78

… clip_op

Remove const cast for device context

14fb15b

Add nullptr check

743dfd8

JiayiFeng reviewed Sep 19, 2017

View reviewed changes

wanghaoshuang added 3 commits September 20, 2017 13:46

Use Transform instead of eigen

2cde56c

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

3f3848c

… clip_op

Update transform invocation

1fdad1a

JiayiFeng reviewed Sep 21, 2017

View reviewed changes

JiayiFeng previously approved these changes Sep 21, 2017

View reviewed changes

JiayiFeng reviewed Sep 21, 2017

View reviewed changes

wanghaoshuang added 2 commits September 21, 2017 09:51

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

c7b6d2c

… clip_op

Fix some inssues

1244050

wanghaoshuang dismissed JiayiFeng’s stale review via 1244050 September 21, 2017 01:56

JiayiFeng reviewed Sep 21, 2017

View reviewed changes

Fix ptr type

9569255

JiayiFeng approved these changes Sep 21, 2017

View reviewed changes

wanghaoshuang merged commit e53dc8a into PaddlePaddle:develop Sep 21, 2017

wanghaoshuang deleted the clip_op branch May 20, 2022 03:56

Add clip op #3937

Add clip op #3937

Conversation

wanghaoshuang commented Sep 6, 2017

Choose a reason for hiding this comment

reyoung left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

JiayiFeng left a comment

Choose a reason for hiding this comment

JiayiFeng Sep 21, 2017 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

JiayiFeng Sep 21, 2017 •

edited

Loading