From 513cd29eda4627a5f1be843c0ef41303bd1175ce Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sat, 26 Aug 2023 15:25:02 +0200 Subject: [PATCH] nir: make num_workgroups 32 bit only Signed-off-by: Karol Herbst Reviewed-by: Alyssa Rosenzweig Part-of: --- src/amd/common/ac_nir_lower_taskmesh_io_to_mem.c | 2 +- src/compiler/nir/nir_intrinsics.py | 2 +- src/compiler/nir/nir_lower_system_values.c | 8 +++++--- 3 files changed, 7 insertions(+), 5 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 3877b3e984c..334e889a99f 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 @@ -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 b1b9eaed6ec..05370121c92 100644 --- 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 ace1f9d4c35..9ceb9d72f71 100644 --- 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);