diff --git a/src/intel/compiler/brw_compiler.h b/src/intel/compiler/brw_compiler.h
index 96d771e4b59..c8e6f05c38f 100644
--- a/src/intel/compiler/brw_compiler.h
+++ b/src/intel/compiler/brw_compiler.h
@@ -1641,8 +1641,10 @@ brw_stage_has_packed_dispatch(ASSERTED const struct intel_device_info *devinfo,
    /* The code below makes assumptions about the hardware's thread dispatch
     * behavior that could be proven wrong in future generations -- Make sure
     * to do a full test run with brw_fs_test_dispatch_packing() hooked up to
-    * the NIR front-end before changing this assertion.
+    * the NIR front-end before changing this assertion. It can be temporarily
+    * enabled by setting the macro below to true.
     */
+   #define ENABLE_FS_TEST_DISPATCH_PACKING false
    assert(devinfo->ver <= 12);
 
    switch (stage) {
diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp
index 12514838f94..d3c9ef0f67e 100644
--- a/src/intel/compiler/brw_fs.cpp
+++ b/src/intel/compiler/brw_fs.cpp
@@ -4499,40 +4499,6 @@ brw_compile_bs(const struct brw_compiler *compiler,
    return g.get_assembly();
 }
 
-/**
- * Test the dispatch mask packing assumptions of
- * brw_stage_has_packed_dispatch().  Call this from e.g. the top of
- * fs_visitor::emit_nir_code() to cause a GPU hang if any shader invocation is
- * executed with an unexpected dispatch mask.
- */
-static UNUSED void
-brw_fs_test_dispatch_packing(const fs_builder &bld)
-{
-   const fs_visitor *shader = bld.shader;
-   const gl_shader_stage stage = shader->stage;
-   const bool uses_vmask =
-      stage == MESA_SHADER_FRAGMENT &&
-      brw_wm_prog_data(shader->prog_data)->uses_vmask;
-
-   if (brw_stage_has_packed_dispatch(shader->devinfo, stage,
-                                     shader->max_polygons,
-                                     shader->prog_data)) {
-      const fs_builder ubld = bld.exec_all().group(1, 0);
-      const fs_reg tmp = component(bld.vgrf(BRW_TYPE_UD), 0);
-      const fs_reg mask = uses_vmask ? brw_vmask_reg() : brw_dmask_reg();
-
-      ubld.ADD(tmp, mask, brw_imm_ud(1));
-      ubld.AND(tmp, mask, tmp);
-
-      /* This will loop forever if the dispatch mask doesn't have the expected
-       * form '2^n-1', in which case tmp will be non-zero.
-       */
-      bld.emit(BRW_OPCODE_DO);
-      bld.CMP(bld.null_reg_ud(), tmp, brw_imm_ud(0), BRW_CONDITIONAL_NZ);
-      set_predicate(BRW_PREDICATE_NORMAL, bld.emit(BRW_OPCODE_WHILE));
-   }
-}
-
 unsigned
 fs_visitor::workgroup_size() const
 {
diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp
index 992fdb3bd7f..3d7463af6d3 100644
--- a/src/intel/compiler/brw_fs_nir.cpp
+++ b/src/intel/compiler/brw_fs_nir.cpp
@@ -8446,6 +8446,40 @@ emit_shader_float_controls_execution_mode(nir_to_brw_state &ntb)
              brw_imm_d(mode), brw_imm_d(mask));
 }
 
+/**
+ * Test the dispatch mask packing assumptions of
+ * brw_stage_has_packed_dispatch().  Call this from e.g. the top of
+ * nir_to_brw() to cause a GPU hang if any shader invocation is
+ * executed with an unexpected dispatch mask.
+ */
+static UNUSED void
+brw_fs_test_dispatch_packing(const fs_builder &bld)
+{
+   const fs_visitor *shader = bld.shader;
+   const gl_shader_stage stage = shader->stage;
+   const bool uses_vmask =
+      stage == MESA_SHADER_FRAGMENT &&
+      brw_wm_prog_data(shader->prog_data)->uses_vmask;
+
+   if (brw_stage_has_packed_dispatch(shader->devinfo, stage,
+                                     shader->max_polygons,
+                                     shader->prog_data)) {
+      const fs_builder ubld = bld.exec_all().group(1, 0);
+      const fs_reg tmp = component(bld.vgrf(BRW_TYPE_UD), 0);
+      const fs_reg mask = uses_vmask ? brw_vmask_reg() : brw_dmask_reg();
+
+      ubld.ADD(tmp, mask, brw_imm_ud(1));
+      ubld.AND(tmp, mask, tmp);
+
+      /* This will loop forever if the dispatch mask doesn't have the expected
+       * form '2^n-1', in which case tmp will be non-zero.
+       */
+      bld.emit(BRW_OPCODE_DO);
+      bld.CMP(bld.null_reg_ud(), tmp, brw_imm_ud(0), BRW_CONDITIONAL_NZ);
+      set_predicate(BRW_PREDICATE_NORMAL, bld.emit(BRW_OPCODE_WHILE));
+   }
+}
+
 void
 nir_to_brw(fs_visitor *s)
 {
@@ -8457,6 +8491,9 @@ nir_to_brw(fs_visitor *s)
       .bld     = fs_builder(s).at_end(),
    };
 
+   if (ENABLE_FS_TEST_DISPATCH_PACKING)
+      brw_fs_test_dispatch_packing(ntb.bld);
+
    for (unsigned i = 0; i < s->nir->printf_info_count; i++) {
       brw_stage_prog_data_add_printf(s->prog_data,
                                      s->mem_ctx,