nir/lower_system_values change num_workgroups to uint32_t

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24882>

diff --git a/src/compiler/nir/nir.h b/src/compiler/nir/nir.h
index 9850d33..d9ec891 100644 (file)
--- a/src/compiler/nir/nir.h
+++ b/src/compiler/nir/nir.h
@@ -5375,7 +5375,7 @@ typedef struct nir_lower_compute_system_values_options {
     * and compute it quickly. Fall back to slow computation if not.
     */
    bool shortcut_1d_workgroup_id : 1;
-   uint16_t num_workgroups[3]; /* Compile-time-known dispatch sizes, or 0 if unknown. */
+   uint32_t num_workgroups[3]; /* Compile-time-known dispatch sizes, or 0 if unknown. */
 } nir_lower_compute_system_values_options;
 
 bool nir_lower_compute_system_values(nir_shader *shader,

diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c
index 9ceb9d7..9412a09 100644 (file)
--- a/src/compiler/nir/nir_lower_system_values.c
+++ b/src/compiler/nir/nir_lower_system_values.c
@@ -384,7 +384,7 @@ id_to_index_no_umod_slow(nir_builder *b, nir_def *index,
 static nir_def *
 lower_id_to_index_no_umod(nir_builder *b, nir_def *index,
                           nir_def *size, unsigned bit_size,
-                          const uint16_t *size_imm,
+                          const uint32_t *size_imm,
                           bool shortcut_1d)
 {
    nir_def *size_x, *size_y;
@@ -465,7 +465,7 @@ lower_compute_system_value_filter(const nir_instr *instr, const void *_state)
 }
 
 static nir_def *
-try_lower_id_to_index_1d(nir_builder *b, nir_def *index, const uint16_t *size)
+try_lower_id_to_index_1d(nir_builder *b, nir_def *index, const uint32_t *size)
 {
    /* size_x = 1, size_y = 1, therefore Z = local index */
    if (size[0] == 1 && size[1] == 1)
@@ -512,8 +512,10 @@ lower_compute_system_value_instr(nir_builder *b,
              * this way we don't leave behind extra ALU instrs.
              */
 
-            nir_def *val = try_lower_id_to_index_1d(b, local_index,
-                                                    b->shader->info.workgroup_size);
+            uint32_t wg_size[3] = {b->shader->info.workgroup_size[0],
+                                   b->shader->info.workgroup_size[1],
+                                   b->shader->info.workgroup_size[2]};
+            nir_def *val = try_lower_id_to_index_1d(b, local_index, wg_size);
             if (val)
                return val;
          }
@@ -734,7 +736,7 @@ lower_compute_system_value_instr(nir_builder *b,
       if (!options)
          return NULL;
 
-      const uint16_t *num_wgs_imm = options->num_workgroups;
+      const uint32_t *num_wgs_imm = options->num_workgroups;
 
       /* Exit early when none of the num workgroups components are known at
        * compile time.