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>

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);

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],

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));

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);
 

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]);
 

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 */

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.

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;

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);

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);
 }

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);