diff --git a/.pick_status.json b/.pick_status.json index 7075a3cd3fd..b2ac4eb4305 100644 --- a/.pick_status.json +++ b/.pick_status.json @@ -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 diff --git a/src/compiler/nir/nir.c b/src/compiler/nir/nir.c index 98e645b158d..f48cd019bb0 100644 --- a/src/compiler/nir/nir.c +++ b/src/compiler/nir/nir.c @@ -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; +} diff --git a/src/compiler/nir/nir.h b/src/compiler/nir/nir.h index 8f45c1aac2d..7f08e20b44f 100644 --- a/src/compiler/nir/nir.h +++ b/src/compiler/nir/nir.h @@ -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) \ diff --git a/src/compiler/nir/nir_opt_if.c b/src/compiler/nir/nir_opt_if.c index 4b0aee05f77..4734d5c004b 100644 --- a/src/compiler/nir/nir_opt_if.c +++ b/src/compiler/nir/nir_opt_if.c @@ -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. */ diff --git a/src/compiler/nir/nir_opt_loop.c b/src/compiler/nir/nir_opt_loop.c index f52f23843aa..a5a54f41975 100644 --- a/src/compiler/nir/nir_opt_loop.c +++ b/src/compiler/nir/nir_opt_loop.c @@ -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