nir: make workgroup_id 32 bit only

No backend supports 64 bit values natively anyway.

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 f055ae0..3877b3e 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
@@ -26,7 +26,7 @@ static nir_def *
 task_workgroup_index(nir_builder *b,
                      lower_tsms_io_state *s)
 {
-   nir_def *id = nir_load_workgroup_id(b, 32);
+   nir_def *id = nir_load_workgroup_id(b);
 
    nir_def *x = nir_channel(b, id, 0);
    nir_def *y = nir_channel(b, id, 1);

diff --git a/src/amd/vulkan/meta/radv_meta.c b/src/amd/vulkan/meta/radv_meta.c
index eb8aa3d..35ae14f 100644 (file)
--- a/src/amd/vulkan/meta/radv_meta.c
+++ b/src/amd/vulkan/meta/radv_meta.c
@@ -680,7 +680,7 @@ get_global_ids(nir_builder *b, unsigned num_components)
    unsigned mask = BITFIELD_MASK(num_components);
 
    nir_def *local_ids = nir_channels(b, nir_load_local_invocation_id(b), mask);
-   nir_def *block_ids = nir_channels(b, nir_load_workgroup_id(b, 32), mask);
+   nir_def *block_ids = nir_channels(b, nir_load_workgroup_id(b), mask);
    nir_def *block_size =
       nir_channels(b,
                    nir_imm_ivec4(b, b->shader->info.workgroup_size[0], b->shader->info.workgroup_size[1],

diff --git a/src/amd/vulkan/meta/radv_meta_buffer.c b/src/amd/vulkan/meta/radv_meta_buffer.c
index 7415618..4f95749 100644 (file)
--- a/src/amd/vulkan/meta/radv_meta_buffer.c
+++ b/src/amd/vulkan/meta/radv_meta_buffer.c
@@ -16,7 +16,7 @@ build_buffer_fill_shader(struct radv_device *dev)
    nir_def *data = nir_swizzle(&b, nir_channel(&b, pconst, 3), (unsigned[]){0, 0, 0, 0}, 4);
 
    nir_def *global_id = nir_iadd(
-      &b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b, 32), 0), b.shader->info.workgroup_size[0]),
+      &b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
       nir_load_local_invocation_index(&b));
 
    nir_def *offset = nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset);
@@ -38,7 +38,7 @@ build_buffer_copy_shader(struct radv_device *dev)
    nir_def *dst_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b1100));
 
    nir_def *global_id = nir_iadd(
-      &b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b, 32), 0), b.shader->info.workgroup_size[0]),
+      &b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
       nir_load_local_invocation_index(&b));
 
    nir_def *offset = nir_u2u64(&b, nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset));

diff --git a/src/amd/vulkan/meta/radv_meta_decompress.c b/src/amd/vulkan/meta/radv_meta_decompress.c
index c9b1a95..4fd295a 100644 (file)
--- a/src/amd/vulkan/meta/radv_meta_decompress.c
+++ b/src/amd/vulkan/meta/radv_meta_decompress.c
@@ -52,7 +52,7 @@ build_expand_depth_stencil_compute_shader(struct radv_device *dev)
    output_img->data.binding = 1;
 
    nir_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_def *wg_id = nir_load_workgroup_id(&b, 32);
+   nir_def *wg_id = nir_load_workgroup_id(&b);
    nir_def *block_size = nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1],
                                        b.shader->info.workgroup_size[2], 0);
 

diff --git a/src/amd/vulkan/meta/radv_meta_fmask_copy.c b/src/amd/vulkan/meta/radv_meta_fmask_copy.c
index 5d2f2d1..89b1b7a 100644 (file)
--- a/src/amd/vulkan/meta/radv_meta_fmask_copy.c
+++ b/src/amd/vulkan/meta/radv_meta_fmask_copy.c
@@ -43,7 +43,7 @@ build_fmask_copy_compute_shader(struct radv_device *dev, int samples)
    output_img->data.binding = 1;
 
    nir_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_def *wg_id = nir_load_workgroup_id(&b, 32);
+   nir_def *wg_id = nir_load_workgroup_id(&b);
    nir_def *block_size = nir_imm_ivec3(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1],
                                        b.shader->info.workgroup_size[2]);
 

diff --git a/src/broadcom/vulkan/v3dv_query.c b/src/broadcom/vulkan/v3dv_query.c
index 6fbae5b..58c1609 100644 (file)
--- a/src/broadcom/vulkan/v3dv_query.c
+++ b/src/broadcom/vulkan/v3dv_query.c
@@ -1382,7 +1382,7 @@ get_set_query_availability_cs()
     * ever change any of these parameters we need to update how we compute the
     * query index here.
     */
-   nir_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b, 32), 0);
+   nir_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b), 0);
 
    nir_def *offset =
       nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 0, .range = 4);
@@ -1446,7 +1446,7 @@ get_reset_occlusion_query_cs()
     * ever change any of these parameters we need to update how we compute the
     * query index here.
     */
