From f4b03ea6dc0ac0f02f780ba809c3ea551a72f947 Mon Sep 17 00:00:00 2001 From: Alyssa Rosenzweig Date: Sun, 18 Sep 2022 21:37:07 -0400 Subject: [PATCH] nir/lower_system_values: Fix cs_local_index_to_id with variable workgroups MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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: 998d84fca56 ("nir/lower_system_values: Support lowering more intrinsics") Signed-off-by: Alyssa Rosenzweig Reviewed-by: Timur Kristóf Reviewed-by: Jason Ekstrand Reviewed-by: Karol Herbst Part-of: --- src/compiler/nir/nir_lower_system_values.c | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c index aed2f73e25d..bd7a747c6f7 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/src/compiler/nir/nir_lower_system_values.c @@ -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