NeroReflex/mesa
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity

nir: make num_workgroups 32 bit only

Browse source 
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>
This commit is contained in:
Karol Herbst 2023-08-26 15:25:02 +02:00 • committed by Marge Bot
parent 1b22b67199
commit 513cd29eda
3 changed files with 7 additions and 5 deletions
Download patch file Download diff file Expand all files Collapse all files

2
src/amd/common/ac_nir_lower_taskmesh_io_to_mem.c
View file

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

2
src/compiler/nir/nir_intrinsics.py
View file

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

8
src/compiler/nir/nir_lower_system_values.c
View file

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