-   nir_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b, 32), 0);
+   nir_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b), 0);
 
    nir_def *avail_offset =
       nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 0, .range = 4);
@@ -1523,7 +1523,7 @@ get_copy_query_results_cs(VkQueryResultFlags flags)
     * ever change any of these parameters we need to update how we compute the
     * query index here.
     */
-   nir_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b, 32), 0);
+   nir_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b), 0);
    nir_def *query_idx = nir_iadd(&b, base_query_idx, wg_id);
 
    /* Read query availability if needed */

diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py
index 36d3db1..b1b9eae 100644 (file)
--- a/src/compiler/nir/nir_intrinsics.py
+++ b/src/compiler/nir/nir_intrinsics.py
@@ -850,7 +850,7 @@ system_value("local_invocation_id", 3)
 system_value("local_invocation_index", 1)
 # zero_base indicates it starts from 0 for the current dispatch
 # non-zero_base indicates the base is included
-system_value("workgroup_id", 3, bit_sizes=[32, 64])
+system_value("workgroup_id", 3)
 system_value("workgroup_id_zero_base", 3)
 # The workgroup_index is intended for situations when a 3 dimensional
 # workgroup_id is not available on the HW, but a 1 dimensional index is.

diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c
index c020c24..ace1f9d 100644 (file)
--- a/src/compiler/nir/nir_lower_system_values.c
+++ b/src/compiler/nir/nir_lower_system_values.c
@@ -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_workgroup_id:
    case nir_intrinsic_load_workgroup_size:
       return sanitize_32bit_sysval(b, intrin);
 
@@ -666,10 +667,11 @@ lower_compute_system_value_instr(nir_builder *b,
       if ((options && options->has_base_workgroup_id) ||
           !b->shader->options->has_cs_global_id) {
          nir_def *group_size = nir_load_workgroup_size(b);
-         nir_def *group_id = nir_load_workgroup_id(b, bit_size);
+         nir_def *group_id = nir_load_workgroup_id(b);
          nir_def *local_id = nir_load_local_invocation_id(b);
 
-         return nir_iadd(b, nir_imul(b, group_id, nir_u2uN(b, group_size, bit_size)),
+         return nir_iadd(b, nir_imul(b, nir_u2uN(b, group_id, bit_size),
+                         nir_u2uN(b, group_size, bit_size)),
                          nir_u2uN(b, local_id, bit_size));
       } else {
          return NULL;

diff --git a/src/gallium/auxiliary/nir/tgsi_to_nir.c b/src/gallium/auxiliary/nir/tgsi_to_nir.c
index b941d37..5f6ed66 100644 (file)
--- a/src/gallium/auxiliary/nir/tgsi_to_nir.c
+++ b/src/gallium/auxiliary/nir/tgsi_to_nir.c
@@ -614,7 +614,7 @@ ttn_src_for_file_and_index(struct ttn_compile *c, unsigned file, unsigned index,
          load = nir_load_local_invocation_id(b);
          break;
       case TGSI_SEMANTIC_BLOCK_ID:
-         load = nir_load_workgroup_id(b, 32);
+         load = nir_load_workgroup_id(b);
          break;
       case TGSI_SEMANTIC_BLOCK_SIZE:
          load = nir_load_workgroup_size(b);

diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
index 974ad99..b78ed0d 100644 (file)
--- a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
+++ b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
@@ -44,7 +44,7 @@ static nir_def *get_global_ids(nir_builder *b, unsigned num_components)
    unsigned mask = BITFIELD_MASK(num_components);
 
    nir_def *local_ids = nir_channels(b, nir_load_local_invocation_id(b), mask);
-   nir_def *block_ids = nir_channels(b, nir_load_workgroup_id(b, 32), mask);
+   nir_def *block_ids = nir_channels(b, nir_load_workgroup_id(b), mask);
    nir_def *block_size = nir_channels(b, nir_load_workgroup_size(b), mask);
    return nir_iadd(b, nir_imul(b, block_ids, block_size), local_ids);
 }

diff --git a/src/mesa/state_tracker/st_pbo_compute.c b/src/mesa/state_tracker/st_pbo_compute.c
index 223ad5a..6e95b45 100644 (file)
--- a/src/mesa/state_tracker/st_pbo_compute.c
+++ b/src/mesa/state_tracker/st_pbo_compute.c
@@ -650,7 +650,7 @@ create_conversion_shader(struct st_context *st, enum pipe_texture_target target,
                                       b.shader->info.workgroup_size[1],
                                       b.shader->info.workgroup_size[2],
                                       0);
-   nir_def *wid = nir_load_workgroup_id(&b, 32);
+   nir_def *wid = nir_load_workgroup_id(&b);
    nir_def *iid = nir_load_local_invocation_id(&b);
    nir_def *tile = nir_imul(&b, wid, bsize);
    nir_def *global_id = nir_iadd(&b, tile, iid);