diff --git a/src/freedreno/common/freedreno_dev_info.h b/src/freedreno/common/freedreno_dev_info.h
index 4519bbbebc1..3f4784eb3cd 100644
--- a/src/freedreno/common/freedreno_dev_info.h
+++ b/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;
diff --git a/src/freedreno/common/freedreno_devices.py b/src/freedreno/common/freedreno_devices.py
index 793472f4308..2b5a3fcdf77 100644
--- a/src/freedreno/common/freedreno_devices.py
+++ b/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(
diff --git a/src/freedreno/ir3/ir3_compiler.c b/src/freedreno/ir3/ir3_compiler.c
index 79a77de8ada..2fbdff823d2 100644
--- a/src/freedreno/ir3/ir3_compiler.c
+++ b/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;
diff --git a/src/freedreno/ir3/ir3_compiler.h b/src/freedreno/ir3/ir3_compiler.h
index daa4fadc99f..0d702a6ec16 100644
--- a/src/freedreno/ir3/ir3_compiler.h
+++ b/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;
diff --git a/src/freedreno/ir3/ir3_legalize.c b/src/freedreno/ir3/ir3_legalize.c
index e4809922d9a..e13aaeb890f 100644
--- a/src/freedreno/ir3/ir3_legalize.c
+++ b/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);