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.

Sign up for GitHub

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

Jump to bottom

Add -fno-hip-uniform-block flag for the compilation when necessary #1986

Closed
umangyadav opened this issue Jul 20, 2023 · 1 comment · Fixed by #2167
Closed

Add -fno-hip-uniform-block flag for the compilation when necessary #1986

umangyadav opened this issue Jul 20, 2023 · 1 comment · Fixed by #2167
Assignees
@umangyadav

Comments

@umangyadav
Copy link
Member

HIP requires uniform WG sizes otherwise HIP runtime would through errors or it may result in undefined behaviors.

Using uniform WG sizes always however may result in performance degradation.

HIP is going to add -fno-hip-uniform-block flag to allow users to use non-uniform WG sizes.

https://reviews.llvm.org/D155213

it is supposed to become available in either rocm-5.7 or rocm-6.0. When it does, make use of it.

Related issue :
#1976
SWDEV-406592
@umangyadav umangyadav self-assigned this Jul 20, 2023
@umangyadav
Copy link
Member Author

Related issues :
ROCm/MIOpen#2307

https://github.com/ROCmSoftwarePlatform/composable_kernel/pull/838/files

@umangyadav umangyadav mentioned this issue Sep 11, 2023
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@umangyadav umangyadav

Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging a pull request may close this issue.

Add flag to accept non-uniform WG sizes ROCm/AMDMIGraphX
1 participant
@umangyadav