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

nir: extract try_lower_id_to_index_1d

Browse Source 
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/22334>
This commit is contained in:
Marcin Ślusarz
2023-03-31 11:51:16 +02:00
committed by Marge Bot
parent 1f8ecd3ae0
commit b5792c1a34
1 changed files with 22 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -423,6 +423,24 @@ lower_compute_system_value_filter(const nir_instr *instr, const void *_state)
return instr->type == nir_instr_type_intrinsic; return instr->type == nir_instr_type_intrinsic;
} }
static nir_ssa_def *
try_lower_id_to_index_1d(nir_builder *b, nir_ssa_def *index, const uint16_t *size)
{
/* size_x = 1, size_y = 1, therefore Z = local index */
if (size[0] == 1 && size[1] == 1)
return nir_vec3(b, nir_imm_int(b, 0), nir_imm_int(b, 0), index);
/* size_x = 1, size_z = 1, therefore Y = local index */
if (size[0] == 1 && size[2] == 1)
return nir_vec3(b, nir_imm_int(b, 0), index, nir_imm_int(b, 0));
/* size_y = 1, size_z = 1, therefore X = local index */
if (size[1] == 1 && size[2] == 1)
return nir_vec3(b, index, nir_imm_int(b, 0), nir_imm_int(b, 0));
return NULL;
}
static nir_ssa_def * static nir_ssa_def *
lower_compute_system_value_instr(nir_builder *b, lower_compute_system_value_instr(nir_builder *b,
nir_instr *instr, void *_state) nir_instr *instr, void *_state)
@@ -454,20 +472,10 @@ lower_compute_system_value_instr(nir_builder *b,
* this way we don't leave behind extra ALU instrs. * this way we don't leave behind extra ALU instrs.
*/ */
/* size_x = 1, size_y = 1, therefore Z = local index */ nir_ssa_def *val = try_lower_id_to_index_1d(b, local_index,
if (b->shader->info.workgroup_size[0] == 1 && b->shader->info.workgroup_size);
b->shader->info.workgroup_size[1] == 1) if (val)
return nir_vec3(b, nir_imm_int(b, 0), nir_imm_int(b, 0), local_index); return val;
/* size_x = 1, size_z = 1, therefore Y = local index */
if (b->shader->info.workgroup_size[0] == 1 &&
b->shader->info.workgroup_size[2] == 1)
return nir_vec3(b, nir_imm_int(b, 0), local_index, nir_imm_int(b, 0));
/* size_y = 1, size_z = 1, therefore X = local index */
if (b->shader->info.workgroup_size[1] == 1 &&
b->shader->info.workgroup_size[2] == 1)
return nir_vec3(b, local_index, nir_imm_int(b, 0), nir_imm_int(b, 0));
} }
nir_ssa_def *local_size = nir_load_workgroup_size(b); nir_ssa_def *local_size = nir_load_workgroup_size(b);