From c2c0e13b78c0b5cf4a75da3fb82b06966f39ed57 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Tue, 29 Jul 2025 00:01:04 +0200 Subject: [PATCH] nak: optimize load_subgroup_id If a dimension in workgroup_size is 1 we can simply emit a 0 instead of loading a system value. Helps a couple of compute shaders. Totals: CodeSize: 934788544 -> 934788048 (-0.00%) Static cycle count: 222856984 -> 222856908 (-0.00%) Totals from 8 (0.01% of 87622) affected shaders: CodeSize: 34992 -> 34496 (-1.42%) Static cycle count: 8593 -> 8517 (-0.88%) Reviewed-by: Faith Ekstrand Part-of: --- src/nouveau/compiler/nak_nir.c | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/src/nouveau/compiler/nak_nir.c b/src/nouveau/compiler/nak_nir.c index 1a3ae4d5c70..9f59338a40a 100644 --- a/src/nouveau/compiler/nak_nir.c +++ b/src/nouveau/compiler/nak_nir.c @@ -582,12 +582,16 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, val = nir_imm_int(b, 0); } else { assert(!b->shader->info.workgroup_size_variable); - nir_def *tid_x = nak_nir_load_sysval(b, NAK_SV_TID_X, - ACCESS_CAN_REORDER); - nir_def *tid_y = nak_nir_load_sysval(b, NAK_SV_TID_Y, - ACCESS_CAN_REORDER); - nir_def *tid_z = nak_nir_load_sysval(b, NAK_SV_TID_Z, - ACCESS_CAN_REORDER); + + nir_def *tid_x = b->shader->info.workgroup_size[0] > 1 + ? nak_nir_load_sysval(b, NAK_SV_TID_X, ACCESS_CAN_REORDER) + : nir_imm_int(b, 0); + nir_def *tid_y = b->shader->info.workgroup_size[1] > 1 + ? nak_nir_load_sysval(b, NAK_SV_TID_Y, ACCESS_CAN_REORDER) + : nir_imm_int(b, 0); + nir_def *tid_z = b->shader->info.workgroup_size[2] > 1 + ? nak_nir_load_sysval(b, NAK_SV_TID_Z, ACCESS_CAN_REORDER) + : nir_imm_int(b, 0); const uint16_t *wg_size = b->shader->info.workgroup_size; nir_def *tid =