[GPU] Become more honest about barrier semantics (and drop stale code) #23335
Conversation
This file contains hidden or 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
Recently, I added an upstream PR ( llvm/llvm-project#177425 ) that adds the ability for gpu barriers to declare what address spaces they fence on in addition to allowing the __syncthreads() semantics. This PR uses this support pervasively in IREE's GPU code generation, either by explicitly specifying the address spaces fenced by barrier ops we create or by inhereting it from address spaces on memrefs (which we know will b in shared memory). We were **already** using such barriers on AMD GPU targets, where we had a rewrite pattern to unconditionally rewrite `gpu.barrie' to `amdgpu.lds_barrier`, which always only fences workgroup memory and doesn't synthronize global memory. This commit removes that rewrite, which allows us to be more explicit abuht the semantics we're using **and** will allow us to synchronize global memory accesses within a workgrop if we ever need to do so. This PR also adds explanatory notes explaining the memory semantics we're expecting where needed. While I'm here, this PR removes an unmanitained and stale copy of the gpu barrier elimination logic from the IREE transform ops and delegates to the upstream version instead.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
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.
Draft note: I'm stashing my work so I can get back to it next week. I intend to amend our tests to make sure the changes were effective.
Recently, I added an upstream PR
( llvm/llvm-project#177425 ) that adds the ability for gpu barriers to declare what address spaces they fence on in addition to allowing the __syncthreads() semantics.
This PR uses this support pervasively in IREE's GPU code generation, either by explicitly specifying the address spaces fenced by barrier ops we create or by inhereting it from address spaces on memrefs (which we know will b in shared memory).
We were already using such barriers on AMD GPU targets, where we had a rewrite pattern to unconditionally rewrite
gpu.barrie' toamdgpu.lds_barrier`, which always only fences workgroup memory and doesn't synthronize global memory. This commit removes that rewrite, which allows us to be more explicit abuht the semantics we're using and will allow us to synchronize global memory accesses within a workgrop if we ever need to do so.This PR also adds explanatory notes explaining the memory semantics we're expecting where needed.
While I'm here, this PR removes an unmanitained and stale copy of the gpu barrier elimination logic from the IREE transform ops and delegates to the upstream version instead.