fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-25 02:10:11 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

intel/compiler: optimize away local_inv_index and local_inv_id if workgroup size is 1

Browse source 
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/20292>
This commit is contained in:
Marcin Ślusarz 2022-12-02 14:47:19 +01:00 • committed by Marge Bot
parent 85b1c89e20
commit 3a60112ce5
1 changed files with 16 additions and 7 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

23
src/intel/compiler/brw_nir_lower_cs_intrinsics.c
View file

@ -204,16 +204,25 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
case nir_intrinsic_load_local_invocation_index:
case nir_intrinsic_load_local_invocation_id: {
if (nir->info.stage == MESA_SHADER_TASK ||
nir->info.stage == MESA_SHADER_MESH) {
/* Will be lowered by nir_emit_task_mesh_intrinsic() using
* information from the payload.
*/
continue;
if (!local_index && !nir->info.workgroup_size_variable) {
const uint16_t *ws = nir->info.workgroup_size;
if (ws[0] * ws[1] * ws[2] == 1) {
nir_ssa_def *zero = nir_imm_int(b, 0);
local_index = zero;
local_id = nir_vec3(b, zero, zero, zero);
}
}
/* First time we are using those, so let's calculate them. */
if (!local_index) {
if (nir->info.stage == MESA_SHADER_TASK ||
nir->info.stage == MESA_SHADER_MESH) {
/* Will be lowered by nir_emit_task_mesh_intrinsic() using
* information from the payload.
*/
continue;
}
/* First time we are using those, so let's calculate them. */
assert(!local_id);
compute_local_index_id(b, nir, &local_index, &local_id);
}