brw: Include some NIR states in the debug archive

Acked-by: Kenneth Graunke <kenneth@whitecape.org>
Acked-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29146>

src/intel/compiler/brw_compile_bs.cpp

@@ -84,7 +84,11 @@ compile_single_bs(const struct brw_compiler *compiler,
    const unsigned required_width = compiler->devinfo->ver >= 20 ? 16u : 8u;
 
    brw_nir_apply_key(shader, compiler, &key->base, required_width);
-   brw_postprocess_nir(shader, compiler, debug_enabled,
+
+   brw_debug_archive_nir(params->base.archiver, shader, required_width, "first");
+
+   brw_postprocess_nir(shader, compiler, required_width,
+                       params->base.archiver, debug_enabled,
                        key->base.robust_flags);
 
    const brw_shader_params shader_params = {

src/intel/compiler/brw_compile_cs.cpp

@@ -199,6 +199,8 @@ brw_compile_cs(const struct brw_compiler *compiler,
       const unsigned dispatch_width = 8u << simd;
 
       nir_shader *shader = nir_shader_clone(params->base.mem_ctx, nir);
+      brw_debug_archive_nir(params->base.archiver, shader, dispatch_width, "first");
+
       brw_nir_apply_key(shader, compiler, &key->base,
                         dispatch_width);
 
@@ -208,7 +210,8 @@ brw_compile_cs(const struct brw_compiler *compiler,
       NIR_PASS(_, shader, nir_opt_constant_folding);
       NIR_PASS(_, shader, nir_opt_dce);
 
-      brw_postprocess_nir(shader, compiler, debug_enabled,
+      brw_postprocess_nir(shader, compiler, dispatch_width,
+                          params->base.archiver, debug_enabled,
                           key->base.robust_flags);
 
       const brw_shader_params shader_params = {

src/intel/compiler/brw_compile_fs.cpp

@@ -1499,6 +1499,8 @@ brw_compile_fs(const struct brw_compiler *compiler,
    const unsigned max_subgroup_size = 32;
    unsigned max_polygons = MAX2(1, params->max_polygons);
 
+   brw_debug_archive_nir(params->base.archiver, nir, 0, "first");
+
    brw_nir_apply_key(nir, compiler, &key->base, max_subgroup_size);
 
    if (brw_nir_fragment_shader_needs_wa_18019110168(devinfo, key->mesh_input, nir)) {
@@ -1556,8 +1558,8 @@ brw_compile_fs(const struct brw_compiler *compiler,
       NIR_PASS(_, nir, nir_inline_sysval, nir_intrinsic_load_fs_msaa_intel, f);
    }
 
-   brw_postprocess_nir(nir, compiler, debug_enabled,
-                       key->base.robust_flags);
+   brw_postprocess_nir(nir, compiler, 0, params->base.archiver,
+                       debug_enabled, key->base.robust_flags);
 
    int per_primitive_offsets[VARYING_SLOT_MAX];
    memset(per_primitive_offsets, -1, sizeof(per_primitive_offsets));

src/intel/compiler/brw_compile_gs.cpp

@@ -146,6 +146,8 @@ brw_compile_gs(const struct brw_compiler *compiler,
 
    const bool debug_enabled = brw_should_print_shader(nir, DEBUG_GS, params->base.source_hash);
 
+   brw_debug_archive_nir(params->base.archiver, nir, dispatch_width, "first");
+
    brw_prog_data_init(&prog_data->base.base, &params->base);
 
    /* The GLSL linker will have already matched up GS inputs and the outputs
@@ -174,7 +176,8 @@ brw_compile_gs(const struct brw_compiler *compiler,
    brw_nir_apply_key(nir, compiler, &key->base, dispatch_width);
    brw_nir_lower_vue_inputs(nir, &input_vue_map);
    brw_nir_lower_vue_outputs(nir);
-   brw_postprocess_nir(nir, compiler, debug_enabled,
+   brw_postprocess_nir(nir, compiler, dispatch_width,
+                       params->base.archiver, debug_enabled,
                        key->base.robust_flags);
 
    const bool has_clip_cull_dist =

src/intel/compiler/brw_compile_mesh.cpp

@@ -354,6 +354,8 @@ brw_compile_task(const struct brw_compiler *compiler,
    struct brw_task_prog_data *prog_data = params->prog_data;
    const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TASK, params->base.source_hash);
 
+   brw_debug_archive_nir(params->base.archiver, nir, 0, "first");
+
    brw_nir_lower_tue_outputs(nir, &prog_data->map);
 
    NIR_PASS(_, nir, brw_nir_align_launch_mesh_workgroups);
@@ -388,6 +390,8 @@ brw_compile_task(const struct brw_compiler *compiler,
       .required_width = brw_required_dispatch_width(&nir->info),
    };
 
+   brw_debug_archive_nir(params->base.archiver, nir, 0, "before-simd");
+
    std::unique_ptr<brw_shader> v[3];
 
    for (unsigned i = 0; i < 3; i++) {
@@ -401,9 +405,12 @@ brw_compile_task(const struct brw_compiler *compiler,
       nir_shader *shader = nir_shader_clone(params->base.mem_ctx, nir);
       brw_nir_apply_key(shader, compiler, &key->base, dispatch_width);
 
+      brw_debug_archive_nir(params->base.archiver, shader, dispatch_width, "first");
+
       NIR_PASS(_, shader, brw_nir_lower_simd, dispatch_width);
 
-      brw_postprocess_nir(shader, compiler, debug_enabled,
+      brw_postprocess_nir(shader, compiler, dispatch_width,
+                          params->base.archiver, debug_enabled,
                           key->base.robust_flags);
 
       const brw_shader_params shader_params = {
@@ -1164,6 +1171,8 @@ brw_compile_mesh(const struct brw_compiler *compiler,
    struct brw_mesh_prog_data *prog_data = params->prog_data;
    const bool debug_enabled = brw_should_print_shader(nir, DEBUG_MESH, params->base.source_hash);
 
+   brw_debug_archive_nir(params->base.archiver, nir, 0, "first");
+
    brw_prog_data_init(&prog_data->base.base, &params->base);
 
    prog_data->base.local_size[0] = nir->info.workgroup_size[0];
@@ -1233,6 +1242,8 @@ brw_compile_mesh(const struct brw_compiler *compiler,
 
    std::unique_ptr<brw_shader> v[3];
 
+   brw_debug_archive_nir(params->base.archiver, nir, 0, "before-simd");
+
    for (unsigned i = 0; i < 3; i++) {
       const unsigned simd = devinfo->ver >= 30 ? 2 - i : i;
 
@@ -1243,6 +1254,8 @@ brw_compile_mesh(const struct brw_compiler *compiler,
 
       nir_shader *shader = nir_shader_clone(params->base.mem_ctx, nir);
 
+      brw_debug_archive_nir(params->base.archiver, shader, dispatch_width, "first");
+
       /*
        * When Primitive Header is enabled, we may not generates writes to all
        * fields, so let's initialize everything.
@@ -1258,7 +1271,8 @@ brw_compile_mesh(const struct brw_compiler *compiler,
 
       NIR_PASS(_, shader, brw_nir_lower_simd, dispatch_width);
 
-      brw_postprocess_nir(shader, compiler, debug_enabled,
+      brw_postprocess_nir(shader, compiler, dispatch_width,
+                          params->base.archiver, debug_enabled,
                           key->base.robust_flags);
 
       const brw_shader_params shader_params = {

src/intel/compiler/brw_compile_tcs.cpp

@@ -192,6 +192,8 @@ brw_compile_tcs(const struct brw_compiler *compiler,
 
    const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TCS, params->base.source_hash);
 
+   brw_debug_archive_nir(params->base.archiver, nir, dispatch_width, "first");
+
    brw_prog_data_init(&prog_data->base.base, &params->base);
 
    nir->info.outputs_written = key->outputs_written;
@@ -211,7 +213,8 @@ brw_compile_tcs(const struct brw_compiler *compiler,
                              key->_tes_primitive_mode);
    intel_nir_lower_patch_vertices_in(nir, key->input_vertices);
 
-   brw_postprocess_nir(nir, compiler, debug_enabled,
+   brw_postprocess_nir(nir, compiler, dispatch_width,
+                       params->base.archiver, debug_enabled,
                        key->base.robust_flags);
 
    bool has_primitive_id =

src/intel/compiler/brw_compile_tes.cpp

@@ -72,6 +72,8 @@ brw_compile_tes(const struct brw_compiler *compiler,
 
    const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TES, params->base.source_hash);
 
+   brw_debug_archive_nir(params->base.archiver, nir, dispatch_width, "first");
+
    brw_prog_data_init(&prog_data->base.base, &params->base);
 
    if (params->input_vue_map != NULL) {
@@ -91,8 +93,8 @@ brw_compile_tes(const struct brw_compiler *compiler,
    brw_nir_lower_tes_inputs(nir, &input_vue_map);
    brw_nir_lower_vue_outputs(nir);
    NIR_PASS(_, nir, intel_nir_lower_patch_vertices_tes);
-   brw_postprocess_nir(nir, compiler, debug_enabled,
-                       key->base.robust_flags);
+   brw_postprocess_nir(nir, compiler, dispatch_width, params->base.archiver,
+                       debug_enabled, key->base.robust_flags);
 
    const uint32_t pos_slots =
       (nir->info.per_view_outputs & VARYING_BIT_POS) ?

src/intel/compiler/brw_compile_vs.cpp

@@ -246,6 +246,8 @@ brw_compile_vs(const struct brw_compiler *compiler,
     */
    assert(!key->no_vf_slot_compaction || key->vf_component_packing);
 
+   brw_debug_archive_nir(params->base.archiver, nir, dispatch_width, "first");
+
    brw_prog_data_init(&prog_data->base.base, &params->base);
 
    /* When using Primitive Replication for multiview, each view gets its own
@@ -277,7 +279,8 @@ brw_compile_vs(const struct brw_compiler *compiler,
    if (key->vf_component_packing)
       nr_packed_regs = brw_nir_pack_vs_input(nir, prog_data);
 
-   brw_postprocess_nir(nir, compiler, debug_enabled,
+   brw_postprocess_nir(nir, compiler, dispatch_width,
+                       params->base.archiver, debug_enabled,
                        key->base.robust_flags);
 
    const bool has_clip_cull_dist =

src/intel/compiler/brw_compiler.h

@@ -26,6 +26,7 @@
 #include <stdio.h>
 #include "c11/threads.h"
 #include "dev/intel_device_info.h"
+#include "mda/debug_archiver.h"
 #include "isl/isl.h"
 #include "mda/debug_archiver.h"
 #include "util/macros.h"

src/intel/compiler/brw_nir.c

@@ -23,6 +23,7 @@
 
 #include "intel_nir.h"
 #include "brw_nir.h"
+#include "brw_private.h"
 #include "compiler/glsl_types.h"
 #include "compiler/nir/nir_builder.h"
 #include "dev/intel_debug.h"
@@ -2122,6 +2123,8 @@ nir_shader_has_local_variables(const nir_shader *nir)
  */
 void
 brw_postprocess_nir(nir_shader *nir, const struct brw_compiler *compiler,
+                    unsigned dispatch_width,
+                    debug_archiver *archiver,
                     bool debug_enabled,
                     enum brw_robustness_flags robust_flags)
 {
@@ -2368,15 +2371,20 @@ brw_postprocess_nir(nir_shader *nir, const struct brw_compiler *compiler,
 
    OPT(nir_lower_locals_to_regs, 32);
 
-   if (unlikely(debug_enabled)) {
+   if (unlikely(debug_enabled || archiver)) {
       /* Re-index SSA defs so we print more sensible numbers. */
       nir_foreach_function_impl(impl, nir) {
          nir_index_ssa_defs(impl);
       }
 
-      fprintf(stderr, "NIR (SSA form) for %s shader:\n",
-              _mesa_shader_stage_to_string(nir->info.stage));
-      nir_print_shader(nir, stderr);
+      if (debug_enabled) {
+         fprintf(stderr, "NIR (SSA form) for %s shader:\n",
+                 _mesa_shader_stage_to_string(nir->info.stage));
+         nir_print_shader(nir, stderr);
+      }
+
+      if (unlikely(archiver))
+         brw_debug_archive_nir(archiver, nir, dispatch_width, "ssa");
    }
 
    nir_validate_ssa_dominance(nir, "before nir_convert_from_ssa");
@@ -2412,6 +2420,9 @@ brw_postprocess_nir(nir_shader *nir, const struct brw_compiler *compiler,
               _mesa_shader_stage_to_string(nir->info.stage));
       nir_print_shader(nir, stderr);
    }
+
+   if (unlikely(archiver))
+      brw_debug_archive_nir(archiver, nir, dispatch_width, "out");
 }
 
 static unsigned

src/intel/compiler/brw_nir.h

@@ -234,6 +234,8 @@ bool brw_nir_lower_simd(nir_shader *nir, unsigned dispatch_width);
 
 void brw_postprocess_nir(nir_shader *nir,
                          const struct brw_compiler *compiler,
+                         unsigned dispatch_width,
+                         debug_archiver *archiver,
                          bool debug_enabled,
                          enum brw_robustness_flags robust_flags);
 

src/intel/compiler/brw_private.h

@@ -37,6 +37,15 @@ void brw_alloc_reg_sets(struct brw_compiler *compiler);
 extern const char *const conditional_modifier[16];
 extern const char *const pred_ctrl_align16[16];
 
+#ifndef NDEBUG
+void brw_debug_archive_nir(debug_archiver *archiver, nir_shader *nir,
+                           unsigned dispatch_width, const char *step);
+#else
+static inline void
+brw_debug_archive_nir(debug_archiver *archiver, nir_shader *nir,
+                      unsigned dispatch_width, const char *step) {}
+#endif
+
 #ifdef __cplusplus
 }
 #endif

src/intel/compiler/brw_shader.cpp

@@ -1317,6 +1317,26 @@ brw_allocate_registers(brw_shader &s, bool allow_spilling)
    s.debug_optimizer(nir, "scoreboard", 96, pass_num++);
 }
 
+#ifndef NDEBUG
+void
+brw_debug_archive_nir(debug_archiver *archiver, nir_shader *nir,
+                      unsigned dispatch_width, const char *step)
+{
+   if (!archiver)
+      return;
+
+   const bool prefix_dispatch_width =
+      dispatch_width > 0 && mesa_shader_stage_uses_workgroup(nir->info.stage);
+   const char *filename = prefix_dispatch_width ?
+         ralloc_asprintf(archiver, "NIR%d/%s", dispatch_width, step) :
+         ralloc_asprintf(archiver, "NIR/%s", step);
+
+   FILE *f = debug_archiver_start_file(archiver, filename);
+   nir_print_shader(nir, f);
+   debug_archiver_finish_file(archiver);
+}
+#endif
+
 unsigned
 brw_cs_push_const_total_size(const struct brw_cs_prog_data *cs_prog_data,
                              unsigned threads)