nir/opt_if: fix fighting between split_alu_of_phi and peel_initial_break

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Fixes: 6b4b044739 ("nir/opt_loop: add loop peeling optimization")
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/11822
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/31002>
(cherry picked from commit 4f44a944bb)
This commit is contained in:
Rhys Perry 2024-09-02 11:58:22 +01:00 committed by Eric Engestrom
parent be004d4797
commit 618c4c127d
5 changed files with 30 additions and 28 deletions

View file

@ -1054,7 +1054,7 @@
"description": "nir/opt_if: fix fighting between split_alu_of_phi and peel_initial_break",
"nominated": true,
"nomination_type": 1,
"resolution": 0,
"resolution": 1,
"main_sha": null,
"because_sha": "6b4b04473986c9b0e77c925a116be39f6ff3982f",
"notes": null

View file

@ -3451,3 +3451,18 @@ nir_static_workgroup_size(const nir_shader *s)
return s->info.workgroup_size[0] * s->info.workgroup_size[1] *
s->info.workgroup_size[2];
}
bool
nir_block_contains_work(nir_block *block)
{
if (!nir_cf_node_is_last(&block->cf_node))
return true;
nir_foreach_instr(instr, block) {
if (instr->type != nir_instr_type_alu ||
!nir_op_is_vec_or_mov(nir_instr_as_alu(instr)->op))
return true;
}
return false;
}

View file

@ -3093,6 +3093,8 @@ nir_block_ends_in_break(nir_block *block)
nir_instr_as_jump(instr)->type == nir_jump_break;
}
bool nir_block_contains_work(nir_block *block);
#define nir_foreach_instr(instr, block) \
foreach_list_typed(nir_instr, instr, node, &(block)->instr_list)
#define nir_foreach_instr_reverse(instr, block) \

View file

@ -304,13 +304,6 @@ is_trivial_bcsel(const nir_instr *instr, bool allow_non_phi_src)
return true;
}
static bool
is_block_empty(nir_block *block)
{
return nir_cf_node_is_last(&block->cf_node) &&
exec_list_is_empty(&block->instr_list);
}
/**
* Splits ALU instructions that have a source that is a phi node
*
@ -395,8 +388,10 @@ opt_split_alu_of_phi(nir_builder *b, nir_loop *loop, nir_opt_if_options options)
if (continue_block == header_block)
return false;
/* If the continue block is otherwise empty, leave it that way. */
if (is_block_empty(continue_block))
/* If the continue block is otherwise empty, leave it that way. This must match
* opt_loop_peel_initial_break so that this optimization doesn't fight that one.
*/
if (!nir_block_contains_work(continue_block))
return false;
nir_foreach_instr_safe(instr, header_block) {
@ -701,6 +696,13 @@ opt_simplify_bcsel_of_phi(nir_builder *b, nir_loop *loop)
return progress;
}
static bool
is_block_empty(nir_block *block)
{
return nir_cf_node_is_last(&block->cf_node) &&
exec_list_is_empty(&block->instr_list);
}
/* Walk all the phis in the block immediately following the if statement and
* swap the blocks.
*/

View file

@ -297,23 +297,6 @@ opt_loop_last_block(nir_block *block, bool is_trivial_continue, bool is_trivial_
return progress;
}
static bool
block_contains_work(nir_block *block)
{
if (!nir_cf_node_is_last(&block->cf_node))
return true;
if (exec_list_is_empty(&block->instr_list))
return false;
/* Return false if the block contains only move-instructions. */
nir_foreach_instr(instr, block) {
if (instr->type != nir_instr_type_alu ||
!nir_op_is_vec_or_mov(nir_instr_as_alu(instr)->op))
return true;
}
return false;
}
static bool
can_constant_fold(nir_scalar scalar, nir_block *loop_header)
{
@ -403,7 +386,7 @@ opt_loop_peel_initial_break(nir_loop *loop)
return false;
/* Check that there is actual work to be done after the initial break. */
if (!block_contains_work(nir_cf_node_cf_tree_next(if_node)))
if (!nir_block_contains_work(nir_cf_node_cf_tree_next(if_node)))
return false;
/* For now, we restrict this optimization to cases where the outer IF