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

nir/lower_system_values: replace local_invocation_id components with zero

Browse Source 
fossil-db (Sienna Cichlid):
Totals from 360 (0.28% of 128647) affected shaders:
VGPRs: 7912 -> 7272 (-8.09%); split: -8.59%, +0.51%
CodeSize: 542456 -> 544688 (+0.41%); split: -0.32%, +0.73%
MaxWaves: 10866 -> 10952 (+0.79%)
Instrs: 95973 -> 96010 (+0.04%); split: -0.34%, +0.38%
Latency: 4366023 -> 4344664 (-0.49%); split: -0.90%, +0.41%
InvThroughput: 19656659 -> 18297185 (-6.92%); split: -6.92%, +0.00%
VClause: 3242 -> 3116 (-3.89%); split: -4.04%, +0.15%
SClause: 3422 -> 3504 (+2.40%); split: -0.20%, +2.60%
Copies: 8854 -> 9376 (+5.90%); split: -0.89%, +6.79%
Branches: 2329 -> 2326 (-0.13%); split: -0.39%, +0.26%
PreSGPRs: 7620 -> 7841 (+2.90%); split: -0.43%, +3.33%
PreVGPRs: 5765 -> 5504 (-4.53%)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel-schuermann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13757>
This commit is contained in:
Rhys Perry
2021-11-11 14:26:00 +00:00
committed by Marge Bot
parent a9b4a507fe
commit 719b48f85d
1 changed files with 15 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -387,6 +387,21 @@ lower_compute_system_value_instr(nir_builder *b,
return nir_vec3(b, x, y, z);
}
/* If a workgroup size dimension is 1, then the local invocation id must be zero. */
nir_component_mask_t is_zero = 0;
is_zero |= b->shader->info.workgroup_size[0] == 1 ? 0x1 : 0x0;
is_zero |= b->shader->info.workgroup_size[1] == 1 ? 0x2 : 0x0;
is_zero |= b->shader->info.workgroup_size[2] == 1 ? 0x4 : 0x0;
if (!b->shader->info.workgroup_size_variable && is_zero) {
nir_ssa_def *defs[3];
for (unsigned i = 0; i < 3; i++) {
defs[i] = is_zero & (1 << i) ? nir_imm_zero(b, 1, 32) :
nir_channel(b, &intrin->dest.ssa, i);
}
return nir_vec(b, defs, 3);
}
return NULL;
case nir_intrinsic_load_local_invocation_index: