[llvm] [AMDGPU][SIMemoryLegalizer] Ensure LDS -> DMA hazards respect fences (PR #220288)

[Krzysztof Drewniak via llvm-commits]

================
@@ -8140,8 +8140,20 @@ in table :ref:`amdgpu-amdhsa-memory-model-code-sequences-gfx6-gfx9-table`.
                                                              is being released.
 
                                                          2. buffer/global/flat_atomic
-     fence        release      - singlethread *none*     *none*
+     fence        release      - singlethread *none*     1. s_waitcnt lgkmcnt(0)
----------------
krzysz00 wrote:

I do think the demotion's safe in cases that rhyme with
```
[reqd_workgroup_size=32, 1, 1]
...
store LDS
fence workgroup release
barrier ; also deletable
fence workgroup acquire
load LDS
store LDS
fence workgroup relesae
barrier
fence workgroup acquire
DMA LDS => global
asyncmark
asyncwait
...
```
can become
```
[reqd_workgroup_size = 32, 1, 1]
store LDS
load LDS
store LDS
fence workgroup release
fence wornkgoup acquire
DMA LDS => global
asyncmark
asyncwait
```

because *some* of these fences are redundant

https://github.com/llvm/llvm-project/pull/220288