nir: add a new nir_jump_abort

Some backends need to emit different code between nir_jump_halt and
nir_jump_abort.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/42162>

src/broadcom/compiler/nir_to_vir.c

@@ -4501,6 +4501,7 @@ ntq_emit_jump(struct v3d_compile *c, nir_jump_instr *jump)
                 break;
 
         case nir_jump_halt:
+        case nir_jump_abort:
         case nir_jump_goto:
         case nir_jump_goto_if:
                 UNREACHABLE("not supported\n");
@@ -4528,6 +4529,7 @@ ntq_emit_uniform_jump(struct v3d_compile *c, nir_jump_instr *jump)
                 break;
 
         case nir_jump_halt:
+        case nir_jump_abort:
         case nir_jump_goto:
         case nir_jump_goto_if:
                 UNREACHABLE("not supported\n");

src/compiler/nir/nir.h

@@ -2920,6 +2920,20 @@ typedef enum {
     */
    nir_jump_halt,
 
+   /** Immediately exit the current shader due to a fatal error
+    *
+    * This has the same CFG semantics as nir_jump_halt — it jumps to the end
+    * block of the shader entrypoint — but carries the additional semantic
+    * that the invocation is terminating because a fatal error was detected
+    * (e.g. nir_intrinsic_abort wrote a message to the abort buffer).
+    *
+    * Backends that support hardware-level abort signalling (device fault,
+    * special sendmsg, etc.) may generate different code for nir_jump_abort
+    * vs nir_jump_halt.  Backends that do not distinguish the two can lower
+    * nir_jump_abort to nir_jump_halt before instruction selection.
+    */
+   nir_jump_abort,
+
    /** Break out of the inner-most loop
     *
     * This has the same semantics as C's "break" statement.

src/compiler/nir/nir_control_flow.c

@@ -485,6 +485,7 @@ nir_handle_add_jump(nir_block *block)
    switch (jump_instr->type) {
    case nir_jump_return:
    case nir_jump_halt:
+   case nir_jump_abort:
       link_blocks(block, impl->end_block, NULL);
       break;
 

src/compiler/nir/nir_divergence_analysis.c

@@ -1280,6 +1280,7 @@ visit_jump(nir_jump_instr *jump, struct divergence_state *state)
          state->divergent_loop_break = true;
       return state->divergent_loop_break;
    case nir_jump_halt:
+   case nir_jump_abort:
       /* This totally kills invocations so it doesn't add divergence */
       break;
    case nir_jump_return:

src/compiler/nir/nir_lower_tex.c

@@ -1578,7 +1578,8 @@ nir_lower_tex_block(nir_block *block, nir_builder *b,
          continue;
       } else if (instr->type == nir_instr_type_jump) {
          if (nir_instr_as_jump(instr)->type == nir_jump_halt ||
-             nir_instr_as_jump(instr)->type == nir_jump_return)
+             nir_instr_as_jump(instr)->type == nir_jump_return ||
+             nir_instr_as_jump(instr)->type == nir_jump_abort)
             *prev_terminate_return = instr->index;
          continue;
       } else if (instr->type != nir_instr_type_tex) {

src/compiler/nir/nir_lower_workgroup_size.c

@@ -325,6 +325,7 @@ nlwgs_instr_splits_augmented_block(nir_instr *instr)
          return true;
       case nir_jump_halt:
       case nir_jump_return:
+      case nir_jump_abort:
       case nir_jump_goto:
       case nir_jump_goto_if:
          UNREACHABLE("halt/return/goto should have been already lowered");

src/compiler/nir/nir_opt_dead_cf.c

@@ -205,7 +205,8 @@ node_is_dead(nir_cf_node *node)
          if (instr->type == nir_instr_type_jump &&
              (!inside_loop ||
               nir_instr_as_jump(instr)->type == nir_jump_return ||
-              nir_instr_as_jump(instr)->type == nir_jump_halt))
+              nir_instr_as_jump(instr)->type == nir_jump_halt ||
+              nir_instr_as_jump(instr)->type == nir_jump_abort))
             return false;
 
          if (instr->type == nir_instr_type_intrinsic) {

src/compiler/nir/nir_print.c

@@ -2285,6 +2285,10 @@ print_jump_instr(nir_jump_instr *instr, print_state *state)
       fprintf(fp, "halt");
       break;
 
+   case nir_jump_abort:
+      fprintf(fp, "abort");
+      break;
+
    case nir_jump_goto:
       fprintf(fp, "goto b%u",
               instr->target ? instr->target->index : -1);

src/compiler/nir/nir_validate.c

@@ -1278,6 +1278,7 @@ validate_jump_instr(nir_jump_instr *instr, validate_state *state)
    switch (instr->type) {
    case nir_jump_return:
    case nir_jump_halt:
+   case nir_jump_abort:
       validate_assert(state, block->successors[0] == state->impl->end_block);
       validate_assert(state, block->successors[1] == NULL);
       validate_assert(state, instr->target == NULL);

src/kosmickrisp/compiler/nir_to_msl.c

@@ -1861,7 +1861,8 @@ jump_instr_to_msl(struct nir_to_msl_ctx *ctx, nir_jump_instr *jump)
 {
    switch (jump->type) {
    case nir_jump_halt:
-      P_IND(ctx, "TODO: halt\n");
+   case nir_jump_abort:
+      P_IND(ctx, "TODO: halt/abort\n");
       assert(!"Unimplemented");
       break;
    case nir_jump_break: