ir3: add workaround for predication hardware bug

Predication instructions sometimes need extra nops to workaround what
seems to be a hardware bug: prede needs 6 nops and the second
predt/predf of a predt/predf pair needs 4 nops.

The prede workaround is enabled starting from a6xx gen3 and the
predf/predt workaround from a6xx gen4, following the blob.

Fixes rendering corruption in God of War (2018).

Signed-off-by: Job Noorman <job@noorman.info>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32366>

src/freedreno/common/freedreno_dev_info.h

@@ -193,6 +193,16 @@ struct fd_dev_info {
        */
       bool has_ubwc_linear_mipmap_fallback;
 
+      /* Whether 4 nops are needed after the second pred[tf] of a
+       * pred[tf]/pred[ft] pair to work around a hardware issue.
+       */
+      bool predtf_nop_quirk;
+
+      /* Whether 6 nops are needed after prede to work around a hardware
+       * issue.
+       */
+      bool prede_nop_quirk;
+
       struct {
          uint32_t PC_POWER_CNTL;
          uint32_t TPL1_DBG_ECO_CNTL;

src/freedreno/common/freedreno_devices.py

@@ -400,6 +400,7 @@ a6xx_gen3 = A6XXProps(
         has_per_view_viewport = True,
         has_scalar_alu = True,
         has_early_preamble = True,
+        prede_nop_quirk = True,
     )
 
 a6xx_gen4 = A6XXProps(
@@ -431,6 +432,8 @@ a6xx_gen4 = A6XXProps(
         # TODO: there seems to be a quirk where at least rcp can't be in an
         # early preamble. a660 at least is affected.
         #has_early_preamble = True,
+        prede_nop_quirk = True,
+        predtf_nop_quirk = True,
     )
 
 add_gpus([
@@ -859,6 +862,8 @@ a7xx_base = A6XXProps(
         has_early_preamble = True,
         has_attachment_shading_rate = True,
         has_ubwc_linear_mipmap_fallback = True,
+        prede_nop_quirk = True,
+        predtf_nop_quirk = True,
     )
 
 a7xx_gen1 = A7XXProps(

src/freedreno/ir3/ir3_compiler.c

@@ -216,6 +216,8 @@ ir3_compiler_create(struct fd_device *dev, const struct fd_dev_id *dev_id,
       compiler->bitops_can_write_predicates = true;
       compiler->has_branch_and_or = true;
       compiler->has_predication = true;
+      compiler->predtf_nop_quirk = dev_info->a6xx.predtf_nop_quirk;
+      compiler->prede_nop_quirk = dev_info->a6xx.prede_nop_quirk;
       compiler->has_scalar_alu = dev_info->a6xx.has_scalar_alu;
       compiler->has_isam_v = dev_info->a6xx.has_isam_v;
       compiler->has_ssbo_imm_offsets = dev_info->a6xx.has_ssbo_imm_offsets;

src/freedreno/ir3/ir3_compiler.h

@@ -218,6 +218,8 @@ struct ir3_compiler {
 
    /* True if predt/predf/prede are supported. */
    bool has_predication;
+   bool predtf_nop_quirk;
+   bool prede_nop_quirk;
 
    /* MAX_COMPUTE_VARIABLE_GROUP_INVOCATIONS_ARB */
    uint32_t max_variable_workgroup_size;

src/freedreno/ir3/ir3_legalize.c

@@ -1206,6 +1206,31 @@ block_sched(struct ir3 *ir)
    }
 }
 
+/* Some gens have a hardware issue that needs to be worked around by 1)
+ * inserting 4 nops after the second pred[tf] of a pred[tf]/pred[ft] pair and/or
+ * inserting 6 nops after prede.
+ *
+ * This function should be called with the second pred[tf] of such a pair and
+ * NULL if there is only one pred[tf].
+ */
+static void
+add_predication_workaround(struct ir3_compiler *compiler,
+                           struct ir3_instruction *predtf,
+                           struct ir3_instruction *prede)
+{
+   if (predtf && compiler->predtf_nop_quirk) {
+      struct ir3_instruction *nop = ir3_NOP(predtf->block);
+      nop->repeat = 4;
+      ir3_instr_move_after(nop, predtf);
+   }
+
+   if (compiler->prede_nop_quirk) {
+      struct ir3_instruction *nop = ir3_NOP(prede->block);
+      nop->repeat = 6;
+      ir3_instr_move_after(nop, prede);
+   }
+}
+
 static void
 prede_sched(struct ir3 *ir)
 {
@@ -1275,7 +1300,8 @@ prede_sched(struct ir3 *ir)
        *        |----------|
        */
       if (!list_is_empty(&succ1->instr_list)) {
-         ir3_PREDE(succ1);
+         struct ir3_instruction *prede = ir3_PREDE(succ1);
+         add_predication_workaround(ir->compiler, succ0_terminator, prede);
          continue;
       }
 
@@ -1295,7 +1321,8 @@ prede_sched(struct ir3 *ir)
        *        |----------|
        */
       list_delinit(&succ0_terminator->node);
-      ir3_PREDE(succ0);
+      struct ir3_instruction *prede = ir3_PREDE(succ0);
+      add_predication_workaround(ir->compiler, NULL, prede);
       remove_unused_block(succ1);
       block->successors[1] = succ0->successors[0];
       ir3_block_add_predecessor(succ0->successors[0], block);