eden-emu
/
eden
mirror of https://git.eden-emu.dev/eden-emu/eden.git
1
0
Fork 0
Code Releases Wiki Activity
Browse Source

[debug] Added extra logging/ address for shader info -> FP32Mul Optimize Path

lizzie/vkexperiments1-highp-fucked
CamilleLaVey 6 days ago
committed by lizzie
parent
5b12a7725d
commit
108bb3d28b
6 changed files with 198 additions and 3 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 109
      src/shader_recompiler/backend/spirv/emit_spirv.cpp
  2. 45
      src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp
  3. 39
      src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
  4. 2
      src/shader_recompiler/backend/spirv/spirv_emit_context.h
  5. 4
      src/shader_recompiler/frontend/ir/program.h
  6. 2
      src/shader_recompiler/frontend/maxwell/translate_program.cpp

109
src/shader_recompiler/backend/spirv/emit_spirv.cpp
View File

@ -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-License-Identifier: GPL-3.0-or-later
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
@ -11,6 +11,7 @@
#include <vector> #include <vector>
#include <spirv-tools/optimizer.hpp> #include <spirv-tools/optimizer.hpp>
#include "common/logging/log.h"
#include "common/settings.h" #include "common/settings.h"
#include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv.h"
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
@ -20,6 +21,100 @@
namespace Shader::Backend::SPIRV { namespace Shader::Backend::SPIRV {
namespace { 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<IR::FpControl>().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 <class Func> template <class Func>
struct FuncTraits {}; struct FuncTraits {};
thread_local std::unique_ptr<spvtools::Optimizer> thread_optimizer; thread_local std::unique_ptr<spvtools::Optimizer> thread_optimizer;
@ -503,6 +598,7 @@ void PatchPhiNodes(IR::Program& program, EmitContext& ctx) {
std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info, std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info,
IR::Program& program, Bindings& bindings, bool optimize) { IR::Program& program, Bindings& bindings, bool optimize) {
LogRzBackendSummary(profile, program, optimize);
EmitContext ctx{profile, runtime_info, program, bindings}; EmitContext ctx{profile, runtime_info, program, bindings};
const Id main{DefineMain(ctx, program)}; const Id main{DefineMain(ctx, program)};
DefineEntryPoint(program, ctx, main); DefineEntryPoint(program, ctx, main);
@ -516,6 +612,12 @@ std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_in
PatchPhiNodes(program, ctx); PatchPhiNodes(program, ctx);
if (!optimize) { if (!optimize) {
if (Settings::values.renderer_debug && ctx.log_rz_fp_controls) {
const std::vector<u32> 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(); return ctx.Assemble();
} else { } else {
std::vector<u32> spirv = ctx.Assemble(); std::vector<u32> spirv = ctx.Assemble();
@ -535,6 +637,11 @@ std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_in
"Failed to optimize SPIRV shader output, continuing without optimization"); "Failed to optimize SPIRV shader output, continuing without optimization");
result = std::move(spirv); 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; return result;
} }
} }

45
src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp
View File

@ -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-License-Identifier: GPL-3.0-or-later
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later // 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/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/backend/spirv/spirv_emit_context.h"
#include "shader_recompiler/frontend/ir/modifiers.h" #include "shader_recompiler/frontend/ir/modifiers.h"
namespace Shader::Backend::SPIRV { namespace Shader::Backend::SPIRV {
namespace { 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) { Id Decorate(EmitContext& ctx, IR::Inst* inst, Id op) {
const auto flags{inst->Flags<IR::FpControl>()}; const auto flags{inst->Flags<IR::FpControl>()};
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) { if (flags.no_contraction) {
ctx.Decorate(op, spv::Decoration::NoContraction); ctx.Decorate(op, spv::Decoration::NoContraction);
} }

39
src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
View File

@ -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_, EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_info_,
IR::Program& program, Bindings& bindings) IR::Program& program, Bindings& bindings)
: Sirit::Module(profile_.supported_spirv), profile{profile_}, runtime_info{runtime_info_}, : 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<IR::FpControl>().rounding == IR::FpRounding::RZ;
default:
return false;
}
});
})}, texture_rescaling_index{bindings.texture_scaling_index},
image_rescaling_index{bindings.image_scaling_index} { image_rescaling_index{bindings.image_scaling_index} {
const bool is_unified{profile.unified_descriptor_binding}; const bool is_unified{profile.unified_descriptor_binding};
u32& uniform_binding{is_unified ? bindings.unified : bindings.uniform_buffer}; u32& uniform_binding{is_unified ? bindings.unified : bindings.uniform_buffer};

2
src/shader_recompiler/backend/spirv/spirv_emit_context.h
View File

@ -216,6 +216,8 @@ public:
const Profile& profile; const Profile& profile;
const RuntimeInfo& runtime_info; const RuntimeInfo& runtime_info;
Stage stage{}; Stage stage{};
u32 start_address{};
bool log_rz_fp_controls{};
Id void_id{}; Id void_id{};
Id U1{}; Id U1{};

4
src/shader_recompiler/frontend/ir/program.h
View File

@ -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-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later // SPDX-License-Identifier: GPL-2.0-or-later
@ -20,6 +23,7 @@ struct Program {
BlockList post_order_blocks; BlockList post_order_blocks;
Info info; Info info;
Stage stage{}; Stage stage{};
u32 start_address{};
std::array<u32, 3> workgroup_size{}; std::array<u32, 3> workgroup_size{};
OutputTopology output_topology{}; OutputTopology output_topology{};
u32 output_vertices{}; u32 output_vertices{};

2
src/shader_recompiler/frontend/maxwell/translate_program.cpp
View File

@ -458,6 +458,7 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
program.blocks = GenerateBlocks(program.syntax_list); program.blocks = GenerateBlocks(program.syntax_list);
program.post_order_blocks = PostOrder(program.syntax_list.front()); program.post_order_blocks = PostOrder(program.syntax_list.front());
program.stage = env.ShaderStage(); program.stage = env.ShaderStage();
program.start_address = env.StartAddress();
program.local_memory_size = env.LocalMemorySize(); program.local_memory_size = env.LocalMemorySize();
switch (program.stage) { switch (program.stage) {
case Stage::TessellationControl: { case Stage::TessellationControl: {
@ -554,6 +555,7 @@ IR::Program MergeDualVertexPrograms(IR::Program& vertex_a, IR::Program& vertex_b
result.post_order_blocks.push_back(block); result.post_order_blocks.push_back(block);
} }
result.stage = Stage::VertexB; result.stage = Stage::VertexB;
result.start_address = env_vertex_b.StartAddress();
result.info = vertex_a.info; result.info = vertex_a.info;
result.local_memory_size = (std::max)(vertex_a.local_memory_size, vertex_b.local_memory_size); result.local_memory_size = (std::max)(vertex_a.local_memory_size, vertex_b.local_memory_size);
result.info.loads.mask |= vertex_b.info.loads.mask; result.info.loads.mask |= vertex_b.info.loads.mask;

Write Preview
Loading…
Cancel
Save