OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

nir/lower_system_values: Fix cs_local_index_to_id with variable workgroups

Browse Source 
In that case we need to use the sysval. That sysval can be optimized anyway in
the nonvariable case. Fixes test_basic.get_linear_ids on panfrost.

Fixes: 998d84fca5 ("nir/lower_system_values: Support lowering more intrinsics")
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18662>
This commit is contained in:
Alyssa Rosenzweig
2022-09-18 21:37:07 -04:00
committed by Marge Bot
parent 2dfab687ec
commit f4b03ea6dc
1 changed files with 3 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -487,11 +487,9 @@ lower_compute_system_value_instr(nir_builder *b,
* gl_WorkGroupSize.x + gl_LocalInvocationID.x"
*/
nir_ssa_def *local_id = nir_load_local_invocation_id(b);
nir_ssa_def *size_x =
nir_imm_int(b, b->shader->info.workgroup_size[0]);
nir_ssa_def *size_y =
nir_imm_int(b, b->shader->info.workgroup_size[1]);
nir_ssa_def *local_size = nir_load_workgroup_size(b);
nir_ssa_def *size_x = nir_channel(b, local_size, 0);
nir_ssa_def *size_y = nir_channel(b, local_size, 1);
/* Because no hardware supports a local workgroup size greater than
* about 1K, this calculation can be done in 32-bit and can save some