nir: make num_workgroups 32 bit only

Signed-off-by: Karol Herbst <git@karolherbst.de>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24905>

diff --git a/src/amd/common/ac_nir_lower_taskmesh_io_to_mem.c b/src/amd/common/ac_nir_lower_taskmesh_io_to_mem.c
index 3877b3e..334e889 100644 (file)
--- a/src/amd/common/ac_nir_lower_taskmesh_io_to_mem.c
+++ b/src/amd/common/ac_nir_lower_taskmesh_io_to_mem.c
@@ -32,7 +32,7 @@ task_workgroup_index(nir_builder *b,
    nir_def *y = nir_channel(b, id, 1);
    nir_def *z = nir_channel(b, id, 2);
 
-   nir_def *grid_size = nir_load_num_workgroups(b, 32);
+   nir_def *grid_size = nir_load_num_workgroups(b);
    nir_def *grid_size_x = nir_channel(b, grid_size, 0);
    nir_def *grid_size_y = nir_channel(b, grid_size, 1);
 

diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py
index b1b9eae..0537012 100644 (file)
--- a/src/compiler/nir/nir_intrinsics.py
+++ b/src/compiler/nir/nir_intrinsics.py
@@ -857,7 +857,7 @@ system_value("workgroup_id_zero_base", 3)
 system_value("workgroup_index", 1)
 system_value("base_workgroup_id", 3, bit_sizes=[32, 64])
 system_value("user_clip_plane", 4, indices=[UCP_ID])
-system_value("num_workgroups", 3, bit_sizes=[32, 64])
+system_value("num_workgroups", 3)
 system_value("num_vertices", 1)
 system_value("helper_invocation", 1, bit_sizes=[1, 32])
 system_value("layer_id", 1)

diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c
index ace1f9d..9ceb9d7 100644 (file)
--- a/src/compiler/nir/nir_lower_system_values.c
+++ b/src/compiler/nir/nir_lower_system_values.c
@@ -54,9 +54,9 @@ static nir_def *
 build_global_group_size(nir_builder *b, unsigned bit_size)
 {
    nir_def *group_size = nir_load_workgroup_size(b);
-   nir_def *num_workgroups = nir_load_num_workgroups(b, bit_size);
+   nir_def *num_workgroups = nir_load_num_workgroups(b);
    return nir_imul(b, nir_u2uN(b, group_size, bit_size),
-                   num_workgroups);
+                   nir_u2uN(b, num_workgroups, bit_size));
 }
 
 static bool
@@ -110,6 +110,7 @@ lower_system_value_instr(nir_builder *b, nir_instr *instr, void *_state)
 
    case nir_intrinsic_load_local_invocation_id:
    case nir_intrinsic_load_local_invocation_index:
+   case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_workgroup_id:
    case nir_intrinsic_load_workgroup_size:
       return sanitize_32bit_sysval(b, intrin);
@@ -718,8 +719,9 @@ lower_compute_system_value_instr(nir_builder *b,
          if (val)
             return val;
 
+         nir_def *num_workgroups = nir_load_num_workgroups(b);
          return lower_id_to_index_no_umod(b, wg_idx,
-                                          nir_load_num_workgroups(b, bit_size),
+                                          nir_u2uN(b, num_workgroups, bit_size),
                                           bit_size,
                                           options->num_workgroups,
                                           options->shortcut_1d_workgroup_id);