From cdef824b7ad904366dc92deb4f94d707b47a51b7 Mon Sep 17 00:00:00 2001 From: Caio Oliveira Date: Fri, 10 May 2024 13:44:44 -0700 Subject: [PATCH] brw: Include some NIR states in the debug archive Acked-by: Kenneth Graunke Acked-by: Lionel Landwerlin Part-of: --- src/intel/compiler/brw_compile_bs.cpp | 6 +++++- src/intel/compiler/brw_compile_cs.cpp | 5 ++++- src/intel/compiler/brw_compile_fs.cpp | 6 ++++-- src/intel/compiler/brw_compile_gs.cpp | 5 ++++- src/intel/compiler/brw_compile_mesh.cpp | 18 ++++++++++++++++-- src/intel/compiler/brw_compile_tcs.cpp | 5 ++++- src/intel/compiler/brw_compile_tes.cpp | 6 ++++-- src/intel/compiler/brw_compile_vs.cpp | 5 ++++- src/intel/compiler/brw_compiler.h | 1 + src/intel/compiler/brw_nir.c | 19 +++++++++++++++---- src/intel/compiler/brw_nir.h | 2 ++ src/intel/compiler/brw_private.h | 9 +++++++++ src/intel/compiler/brw_shader.cpp | 20 ++++++++++++++++++++ 13 files changed, 92 insertions(+), 15 deletions(-) diff --git a/src/intel/compiler/brw_compile_bs.cpp b/src/intel/compiler/brw_compile_bs.cpp index ae23d25c097..7a71f863c48 100644 --- a/src/intel/compiler/brw_compile_bs.cpp +++ b/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 = { diff --git a/src/intel/compiler/brw_compile_cs.cpp b/src/intel/compiler/brw_compile_cs.cpp index f018fe9647e..062d444eb9d 100644 --- a/src/intel/compiler/brw_compile_cs.cpp +++ b/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 = { diff --git a/src/intel/compiler/brw_compile_fs.cpp b/src/intel/compiler/brw_compile_fs.cpp index 6eef9eb497c..3c8b500af1b 100644 --- a/src/intel/compiler/brw_compile_fs.cpp +++ b/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)); diff --git a/src/intel/compiler/brw_compile_gs.cpp b/src/intel/compiler/brw_compile_gs.cpp index ec4f57fc233..9cc6446007d 100644 --- a/src/intel/compiler/brw_compile_gs.cpp +++ b/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, ¶ms->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 = diff --git a/src/intel/compiler/brw_compile_mesh.cpp b/src/intel/compiler/brw_compile_mesh.cpp index da3b41d82f7..8c3e0fbe8be 100644 --- a/src/intel/compiler/brw_compile_mesh.cpp +++ b/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 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, ¶ms->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 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 = { diff --git a/src/intel/compiler/brw_compile_tcs.cpp b/src/intel/compiler/brw_compile_tcs.cpp index 9d5e303dbce..6db72faf381 100644 --- a/src/intel/compiler/brw_compile_tcs.cpp +++ b/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, ¶ms->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 = diff --git a/src/intel/compiler/brw_compile_tes.cpp b/src/intel/compiler/brw_compile_tes.cpp index fa1acdc8a6e..c26a293f6f3 100644 --- a/src/intel/compiler/brw_compile_tes.cpp +++ b/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, ¶ms->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) ? diff --git a/src/intel/compiler/brw_compile_vs.cpp b/src/intel/compiler/brw_compile_vs.cpp index 5c8f1361879..58b54d39d34 100644 --- a/src/intel/compiler/brw_compile_vs.cpp +++ b/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, ¶ms->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 = diff --git a/src/intel/compiler/brw_compiler.h b/src/intel/compiler/brw_compiler.h index 191699e478f..c382ab4bd26 100644 --- a/src/intel/compiler/brw_compiler.h +++ b/src/intel/compiler/brw_compiler.h @@ -26,6 +26,7 @@ #include #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" diff --git a/src/intel/compiler/brw_nir.c b/src/intel/compiler/brw_nir.c index 3cac5eee9ea..53084b08475 100644 --- a/src/intel/compiler/brw_nir.c +++ b/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 diff --git a/src/intel/compiler/brw_nir.h b/src/intel/compiler/brw_nir.h index 330d6c4a25c..38a1c98d861 100644 --- a/src/intel/compiler/brw_nir.h +++ b/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); diff --git a/src/intel/compiler/brw_private.h b/src/intel/compiler/brw_private.h index 87cb800e4a8..9a2f04e0bff 100644 --- a/src/intel/compiler/brw_private.h +++ b/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 diff --git a/src/intel/compiler/brw_shader.cpp b/src/intel/compiler/brw_shader.cpp index 24d1d8b36d3..cab0a3d0215 100644 --- a/src/intel/compiler/brw_shader.cpp +++ b/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)