diff --git a/src/amd/compiler/aco_insert_waitcnt.cpp b/src/amd/compiler/aco_insert_waitcnt.cpp
index 7ac2c822ff0..b7edf044629 100644
--- a/src/amd/compiler/aco_insert_waitcnt.cpp
+++ b/src/amd/compiler/aco_insert_waitcnt.cpp
@@ -551,21 +551,6 @@ kill(wait_imm& imm, alu_delay_info& delay, Instruction* instr, wait_ctx& ctx,
       imm.combine(ctx.barrier_imm[ffs(storage_gds) - 1]);
    }
 
-   if (ctx.program->early_rast && instr->opcode == aco_opcode::exp) {
-      if (instr->exp().dest >= V_008DFC_SQ_EXP_POS && instr->exp().dest < V_008DFC_SQ_EXP_PRIM) {
-
-         /* With early_rast, the HW will start clipping and rasterization after the 1st DONE pos
-          * export. Wait for all stores (and atomics) to complete, so PS can read them.
-          * TODO: This only really applies to DONE pos exports.
-          *       Consider setting the DONE bit earlier.
-          */
-         if (ctx.vs_cnt > 0)
-            imm.vs = 0;
-         if (ctx.vm_cnt > 0)
-            imm.vm = 0;
-      }
-   }
-
    if (instr->opcode == aco_opcode::p_barrier)
       perform_barrier(ctx, imm, instr->barrier().sync, semantic_acqrel);
    else
diff --git a/src/amd/compiler/aco_instruction_selection_setup.cpp b/src/amd/compiler/aco_instruction_selection_setup.cpp
index 2d16f39ba6f..d43c5966ed5 100644
--- a/src/amd/compiler/aco_instruction_selection_setup.cpp
+++ b/src/amd/compiler/aco_instruction_selection_setup.cpp
@@ -258,14 +258,6 @@ setup_vs_output_info(isel_context* ctx, nir_shader* nir)
    ctx->num_cull_distances = util_bitcount(outinfo->cull_dist_mask);
 
    assert(ctx->num_clip_distances + ctx->num_cull_distances <= 8);
-
-   /* GFX10+ early rasterization:
-    * When there are no param exports in an NGG (or legacy VS) shader,
-    * RADV sets NO_PC_EXPORT=1, which means the HW will start clipping and rasterization
-    * as soon as it encounters a DONE pos export. When this happens, PS waves can launch
-    * before the NGG (or VS) waves finish.
-    */
-   ctx->program->early_rast = ctx->program->gfx_level >= GFX10 && outinfo->param_exports == 0;
 }
 
 void
diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h
index baf64b04267..1ae5cf35b85 100644
--- a/src/amd/compiler/aco_ir.h
+++ b/src/amd/compiler/aco_ir.h
@@ -2168,7 +2168,6 @@ public:
    uint16_t min_waves = 0;
    unsigned workgroup_size; /* if known; otherwise UINT_MAX */
    bool wgp_mode;
-   bool early_rast = false; /* whether rasterization can start as soon as the 1st DONE pos export */
 
    bool needs_vcc = false;