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

Add checks for unknown workgroup count when inferring boundaries. #9258

Merged
merged 2 commits into from
Jun 2, 2022

Conversation

hanhanW
Copy link
Contributor

@hanhanW hanhanW commented Jun 1, 2022

The workgroup count is set to zero when we're not able to infer the bounds.

Fixes #9244

Copy link
Contributor

@dcaballe dcaballe left a comment

Choose a reason for hiding this comment

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

Thanks for looking into this!

Copy link
Contributor

@MaheshRavishankar MaheshRavishankar left a comment

Choose a reason for hiding this comment

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

Awesome! Thanks!

Copy link
Contributor

@MaheshRavishankar MaheshRavishankar left a comment

Choose a reason for hiding this comment

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

I actually have a question that is maybe unrelated, but might help fix this better. It might be worht looking at the call sites of these methods. The workgroup_count setting to 0 itself is wrong as default.

@hanhanW
Copy link
Contributor Author

hanhanW commented Jun 2, 2022

I actually have a question that is maybe unrelated, but might help fix this better. It might be worht looking at the call sites of these methods. The workgroup_count setting to 0 itself is wrong as default.

I had the same idea in the beginning, and I decided to check zeros.

Below is the implementation of getting workgroup_count. We're able to distribute some loops, but the bound (or say workload_size) can't be inferred. It does not only happen in peeling case, but also in dynamic shapes. We'll need a thing about unknown. Setting them to zero means that the workgroup count is unknown to me.

Ideally, the coding style would suggest to use Optional. It let the type tell us that they are unknown. However, it increases memory usage and code complexity. Using zero in this case sounds reasonable to me. zero is an invalid setting and could be an unknown option here.

https://github.com/google/iree/blob/2187016fa6866437502987ce8eda17988b5f1134/compiler/src/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp#L94-L134

@MaheshRavishankar
Copy link
Contributor

I actually have a question that is maybe unrelated, but might help fix this better. It might be worht looking at the call sites of these methods. The workgroup_count setting to 0 itself is wrong as default.

I had the same idea in the beginning, and I decided to check zeros.

Below is the implementation of getting workgroup_count. We're able to distribute some loops, but the bound (or say workload_size) can't be inferred. It does not only happen in peeling case, but also in dynamic shapes. We'll need a thing about unknown. Setting them to zero means that the workgroup count is unknown to me.

Ideally, the coding style would suggest to use Optional. It let the type tell us that they are unknown. However, it increases memory usage and code complexity. Using zero in this case sounds reasonable to me. zero is an invalid setting and could be an unknown option here.

https://github.com/google/iree/blob/2187016fa6866437502987ce8eda17988b5f1134/compiler/src/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp#L94-L134

(Its now a priority for me to kill this pass! Ill look into how to do this. This code is all just tech debt now).

Its not clear to me how the workgroup_size ends up as zero if it isnt distributed. This all seems like a footgun. I'm fine with landing this, but I think for any exploration this is causing issues with, we should just drop this pass. (FYI @dcaballe )

@hanhanW hanhanW merged commit 697396e into iree-org:main Jun 2, 2022
@hanhanW hanhanW deleted the fix-9244 branch June 2, 2022 04:59
@dcaballe
Copy link
Contributor

dcaballe commented Jun 2, 2022

Its not clear to me how the workgroup_size ends up as zero if it isnt distributed. This all seems like a footgun. I'm fine with landing this, but I think for any exploration this is causing issues with, we should just drop this pass. (FYI @dcaballe )

I tried dropping the pass but that led to other issues. Perhaps there are some dependencies with this pass down the road? Let me see if this fixes the issue.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

RemoveSingleIterationLoop crashes trying to simplify a mod 0 affine expression
4 participants