From: Karol Herbst Date: Sat, 26 Aug 2023 13:25:02 +0000 (+0200) Subject: nir: make num_workgroups 32 bit only X-Git-Tag: upstream/23.3.3~2708 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=513cd29eda4627a5f1be843c0ef41303bd1175ce;p=platform%2Fupstream%2Fmesa.git nir: make num_workgroups 32 bit only Signed-off-by: Karol Herbst Reviewed-by: Alyssa Rosenzweig Part-of: --- 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 --- 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 --- 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 --- 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);