Mesa (main): ir3: Initialize local size earlier

[GitLab Mirror]

Module: Mesa
Branch: main
Commit: decbced48f0dfcc2d4a4a2e3a8b39be2e90f4962
URL:    http://cgit.freedesktop.org/mesa/mesa/commit/?id=decbced48f0dfcc2d4a4a2e3a8b39be2e90f4962

Author: Connor Abbott <cwabbott0 at gmail.com>
Date:   Thu Sep 30 11:53:14 2021 +0200

ir3: Initialize local size earlier

We need the local size in RA for occupancy calculations. Not
initializing these had the unfortunate consequence of
ir3_get_reg_independent_max_waves() returning 0 for compute shaders with
shared variables, disabling the register limiting logic.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13143>

---

 src/freedreno/ir3/ir3_compiler_nir.c | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/src/freedreno/ir3/ir3_compiler_nir.c b/src/freedreno/ir3/ir3_compiler_nir.c
index 18509be2148..e69b9e76c6a 100644
--- a/src/freedreno/ir3/ir3_compiler_nir.c
+++ b/src/freedreno/ir3/ir3_compiler_nir.c
@@ -3992,6 +3992,13 @@ ir3_compile_shader_nir(struct ir3_compiler *compiler,
 
    ir = so->ir = ctx->ir;
 
+   if (so->type == MESA_SHADER_COMPUTE) {
+      so->local_size[0] = ctx->s->info.workgroup_size[0];
+      so->local_size[1] = ctx->s->info.workgroup_size[1];
+      so->local_size[2] = ctx->s->info.workgroup_size[2];
+      so->local_size_variable = ctx->s->info.workgroup_size_variable;
+   }
+
    /* Vertex shaders in a tessellation or geometry pipeline treat END as a
     * NOP and has an epilogue that writes the VS outputs to local storage, to
     * be read by the HS.  Then it resets execution mask (chmask) and chains
@@ -4348,13 +4355,6 @@ ir3_compile_shader_nir(struct ir3_compiler *compiler,
        ctx->s->info.fs.needs_quad_helper_invocations)
       so->need_pixlod = true;
 
-   if (so->type == MESA_SHADER_COMPUTE) {
-      so->local_size[0] = ctx->s->info.workgroup_size[0];
-      so->local_size[1] = ctx->s->info.workgroup_size[1];
-      so->local_size[2] = ctx->s->info.workgroup_size[2];
-      so->local_size_variable = ctx->s->info.workgroup_size_variable;
-   }
-
 out:
    if (ret) {
       if (so->ir)