diff --git a/docs/envvars.rst b/docs/envvars.rst index a1b437123be..0923b79cef5 100644 --- a/docs/envvars.rst +++ b/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: diff --git a/src/gallium/drivers/iris/iris_program_cache.c b/src/gallium/drivers/iris/iris_program_cache.c index 9522830f27d..a9fddd6d74b 100644 --- a/src/gallium/drivers/iris/iris_program_cache.c +++ b/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); + } } } } diff --git a/src/intel/compiler/brw_compile_bs.cpp b/src/intel/compiler/brw_compile_bs.cpp index f90141b4e01..4b82398a2e6 100644 --- a/src/intel/compiler/brw_compile_bs.cpp +++ b/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, ¶ms->base); diff --git a/src/intel/compiler/brw_compile_cs.cpp b/src/intel/compiler/brw_compile_cs.cpp index ab1356d61a4..92fd6f48b22 100644 --- a/src/intel/compiler/brw_compile_cs.cpp +++ b/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, ¶ms->base); prog_data->uses_inline_data = brw_nir_uses_inline_data(nir) || diff --git a/src/intel/compiler/brw_compile_fs.cpp b/src/intel/compiler/brw_compile_fs.cpp index 19e681db4b4..3bf285ae6ab 100644 --- a/src/intel/compiler/brw_compile_fs.cpp +++ b/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, ¶ms->base); diff --git a/src/intel/compiler/brw_compile_gs.cpp b/src/intel/compiler/brw_compile_gs.cpp index b0c0b11a2a3..a49fe60c140 100644 --- a/src/intel/compiler/brw_compile_gs.cpp +++ b/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, ¶ms->base); diff --git a/src/intel/compiler/brw_compile_mesh.cpp b/src/intel/compiler/brw_compile_mesh.cpp index 5939e3a42e0..2f47fecc954 100644 --- a/src/intel/compiler/brw_compile_mesh.cpp +++ b/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, ¶ms->base); diff --git a/src/intel/compiler/brw_compile_tcs.cpp b/src/intel/compiler/brw_compile_tcs.cpp index 96a51bb97e1..10163a37833 100644 --- a/src/intel/compiler/brw_compile_tcs.cpp +++ b/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, ¶ms->base); diff --git a/src/intel/compiler/brw_compile_tes.cpp b/src/intel/compiler/brw_compile_tes.cpp index d6f13baae0d..88f29a84a8a 100644 --- a/src/intel/compiler/brw_compile_tes.cpp +++ b/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, ¶ms->base); diff --git a/src/intel/compiler/brw_compile_vs.cpp b/src/intel/compiler/brw_compile_vs.cpp index f9d180e15ea..43cf6577938 100644 --- a/src/intel/compiler/brw_compile_vs.cpp +++ b/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 diff --git a/src/intel/compiler/brw_generator.cpp b/src/intel/compiler/brw_generator.cpp index b66eaf1f8c8..e160b40ba48 100644 --- a/src/intel/compiler/brw_generator.cpp +++ b/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 diff --git a/src/intel/compiler/brw_nir.c b/src/intel/compiler/brw_nir.c index bf0b2896aa0..4b853766517 100644 --- a/src/intel/compiler/brw_nir.c +++ b/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 == diff --git a/src/intel/compiler/brw_private.h b/src/intel/compiler/brw_private.h index 89a857c9b72..87cb800e4a8 100644 --- a/src/intel/compiler/brw_private.h +++ b/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 diff --git a/src/intel/compiler/brw_shader.cpp b/src/intel/compiler/brw_shader.cpp index 3440d2874fa..1100492c4c3 100644 --- a/src/intel/compiler/brw_shader.cpp +++ b/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)); } diff --git a/src/intel/dev/intel_debug.c b/src/intel/dev/intel_debug.c index e10759dfdbd..26b8a35f393 100644 --- a/src/intel/dev/intel_debug.c +++ b/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)) diff --git a/src/intel/dev/intel_debug.h b/src/intel/dev/intel_debug.h index b861dfe51b1..e22c646fb6c 100644 --- a/src/intel/dev/intel_debug.h +++ b/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))) diff --git a/src/intel/vulkan/anv_blorp.c b/src/intel/vulkan/anv_blorp.c index b35f56c1703..06edff08049 100644 --- a/src/intel/vulkan/anv_blorp.c +++ b/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; diff --git a/src/intel/vulkan/anv_pipeline.c b/src/intel/vulkan/anv_pipeline.c index a1b537e33f0..49c0a835457 100644 --- a/src/intel/vulkan/anv_pipeline.c +++ b/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 = {