intel/debug: shader dump filter

v2: Fixes filtering for various brw shader dump logic

Reviewed-by: José Roberto de Souza <jose.souza@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35061>

docs/envvars.rst

@@ -831,6 +831,11 @@ Intel driver environment variables
       are always dumped if :envvar:`INTEL_SHADER_BIN_DUMP_PATH` variable is
       set.
 
+.. envvar:: INTEL_SHADER_DUMP_FILTER
+
+   Only dump information about shaders that match the specified hexadecimal
+   source hash.
+
 .. envvar:: INTEL_SIMD_DEBUG
 
    a comma-separated list of named flags, which control simd dispatch widths:

src/gallium/drivers/iris/iris_program_cache.c

@@ -220,15 +220,18 @@ iris_upload_shader(struct iris_screen *screen,
    }
 
    if (INTEL_DEBUG(DEBUG_SHADERS_LINENO) && screen->brw) {
-      int start = 0;
-      /* dump each simd variant of shader */
-      while (start < shader->brw_prog_data->program_size) {
-         brw_disassemble_with_lineno(&screen->brw->isa, shader->stage, -1,
-                                    ish ? ish->source_hash : 0, assembly, start,
-                                    res->bo->address + shader->assembly.offset,
-                                    stderr);
-         start += align64(brw_disassemble_find_end(&screen->brw->isa,
-                                                   assembly, start), 64);
+      if (!intel_shader_dump_filter ||
+          (intel_shader_dump_filter && ish && intel_shader_dump_filter == ish->source_hash)) {
+         int start = 0;
+         /* dump each simd variant of shader */
+         while (start < shader->brw_prog_data->program_size) {
+            brw_disassemble_with_lineno(&screen->brw->isa, shader->stage, -1,
+                                        ish ? ish->source_hash : 0, assembly, start,
+                                        res->bo->address + shader->assembly.offset,
+                                        stderr);
+            start += align64(brw_disassemble_find_end(&screen->brw->isa,
+                                                      assembly, start), 64);
+         }
       }
    }
 }

src/intel/compiler/brw_compile_bs.cpp

@@ -73,7 +73,7 @@ compile_single_bs(const struct brw_compiler *compiler,
                   int *prog_offset,
                   uint64_t *bsr)
 {
-   const bool debug_enabled = brw_should_print_shader(shader, DEBUG_RT);
+   const bool debug_enabled = brw_should_print_shader(shader, DEBUG_RT, params->base.source_hash);
 
    prog_data->max_stack_size = MAX2(prog_data->max_stack_size,
                                     shader->scratch_size);
@@ -166,7 +166,7 @@ brw_compile_bs(const struct brw_compiler *compiler,
    struct brw_bs_prog_data *prog_data = params->prog_data;
    unsigned num_resume_shaders = params->num_resume_shaders;
    nir_shader **resume_shaders = params->resume_shaders;
-   const bool debug_enabled = brw_should_print_shader(shader, DEBUG_RT);
+   const bool debug_enabled = brw_should_print_shader(shader, DEBUG_RT, params->base.source_hash);
 
    brw_prog_data_init(&prog_data->base, &params->base);
 

src/intel/compiler/brw_compile_cs.cpp

@@ -131,7 +131,8 @@ brw_compile_cs(const struct brw_compiler *compiler,
 
    const bool debug_enabled =
       brw_should_print_shader(nir, params->base.debug_flag ?
-                                   params->base.debug_flag : DEBUG_CS);
+                                   params->base.debug_flag : DEBUG_CS,
+                                   params->base.source_hash);
 
    brw_prog_data_init(&prog_data->base, &params->base);
    prog_data->uses_inline_data = brw_nir_uses_inline_data(nir) ||

src/intel/compiler/brw_compile_fs.cpp

@@ -1449,7 +1449,8 @@ brw_compile_fs(const struct brw_compiler *compiler,
    bool allow_spilling = params->allow_spilling;
    const bool debug_enabled =
       brw_should_print_shader(nir, params->base.debug_flag ?
-                                   params->base.debug_flag : DEBUG_WM);
+                                   params->base.debug_flag : DEBUG_WM,
+                                   params->base.source_hash);
 
    brw_prog_data_init(&prog_data->base, &params->base);
 

src/intel/compiler/brw_compile_gs.cpp

@@ -146,7 +146,7 @@ brw_compile_gs(const struct brw_compiler *compiler,
    unsigned control_data_bits_per_vertex = 0;
    unsigned control_data_header_size_bits = 0;
 
-   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_GS);
+   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_GS, params->base.source_hash);
 
    brw_prog_data_init(&prog_data->base.base, &params->base);
 

src/intel/compiler/brw_compile_mesh.cpp

@@ -354,7 +354,7 @@ brw_compile_task(const struct brw_compiler *compiler,
    struct nir_shader *nir = params->base.nir;
    const struct brw_task_prog_key *key = params->key;
    struct brw_task_prog_data *prog_data = params->prog_data;
-   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TASK);
+   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TASK, params->base.source_hash);
 
    brw_nir_lower_tue_outputs(nir, &prog_data->map);
 
