commit 3d9eb17d5d037d09479f35b8cd919ea158634d48
author Rhys Perry <pendingchaos02@gmail.com> Wed May 13 16:05:46 2020 +0100
committer Marge Bot <eric+marge@anholt.net> Tue Jul 28 16:56:34 2020 +0000
tree 20a96315d592e37bcc476e39ab7dc043cb04f7a5
parent 3af2b9e3de2ba9878c03b192d3b1574172d1923d

aco: improve workgroup-scope and lower vmem/smem barriers

No fossil-db changes on Navi.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4905>

src/amd/compiler/aco_insert_waitcnt.cpp

diff --git a/src/amd/compiler/aco_insert_waitcnt.cpp b/src/amd/compiler/aco_insert_waitcnt.cpp
index fa2b523..466158f 100644
--- a/src/amd/compiler/aco_insert_waitcnt.cpp
+++ b/src/amd/compiler/aco_insert_waitcnt.cpp
@@ -448,7 +448,7 @@
 {
    wait_imm imm;
    sync_scope subgroup_scope = ctx.program->workgroup_size <= ctx.program->wave_size ? scope_workgroup : scope_subgroup;
-   if (sync.semantics & semantics) {
+   if ((sync.semantics & semantics) && sync.scope > subgroup_scope) {
       unsigned storage = sync.storage;
       while (storage) {
          unsigned idx = u_bit_scan(&storage);
@@ -460,6 +460,10 @@
          if (bar_scope_lds <= subgroup_scope)
             events &= ~event_lds;
 
+         /* in non-WGP, the L1/L0 cache keeps all memory operations in-order for the same workgroup */
+         if (ctx.chip_class < GFX10 && sync.scope <= scope_workgroup)
+            events &= ~(event_vmem | event_vmem_store | event_smem);
+
          if (events)
             imm.combine(ctx.barrier_imm[idx]);
       }