commit db9d11e4c000f693e9a70f7d3100db0c5294db9e
author Krzysztof Drewniak <Krzysztof.Drewniak@amd.com> Fri Feb 06 15:53:04 2026 -0800
committer GitHub <noreply@github.com> Fri Feb 06 15:53:04 2026 -0800
tree 6651dbb86edc8617ef4fc1aa18b1360f63c11cd9
parent 16a2395f4487b32a98746040b65d7ed5a025bc37

[GPU] Become more honest about barrier semantics (and drop stale code) (#23335)

Recently, I added an upstream PR
( https://github.com/llvm/llvm-project/pull/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.