mirror of
https://git.eden-emu.dev/eden-emu/eden
synced 2026-04-10 03:18:55 +02:00
Revert "[debug] Added extra logging/ address for shader info -> FP32Mul Optimize Path"
This commit is contained in:
parent
124c97a88a
commit
73c16c3d45
6 changed files with 1 additions and 241 deletions
|
|
@ -11,7 +11,6 @@
|
|||
#include <vector>
|
||||
#include <spirv-tools/optimizer.hpp>
|
||||
|
||||
#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"
|
||||
|
|
@ -22,148 +21,6 @@
|
|||
|
||||
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";
|
||||
}
|
||||
|
||||
[[nodiscard]] constexpr bool IsFp32RoundingRelevantOpcode(IR::Opcode opcode) noexcept {
|
||||
switch (opcode) {
|
||||
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 true;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
struct Fp32RoundingUsage {
|
||||
u32 rz_count{};
|
||||
bool has_conflicting_rounding{};
|
||||
};
|
||||
|
||||
Fp32RoundingUsage CollectFp32RoundingUsage(const IR::Program& program) {
|
||||
Fp32RoundingUsage usage{};
|
||||
for (const IR::Block* const block : program.post_order_blocks) {
|
||||
for (const IR::Inst& inst : block->Instructions()) {
|
||||
if (!IsFp32RoundingRelevantOpcode(inst.GetOpcode())) {
|
||||
continue;
|
||||
}
|
||||
switch (inst.Flags<IR::FpControl>().rounding) {
|
||||
case IR::FpRounding::RZ:
|
||||
++usage.rz_count;
|
||||
break;
|
||||
case IR::FpRounding::RN:
|
||||
case IR::FpRounding::RM:
|
||||
case IR::FpRounding::RP:
|
||||
usage.has_conflicting_rounding = true;
|
||||
break;
|
||||
case IR::FpRounding::DontCare:
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
return usage;
|
||||
}
|
||||
|
||||
void LogRzBackendSummary(const Profile& profile, const IR::Program& program, bool optimize) {
|
||||
if (!Settings::values.renderer_debug) {
|
||||
return;
|
||||
}
|
||||
const Fp32RoundingUsage usage{CollectFp32RoundingUsage(program)};
|
||||
if (usage.rz_count == 0) {
|
||||
return;
|
||||
}
|
||||
|
||||
LOG_INFO(Shader_SPIRV,
|
||||
"SPV_RZ {} start={:#010x} optimize={} support_float_controls={} separate_denorm_behavior={} separate_rounding_mode={} support_fp32_rounding_rtz={} broken_fp16_float_controls={} fp16_denorm={} fp32_denorm={} signed_nan16={} signed_nan32={} signed_nan64={} rz_inst_count={} mixed_fp32_rounding={}",
|
||||
StageName(program.stage), program.start_address, optimize,
|
||||
profile.support_float_controls, profile.support_separate_denorm_behavior,
|
||||
profile.support_separate_rounding_mode, profile.support_fp32_rounding_rtz,
|
||||
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, usage.rz_count,
|
||||
usage.has_conflicting_rounding);
|
||||
}
|
||||
|
||||
void SetupRoundingControl(const Profile& profile, const IR::Program& program, EmitContext& ctx,
|
||||
Id main_func) {
|
||||
const Fp32RoundingUsage usage{CollectFp32RoundingUsage(program)};
|
||||
if (usage.rz_count == 0) {
|
||||
return;
|
||||
}
|
||||
if (usage.has_conflicting_rounding) {
|
||||
if (Settings::values.renderer_debug) {
|
||||
LOG_INFO(Shader_SPIRV,
|
||||
"SPV_RZ {} start={:#010x} skipping_fp32_rtz_execution_mode reason=mixed_rounding",
|
||||
StageName(program.stage), program.start_address);
|
||||
}
|
||||
return;
|
||||
}
|
||||
if (!profile.support_fp32_rounding_rtz) {
|
||||
if (Settings::values.renderer_debug) {
|
||||
LOG_INFO(Shader_SPIRV,
|
||||
"SPV_RZ {} start={:#010x} skipping_fp32_rtz_execution_mode reason=unsupported_fp32_rtz",
|
||||
StageName(program.stage), program.start_address);
|
||||
}
|
||||
return;
|
||||
}
|
||||
ctx.AddCapability(spv::Capability::RoundingModeRTZ);
|
||||
ctx.AddExecutionMode(main_func, spv::ExecutionMode::RoundingModeRTZ, 32U);
|
||||
}
|
||||
|
||||
template <class Func>
|
||||
struct FuncTraits {};
|
||||
thread_local std::unique_ptr<spvtools::Optimizer> thread_optimizer;
|
||||
|
|
@ -647,7 +504,6 @@ void PatchPhiNodes(IR::Program& program, EmitContext& ctx) {
|
|||
|
||||
std::vector<u32> 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);
|
||||
|
|
@ -662,12 +518,6 @@ std::vector<u32> 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<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();
|
||||
} else {
|
||||
std::vector<u32> spirv = ctx.Assemble();
|
||||
|
|
@ -687,11 +537,6 @@ std::vector<u32> 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;
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,57 +4,14 @@
|
|||
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
|
||||
// SPDX-License-Identifier: GPL-2.0-or-later
|
||||
|
||||
#include "common/logging/log.h"
|
||||
#include "common/settings.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<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 no_contraction={} fmz={} float_controls_ext={}",
|
||||
StageName(ctx.stage), ctx.start_address, inst->GetOpcode(),
|
||||
flags.no_contraction, FmzName(flags.fmz_mode), ctx.profile.support_float_controls);
|
||||
}
|
||||
if (flags.no_contraction) {
|
||||
ctx.Decorate(op, spv::Decoration::NoContraction);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -474,44 +474,7 @@ 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}, 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},
|
||||
stage{program.stage}, 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};
|
||||
|
|
|
|||
|
|
@ -216,8 +216,6 @@ public:
|
|||
const Profile& profile;
|
||||
const RuntimeInfo& runtime_info;
|
||||
Stage stage{};
|
||||
u32 start_address{};
|
||||
bool log_rz_fp_controls{};
|
||||
|
||||
Id void_id{};
|
||||
Id U1{};
|
||||
|
|
|
|||
|
|
@ -23,7 +23,6 @@ struct Program {
|
|||
BlockList post_order_blocks;
|
||||
Info info;
|
||||
Stage stage{};
|
||||
u32 start_address{};
|
||||
std::array<u32, 3> workgroup_size{};
|
||||
OutputTopology output_topology{};
|
||||
u32 output_vertices{};
|
||||
|
|
|
|||
|
|
@ -247,7 +247,6 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
|
|||
program.blocks = GenerateBlocks(program.syntax_list);
|
||||
program.post_order_blocks = PostOrder(program.syntax_list.front());
|
||||
program.stage = env.ShaderStage();
|
||||
program.start_address = env.StartAddress();
|
||||
program.local_memory_size = env.LocalMemorySize();
|
||||
switch (program.stage) {
|
||||
case Stage::TessellationControl: {
|
||||
|
|
@ -339,7 +338,6 @@ IR::Program MergeDualVertexPrograms(IR::Program& vertex_a, IR::Program& vertex_b
|
|||
result.post_order_blocks.push_back(block);
|
||||
}
|
||||
result.stage = Stage::VertexB;
|
||||
result.start_address = env_vertex_b.StartAddress();
|
||||
result.info = vertex_a.info;
|
||||
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;
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue