From 108bb3d28b94b10fef052476155241278a88bbf9 Mon Sep 17 00:00:00 2001 From: CamilleLaVey Date: Fri, 6 Mar 2026 23:24:18 -0400 Subject: [PATCH] [debug] Added extra logging/ address for shader info -> FP32Mul Optimize Path --- .../backend/spirv/emit_spirv.cpp | 109 +++++++++++++++++- .../spirv/emit_spirv_floating_point.cpp | 45 +++++++- .../backend/spirv/spirv_emit_context.cpp | 39 ++++++- .../backend/spirv/spirv_emit_context.h | 2 + src/shader_recompiler/frontend/ir/program.h | 4 + .../frontend/maxwell/translate_program.cpp | 2 + 6 files changed, 198 insertions(+), 3 deletions(-) diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 313a1deb30..719d7a2744 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project // SPDX-License-Identifier: GPL-3.0-or-later // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project @@ -11,6 +11,7 @@ #include #include +#include "common/logging/log.h" #include "common/settings.h" #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" @@ -20,6 +21,100 @@ namespace Shader::Backend::SPIRV { namespace { +[[nodiscard]] constexpr std::string_view StageName(Stage stage) noexcept { + switch (stage) { + case Stage::VertexA: + return "VertexA"; + case Stage::VertexB: + return "VertexB"; + case Stage::TessellationControl: + return "TessellationControl"; + case Stage::TessellationEval: + return "TessellationEval"; + case Stage::Geometry: + return "Geometry"; + case Stage::Fragment: + return "Fragment"; + case Stage::Compute: + return "Compute"; + } + return "Unknown"; +} + +[[nodiscard]] constexpr std::string_view DenormModeName(bool flush, bool preserve) noexcept { + if (flush && preserve) { + return "Flush+Preserve"; + } + if (flush) { + return "Flush"; + } + if (preserve) { + return "Preserve"; + } + return "None"; +} + +void LogRzBackendSummary(const Profile& profile, const IR::Program& program, bool optimize) { + if (!Settings::values.renderer_debug) { + return; + } + u32 rz_count{}; + for (const IR::Block* const block : program.post_order_blocks) { + for (const IR::Inst& inst : block->Instructions()) { + switch (inst.GetOpcode()) { + case IR::Opcode::FPAdd16: + case IR::Opcode::FPFma16: + case IR::Opcode::FPMul16: + case IR::Opcode::FPRoundEven16: + case IR::Opcode::FPFloor16: + case IR::Opcode::FPCeil16: + case IR::Opcode::FPTrunc16: + case IR::Opcode::FPAdd32: + case IR::Opcode::FPFma32: + case IR::Opcode::FPMul32: + case IR::Opcode::FPRoundEven32: + case IR::Opcode::FPFloor32: + case IR::Opcode::FPCeil32: + case IR::Opcode::FPTrunc32: + case IR::Opcode::FPOrdEqual32: + case IR::Opcode::FPUnordEqual32: + case IR::Opcode::FPOrdNotEqual32: + case IR::Opcode::FPUnordNotEqual32: + case IR::Opcode::FPOrdLessThan32: + case IR::Opcode::FPUnordLessThan32: + case IR::Opcode::FPOrdGreaterThan32: + case IR::Opcode::FPUnordGreaterThan32: + case IR::Opcode::FPOrdLessThanEqual32: + case IR::Opcode::FPUnordLessThanEqual32: + case IR::Opcode::FPOrdGreaterThanEqual32: + case IR::Opcode::FPUnordGreaterThanEqual32: + case IR::Opcode::ConvertF16F32: + case IR::Opcode::ConvertF64F32: + rz_count += inst.Flags().rounding == IR::FpRounding::RZ ? 1U : 0U; + break; + default: + break; + } + } + } + if (rz_count == 0) { + return; + } + + LOG_INFO(Shader_SPIRV, + "SPV_RZ {} start={:#010x} optimize={} support_float_controls={} separate_denorm_behavior={} broken_fp16_float_controls={} fp16_denorm={} fp32_denorm={} signed_nan16={} signed_nan32={} signed_nan64={} rz_inst_count={}", + StageName(program.stage), program.start_address, optimize, + profile.support_float_controls, profile.support_separate_denorm_behavior, + profile.has_broken_fp16_float_controls, + DenormModeName(program.info.uses_fp16_denorms_flush, + program.info.uses_fp16_denorms_preserve), + DenormModeName(program.info.uses_fp32_denorms_flush, + program.info.uses_fp32_denorms_preserve), + profile.support_fp16_signed_zero_nan_preserve, + profile.support_fp32_signed_zero_nan_preserve, + profile.support_fp64_signed_zero_nan_preserve, rz_count); +} + template struct FuncTraits {}; thread_local std::unique_ptr thread_optimizer; @@ -503,6 +598,7 @@ void PatchPhiNodes(IR::Program& program, EmitContext& ctx) { std::vector EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info, IR::Program& program, Bindings& bindings, bool optimize) { + LogRzBackendSummary(profile, program, optimize); EmitContext ctx{profile, runtime_info, program, bindings}; const Id main{DefineMain(ctx, program)}; DefineEntryPoint(program, ctx, main); @@ -516,6 +612,12 @@ std::vector EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_in PatchPhiNodes(program, ctx); if (!optimize) { + if (Settings::values.renderer_debug && ctx.log_rz_fp_controls) { + const std::vector spirv{ctx.Assemble()}; + LOG_INFO(Shader_SPIRV, "SPV_RZ {} start={:#010x} assembled_words={} optimized_words={} validator_run=false", + StageName(program.stage), program.start_address, spirv.size(), spirv.size()); + return spirv; + } return ctx.Assemble(); } else { std::vector spirv = ctx.Assemble(); @@ -535,6 +637,11 @@ std::vector EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_in "Failed to optimize SPIRV shader output, continuing without optimization"); result = std::move(spirv); } + if (Settings::values.renderer_debug && ctx.log_rz_fp_controls) { + LOG_INFO(Shader_SPIRV, + "SPV_RZ {} start={:#010x} assembled_words={} optimized_words={} validator_run=false", + StageName(program.stage), program.start_address, spirv.size(), result.size()); + } return result; } } diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp index d921913b4a..d09b2238e5 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp @@ -1,16 +1,59 @@ -// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project // SPDX-License-Identifier: GPL-3.0-or-later + // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include "common/logging/log.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" #include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" namespace Shader::Backend::SPIRV { namespace { +[[nodiscard]] constexpr std::string_view StageName(Stage stage) noexcept { + switch (stage) { + case Stage::VertexA: + return "VertexA"; + case Stage::VertexB: + return "VertexB"; + case Stage::TessellationControl: + return "TessellationControl"; + case Stage::TessellationEval: + return "TessellationEval"; + case Stage::Geometry: + return "Geometry"; + case Stage::Fragment: + return "Fragment"; + case Stage::Compute: + return "Compute"; + } + return "Unknown"; +} + +[[nodiscard]] constexpr std::string_view FmzName(IR::FmzMode fmz_mode) noexcept { + switch (fmz_mode) { + case IR::FmzMode::DontCare: + return "DontCare"; + case IR::FmzMode::FTZ: + return "FTZ"; + case IR::FmzMode::FMZ: + return "FMZ"; + case IR::FmzMode::None: + return "None"; + } + return "Unknown"; +} + Id Decorate(EmitContext& ctx, IR::Inst* inst, Id op) { const auto flags{inst->Flags()}; + if (Settings::values.renderer_debug && ctx.log_rz_fp_controls && + flags.rounding == IR::FpRounding::RZ) { + LOG_INFO(Shader_SPIRV, + "SPV_RZ_EMIT {} start={:#010x} ir_opcode={} spirv_op=OpFMul result_id={} no_contraction={} fmz={} float_controls_ext={}", + StageName(ctx.stage), ctx.start_address, inst->GetOpcode(), op, + flags.no_contraction, FmzName(flags.fmz_mode), ctx.profile.support_float_controls); + } if (flags.no_contraction) { ctx.Decorate(op, spv::Decoration::NoContraction); } diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp index fb66a7962e..7b0181249c 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp @@ -473,7 +473,44 @@ void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_vie EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_info_, IR::Program& program, Bindings& bindings) : Sirit::Module(profile_.supported_spirv), profile{profile_}, runtime_info{runtime_info_}, - stage{program.stage}, texture_rescaling_index{bindings.texture_scaling_index}, + stage{program.stage}, start_address{program.start_address}, + log_rz_fp_controls{std::ranges::any_of(program.post_order_blocks, [](const IR::Block* block) { + return std::ranges::any_of(block->Instructions(), [](const IR::Inst& inst) { + switch (inst.GetOpcode()) { + case IR::Opcode::FPAdd16: + case IR::Opcode::FPFma16: + case IR::Opcode::FPMul16: + case IR::Opcode::FPRoundEven16: + case IR::Opcode::FPFloor16: + case IR::Opcode::FPCeil16: + case IR::Opcode::FPTrunc16: + case IR::Opcode::FPAdd32: + case IR::Opcode::FPFma32: + case IR::Opcode::FPMul32: + case IR::Opcode::FPRoundEven32: + case IR::Opcode::FPFloor32: + case IR::Opcode::FPCeil32: + case IR::Opcode::FPTrunc32: + case IR::Opcode::FPOrdEqual32: + case IR::Opcode::FPUnordEqual32: + case IR::Opcode::FPOrdNotEqual32: + case IR::Opcode::FPUnordNotEqual32: + case IR::Opcode::FPOrdLessThan32: + case IR::Opcode::FPUnordLessThan32: + case IR::Opcode::FPOrdGreaterThan32: + case IR::Opcode::FPUnordGreaterThan32: + case IR::Opcode::FPOrdLessThanEqual32: + case IR::Opcode::FPUnordLessThanEqual32: + case IR::Opcode::FPOrdGreaterThanEqual32: + case IR::Opcode::FPUnordGreaterThanEqual32: + case IR::Opcode::ConvertF16F32: + case IR::Opcode::ConvertF64F32: + return inst.Flags().rounding == IR::FpRounding::RZ; + default: + return false; + } + }); + })}, texture_rescaling_index{bindings.texture_scaling_index}, image_rescaling_index{bindings.image_scaling_index} { const bool is_unified{profile.unified_descriptor_binding}; u32& uniform_binding{is_unified ? bindings.unified : bindings.uniform_buffer}; diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.h b/src/shader_recompiler/backend/spirv/spirv_emit_context.h index 396022eddf..21151bab38 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.h +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h @@ -216,6 +216,8 @@ public: const Profile& profile; const RuntimeInfo& runtime_info; Stage stage{}; + u32 start_address{}; + bool log_rz_fp_controls{}; Id void_id{}; Id U1{}; diff --git a/src/shader_recompiler/frontend/ir/program.h b/src/shader_recompiler/frontend/ir/program.h index 6b4a05c598..1836a18bd3 100644 --- a/src/shader_recompiler/frontend/ir/program.h +++ b/src/shader_recompiler/frontend/ir/program.h @@ -1,3 +1,6 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later + // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later @@ -20,6 +23,7 @@ struct Program { BlockList post_order_blocks; Info info; Stage stage{}; + u32 start_address{}; std::array workgroup_size{}; OutputTopology output_topology{}; u32 output_vertices{}; diff --git a/src/shader_recompiler/frontend/maxwell/translate_program.cpp b/src/shader_recompiler/frontend/maxwell/translate_program.cpp index f52a3e72de..6cca023330 100644 --- a/src/shader_recompiler/frontend/maxwell/translate_program.cpp +++ b/src/shader_recompiler/frontend/maxwell/translate_program.cpp @@ -458,6 +458,7 @@ IR::Program TranslateProgram(ObjectPool& inst_pool, ObjectPool