Optimize layer norm forward when cols is 1024. #39167
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
… fused_dropout_op_support_ln_fp16
|
Thanks for your contribution! |
sneaxiy
reviewed
Jan 24, 2022
| @@ -19,6 +19,9 @@ limitations under the License. */ | |||
| namespace paddle { | |||
| namespace operators { | |||
|
|
|||
| #define DIVUP(x, y) (((x) + ((y)-1)) / (y)) | |||
There was a problem hiding this comment.
I prefer using function instead of macro.
There was a problem hiding this comment.
Done by using std::ceil function.
| @@ -19,6 +19,9 @@ limitations under the License. */ | |||
| namespace paddle { | |||
| namespace operators { | |||
|
|
|||
| #define DIVUP(x, y) (((x) + ((y)-1)) / (y)) | |||
| #define COLS_ 1024 | |||
There was a problem hiding this comment.
How about naming this with a more meaningful name? The name COLS_ is too simple to understand the exact meaning and easy to be conflict with the other macros.
| int ELTS_PER_ROW_PER_CTA = THREADS_PER_ROW *VecSize, | ||
| int LDGS = ELTS_PER_ROW / ELTS_PER_ROW_PER_CTA> | ||
| __global__ __launch_bounds__(THREADS_PER_CTA) void fused_ln_fwd_1024_kernel( | ||
| void *__restrict__ y_, void *__restrict__ residual_out_, |
There was a problem hiding this comment.
The void * is too hard to read. Try to just write T *, U * or anything else.
| void *__restrict__ y_, void *__restrict__ mean_out_, | ||
| void *__restrict__ var_out_, const void *__restrict__ x_, | ||
| const void *__restrict__ gamma_, const void *__restrict__ beta_, | ||
| const float epsilon, int rows, int cols) { |
There was a problem hiding this comment.
Avoid to use void *. Same above.
zkh2016
approved these changes
Jan 26, 2022
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 the performance of layer_norm forward kernel for layer_norm op and fused_dropout_reisdual_layer_norm op.
Performance results: (use nsys to collect kernel times)
(1) layer_norm op:
结论：优化后相比优化前获得2-5x加速；优化后基本打平竞品。在个别case相比竞品略差的原因是：在计算scale * (x-mean)/var + bias时，竞品均使用fp16进行计算，paddle采取转换为fp32，使用fp32进行计算的方法，后者相比前者有一些性能损耗。
(2) fused_dropout_reisudal_layer_norm op:
结论：优化后，fused_dropout_reisudal_layer_norm模块相比竞品，大约获得10%-20%左右加速比。