AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

nak: optimize load_subgroup_id

Browse Source 
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 <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36428>
This commit is contained in:
Karol Herbst
2025-07-29 00:01:04 +02:00
committed by Marge Bot
parent fa8731b859
commit c2c0e13b78
1 changed files with 10 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/nouveau/compiler/nak_nir.c
+10 -6
View File
@@ -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 =