@@ -1143,7 +1143,7 @@ brw_compile_mesh(const struct brw_compiler *compiler,
    struct nir_shader *nir = params->base.nir;
    const struct brw_mesh_prog_key *key = params->key;
    struct brw_mesh_prog_data *prog_data = params->prog_data;
-   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_MESH);
+   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_MESH, params->base.source_hash);
 
    brw_prog_data_init(&prog_data->base.base, &params->base);
 

src/intel/compiler/brw_compile_tcs.cpp

@@ -191,7 +191,7 @@ brw_compile_tcs(const struct brw_compiler *compiler,
    struct brw_vue_prog_data *vue_prog_data = &prog_data->base;
    const unsigned dispatch_width = brw_geometry_stage_dispatch_width(compiler->devinfo);
 
-   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TCS);
+   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TCS, params->base.source_hash);
 
    brw_prog_data_init(&prog_data->base.base, &params->base);
 

src/intel/compiler/brw_compile_tes.cpp

@@ -69,7 +69,7 @@ brw_compile_tes(const struct brw_compiler *compiler,
    struct brw_tes_prog_data *prog_data = params->prog_data;
    const unsigned dispatch_width = brw_geometry_stage_dispatch_width(compiler->devinfo);
 
-   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TES);
+   const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TES, params->base.source_hash);
 
    brw_prog_data_init(&prog_data->base.base, &params->base);
 

src/intel/compiler/brw_compile_vs.cpp

