Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[CK_TILE] Add Various Fusion Functions to RMSNorm #1802

Merged
merged 20 commits into from
Jan 15, 2025

Conversation

ruanjm
Copy link
Contributor

@ruanjm ruanjm commented Jan 8, 2025

Added following functions to RMSNorm. In summary, the interface and function of RMSNorm is very similar to LayerNorm now:

  • Support different strides for each input/output tensors.
  • Fuse add on input.
  • Fuse quant.
  • Customizable epilogue.
  • Tests/examples are generated by example/ck_tile/10_rmsnorm2d/generate.py.
  • Rename XScale to SmoothScale in quant epilogue. This change affects layernorm, smoothquant and moe smoothquant as well.

@ruanjm ruanjm added the feature request New feature or request label Jan 8, 2025
@ruanjm ruanjm force-pushed the amd/dev/jruan/tile_rmsnorm_fusion branch 2 times, most recently from 1714df2 to c91b9b7 Compare January 9, 2025 04:21
@ruanjm ruanjm requested a review from a team as a code owner January 10, 2025 06:15
@ruanjm ruanjm force-pushed the amd/dev/jruan/tile_rmsnorm_fusion branch from 03a4be5 to b14a991 Compare January 10, 2025 06:21
@ruanjm ruanjm requested a review from rocking5566 January 10, 2025 06:22
Copy link
Contributor

@spolifroni-amd spolifroni-amd left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nothing to review for docs.

@ruanjm ruanjm force-pushed the amd/dev/jruan/tile_rmsnorm_fusion branch from a376459 to 180e739 Compare January 13, 2025 15:02
@valarLip
Copy link
Contributor

passed in ater unit test

@carlushuang carlushuang merged commit 04dd314 into develop Jan 15, 2025
10 checks passed
@carlushuang carlushuang deleted the amd/dev/jruan/tile_rmsnorm_fusion branch January 15, 2025 02:23
kylasa pushed a commit to kylasa/composable_kernel that referenced this pull request Jan 19, 2025
* Add shortcut to RMSNorm

* Modify test for adding shortcut for RMSNorm

* Add fused parameter into tests

* 1. Add YDataType. 2. rmsnorm2d_fwd_traits_ from rmsnorm2d_fwd.hpp to rmsnorm2d_fwd_api.cpp and rmsnorm2d_fwd_instance_common.hpp

* 1. Supports various stride and percisions.

* Add support of Epilogue

* Add fuse and epilogue support to rmsnorm ref

* Modify rmsnorm example

* Refactor tests/examples

* Bug fix for newly added tests/examples

* Bug fix for new tests 2

* Modify smoke test scripts

remove dbg code

* Supports non-smooth dyanmic quant

* Update Rmsnorm2dFwd::GetName()

* rename xscale and prec_sx to smoothscale and prec_sm

Bug fix after rename

Remove files

* change example_rmsnorm2d_fwd.cpp

* update performance calculator

* Fix issue in two-pass when fuse add is enabled

* Remove comment of beta

---------

Co-authored-by: rocking <[email protected]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
feature request New feature or request
Projects
None yet
Development

Successfully merging this pull request may close these issues.

5 participants