[pnorm] optimize p_norm for special cases #37685
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
Thanks for your contribution! |
…op for flexible call.
|
@Avin0323 你好，辛苦看下benchmark的CI。这个 PR 改动了 p_norm 的代码，优化了特殊 shape 的性能。从 CI 结果看 p_norm 的测试耗时都降低了。另外，还修改了 cmake 文件，辛苦 review 一下。 |
ZHUI
previously approved these changes
Dec 8, 2021
| HOSTDEVICE explicit inline NonzeroFunctor(int n) {} | ||
| template <typename T> | ||
| HOSTDEVICE inline T operator()(const T& x) const { | ||
| return static_cast<T>(static_cast<double>(x) != 0); |
There was a problem hiding this comment.
为什么先 static_cast<double>(x) cast 为double？
paddle/fluid/operators/p_norm_op.cu
Outdated
| auto xdim = in_x->dims(); | ||
| auto ndim = out_norm->dims(); | ||
| float porder = ctx.Attr<float>("porder"); | ||
| int axis = ctx.Attr<int>("axis"); | ||
| bool asvector = ctx.Attr<bool>("asvector"); | ||
| if (axis < 0) axis = xdim.size() + axis; | ||
| int pre, n, post; | ||
| GetDims(xdim, axis, &pre, &n, &post, asvector); | ||
| std::vector<int> reduce_axis = {axis}; | ||
|
|
||
| auto& dev_ctx = ctx.cuda_device_context(); |
There was a problem hiding this comment.
dev_ctx no usage in the function?
| } else { | ||
| Pnorm<T, block><<<grid, block, 0, dev_ctx.stream()>>>(x, pre, n, post, | ||
| porder, norm); | ||
| framework::Tensor tmp_x; |
There was a problem hiding this comment.
记一下todo，这里的 tmp_x 需要尽早去掉。运行时显存占用提升很多。
| auto negs = dx->constant(static_cast<T>(-1.)); | ||
| auto zeros = dx->constant(static_cast<T>(0.)); | ||
| auto positives = (*x) > zeros; | ||
| dx->device(place) = dy->broadcast(dim) * equals.select(ones, zeros) * |
| @@ -260,32 +254,38 @@ class PnormGradCUDAKernel : public framework::OpKernel<T> { | |||
| float porder = ctx.Attr<float>("porder"); | |||
| T eps = static_cast<T>(ctx.Attr<float>("epsilon")); | |||
| int axis = ctx.Attr<int>("axis"); | |||
| bool reduce_all = ((axis < 0) || (in_norm->numel() == 1)); | |||
paddle/fluid/operators/p_norm_op.cu
Outdated
| bool asvector = ctx.Attr<bool>("asvector"); | ||
| if (axis < 0) axis = xdim.size() + axis; | ||
| int pre, n, post; | ||
| GetDims(xdim, axis, &pre, &n, &post, asvector); | ||
| const std::vector<int> dims = {axis}; | ||
|
|
||
| auto& dev_ctx = ctx.cuda_device_context(); |
Avin0323
approved these changes
Dec 10, 2021
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
PR types
Performance optimization
PR changes
OPs
Describe
Optimize p_norm for two kinds of special cases:
(1) shape=[2, 1000, 1000], reduce axis=0
(2) shape=[1, 2000000, 1], reduce axis=1
The original version is paddlepaddle-gpu == 2.2.1. The
Timedenotes seconds per 1k steps.