From 1b22b671998255686b0f3271987997ac9b1cbbd6 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sat, 26 Aug 2023 15:24:24 +0200 Subject: [PATCH] nir: make workgroup_id 32 bit only No backend supports 64 bit values natively anyway. Signed-off-by: Karol Herbst Reviewed-by: Alyssa Rosenzweig Part-of: --- src/amd/common/ac_nir_lower_taskmesh_io_to_mem.c | 2 +- src/amd/vulkan/meta/radv_meta.c | 2 +- src/amd/vulkan/meta/radv_meta_buffer.c | 4 ++-- src/amd/vulkan/meta/radv_meta_decompress.c | 2 +- src/amd/vulkan/meta/radv_meta_fmask_copy.c | 2 +- src/broadcom/vulkan/v3dv_query.c | 6 +++--- src/compiler/nir/nir_intrinsics.py | 2 +- src/compiler/nir/nir_lower_system_values.c | 6 ++++-- src/gallium/auxiliary/nir/tgsi_to_nir.c | 2 +- src/gallium/drivers/radeonsi/si_shaderlib_nir.c | 2 +- src/mesa/state_tracker/st_pbo_compute.c | 2 +- 11 files changed, 17 insertions(+), 15 deletions(-) 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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); -- 2.7.4