@@ -235,7 +235,8 @@ brw_compile_vs(const struct brw_compiler *compiler,
    struct brw_vs_prog_data *prog_data = params->prog_data;
    const bool debug_enabled =
       brw_should_print_shader(nir, params->base.debug_flag ?
-                                   params->base.debug_flag : DEBUG_VS);
+                                   params->base.debug_flag : DEBUG_VS,
+                                   params->base.source_hash);
    const unsigned dispatch_width = brw_geometry_stage_dispatch_width(compiler->devinfo);
 
    /* We only expect slot compaction to be disabled when using device

src/intel/compiler/brw_generator.cpp

@@ -1418,27 +1418,30 @@ brw_generator::generate_code(const cfg_t *cfg, int dispatch_width,
    }
 
    if (unlikely(debug_flag)) {
-      fprintf(stderr, "Native code for %s (src_hash 0x%08x) (sha1 %s)\n"
-              "SIMD%d shader: %d instructions. %d loops. %u cycles. "
-              "%d:%d spills:fills, %u sends, "
-              "scheduled with mode %s. "
-              "Promoted %u constants. "
-              "Non-SSA regs (after NIR): %u. "
-              "Compacted %d to %d bytes (%.0f%%)\n",
-              shader_name, params->source_hash, sha1buf,
-              dispatch_width,
-              before_size / 16 - nop_count - sync_nop_count,
-              loop_count, perf.latency,
-              shader_stats.spill_count,
-              shader_stats.fill_count,
-              send_count,
-              shader_stats.scheduler_mode,
-              shader_stats.promoted_constants,
-              shader_stats.non_ssa_registers_after_nir,
-              before_size, after_size,
-              100.0f * (before_size - after_size) / before_size);
-      dump_assembly(p->store, start_offset, p->next_insn_offset,
-                    disasm_info, perf.block_latency);
+      if (!intel_shader_dump_filter ||
+          (intel_shader_dump_filter && intel_shader_dump_filter == params->source_hash)) {
+         fprintf(stderr, "Native code for %s (src_hash 0x%08x) (sha1 %s)\n"
+                 "SIMD%d shader: %d instructions. %d loops. %u cycles. "
+                 "%d:%d spills:fills, %u sends, "
+                 "scheduled with mode %s. "
+                 "Promoted %u constants. "
+                 "Non-SSA regs (after NIR): %u. "
+                 "Compacted %d to %d bytes (%.0f%%)\n",
+                 shader_name, params->source_hash, sha1buf,
+                 dispatch_width,
+                 before_size / 16 - nop_count - sync_nop_count,
+                 loop_count, perf.latency,
+                 shader_stats.spill_count,
+                 shader_stats.fill_count,
+                 send_count,
+                 shader_stats.scheduler_mode,
+                 shader_stats.promoted_constants,
+                 shader_stats.non_ssa_registers_after_nir,
+                 before_size, after_size,
+                 100.0f * (before_size - after_size) / before_size);
+         dump_assembly(p->store, start_offset, p->next_insn_offset,
+                       disasm_info, perf.block_latency);
+      }
    }
    ralloc_free(disasm_info);
 #ifndef NDEBUG

src/intel/compiler/brw_nir.c

@@ -25,6 +25,7 @@
 #include "brw_nir.h"
 #include "compiler/glsl_types.h"
 #include "compiler/nir/nir_builder.h"
+#include "dev/intel_debug.h"
 
 /*
  * Returns the minimum number of vec4 (as_vec4 == true) or dvec4 (as_vec4 ==

src/intel/compiler/brw_private.h

@@ -89,6 +89,6 @@ int brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo,
                                        const struct brw_cs_prog_data *prog_data,
                                        const unsigned *sizes);
 
-bool brw_should_print_shader(const nir_shader *shader, uint64_t debug_flag);
+bool brw_should_print_shader(const nir_shader *shader, uint64_t debug_flag, uint32_t source_hash);
 
 #endif // __cplusplus

src/intel/compiler/brw_shader.cpp

@@ -980,7 +980,8 @@ brw_shader::debug_optimizer(const nir_shader *nir,
                             const char *pass_name,
                             int iteration, int pass_num) const
 {
-   if (!brw_should_print_shader(nir, DEBUG_OPTIMIZER))
+   /* source_hash is not readily accessible in this context */
+   if (!brw_should_print_shader(nir, DEBUG_OPTIMIZER, 0))
       return;
 
    char *filename;
@@ -1286,8 +1287,12 @@ brw_shader_phase_update(brw_shader &s, enum brw_shader_phase phase)
    brw_validate(s);
 }
 
-bool brw_should_print_shader(const nir_shader *shader, uint64_t debug_flag)
+bool brw_should_print_shader(const nir_shader *shader, uint64_t debug_flag, uint32_t source_hash)
 {
+   if (intel_shader_dump_filter && intel_shader_dump_filter != source_hash) {
+      return false;
+   }
+
    return INTEL_DEBUG(debug_flag) && (!shader->info.internal || NIR_DEBUG(PRINT_INTERNAL));
 }
 

src/intel/dev/intel_debug.c

@@ -208,6 +208,7 @@ uint64_t intel_debug_batch_frame_stop = -1;
 
 uint32_t intel_debug_bkp_before_draw_count = 0;
 uint32_t intel_debug_bkp_after_draw_count = 0;
