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>

src/compiler/nir/nir.c

@@ -3499,3 +3499,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;
+}

src/compiler/nir/nir.h

@@ -3130,6 +3130,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) \

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.
  */

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