+uint32_t intel_shader_dump_filter = 0;
 
 static void
 parse_debug_bitset(const char *env, const struct debug_control_bitset *tbl)
@@ -260,6 +261,9 @@ process_intel_debug_variable_once(void)
    intel_debug_bkp_after_draw_count =
       debug_get_num_option("INTEL_DEBUG_BKP_AFTER_DRAW_COUNT", 0);
 
+   intel_shader_dump_filter =
+      debug_get_num_option("INTEL_SHADER_DUMP_FILTER", 0);
+
    if (!(intel_simd & DEBUG_FS_SIMD))
       intel_simd |=   DEBUG_FS_SIMD;
    if (!(intel_simd & DEBUG_CS_SIMD))

src/intel/dev/intel_debug.h

@@ -136,6 +136,7 @@ extern uint32_t intel_debug_bkp_before_draw_count;
 extern uint32_t intel_debug_bkp_after_draw_count;
 extern uint64_t intel_debug_batch_frame_start;
 extern uint64_t intel_debug_batch_frame_stop;
+extern uint32_t intel_shader_dump_filter;
 
 #define INTEL_SIMD(type, size)        (!!(intel_simd & (DEBUG_ ## type ## _SIMD ## size)))
 

src/intel/vulkan/anv_blorp.c

@@ -87,9 +87,12 @@ upload_blorp_shader(struct blorp_batch *batch, uint32_t stage,
    anv_shader_bin_unref(device, bin);
 
    if (INTEL_DEBUG(DEBUG_SHADERS_LINENO)) {
-      brw_disassemble_with_lineno(&device->physical->compiler->isa,
-                                  stage, -1, 0, kernel, 0,
-                                  bin->kernel.offset, stderr);
+      /* shader hash is zero in this context */
+      if (!intel_shader_dump_filter) {
+         brw_disassemble_with_lineno(&device->physical->compiler->isa,
+                                     stage, -1, 0, kernel, 0,
+                                     bin->kernel.offset, stderr);
+      }
    }
 
    *kernel_out = bin->kernel.offset;

src/intel/vulkan/anv_pipeline.c

@@ -87,9 +87,12 @@ anv_shader_stage_to_nir(struct anv_device *device,
       return NULL;
 
    if (INTEL_DEBUG(intel_debug_flag_for_shader_stage(stage))) {
-      fprintf(stderr, "NIR (from SPIR-V) for %s shader:\n",
-              gl_shader_stage_name(stage));
-      nir_print_shader(nir, stderr);
+      /* src_hash is unknown at the point */
+      if (!intel_shader_dump_filter) {
+         fprintf(stderr, "NIR (from SPIR-V) for %s shader:\n",
+                 gl_shader_stage_name(stage));
+         nir_print_shader(nir, stderr);
+      }
    }
 
    NIR_PASS_V(nir, nir_lower_io_to_temporaries,
@@ -1685,10 +1688,13 @@ anv_pipeline_add_executable(struct anv_pipeline *pipeline,
    }
 
    if (INTEL_DEBUG(DEBUG_SHADERS_LINENO) && stage->code) {
-      brw_disassemble_with_lineno(&pipeline->device->physical->compiler->isa,
-                                  stage->stage, (int)stats->dispatch_width,
-                                  stage->source_hash, stage->code, code_offset,
-                                  stage->bin->kernel.offset, stderr);
+      if (!intel_shader_dump_filter ||
+          (intel_shader_dump_filter && intel_shader_dump_filter == stage->source_hash)) {
+         brw_disassemble_with_lineno(&pipeline->device->physical->compiler->isa,
+                                     stage->stage, (int)stats->dispatch_width,
+                                     stage->source_hash, stage->code, code_offset,
+                                     stage->bin->kernel.offset, stderr);
+      }
    }
 
    const struct anv_pipeline_executable exe = {