From e2bc05b17d91854cbb9c0ce3647141bf7d33143e Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sat, 20 Feb 2021 03:30:13 -0300
Subject: shader: Add denorm flush support
---
src/shader_recompiler/backend/spirv/emit_spirv.cpp | 63 ++++++++++++++++++++--
src/shader_recompiler/backend/spirv/emit_spirv.h | 4 +-
.../backend/spirv/emit_spirv_floating_point.cpp | 6 +--
3 files changed, 64 insertions(+), 9 deletions(-)
(limited to 'src/shader_recompiler/backend')
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
index 4ce07c2814..2519e446ae 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
@@ -14,8 +14,6 @@
#include "shader_recompiler/frontend/ir/microinstruction.h"
#include "shader_recompiler/frontend/ir/program.h"
-#pragma optimize("", off)
-
namespace Shader::Backend::SPIRV {
namespace {
template <class Func>
@@ -113,9 +111,61 @@ Id TypeId(const EmitContext& ctx, IR::Type type) {
throw NotImplementedException("Phi node type {}", type);
}
}
+
+void SetupDenormControl(const Profile& profile, const IR::Program& program, EmitContext& ctx,
+ Id main_func) {
+ if (!profile.support_float_controls) {
+ return;
+ }
+ const Info& info{program.info};
+ if (!info.uses_fp32_denorms_flush && !info.uses_fp32_denorms_preserve &&
+ !info.uses_fp16_denorms_flush && !info.uses_fp16_denorms_preserve) {
+ return;
+ }
+ ctx.AddExtension("SPV_KHR_float_controls");
+
+ if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) {
+ // LOG_ERROR(HW_GPU, "Fp32 denorm flush and preserve on the same shader");
+ } else if (info.uses_fp32_denorms_flush) {
+ if (profile.support_fp32_denorm_flush) {
+ ctx.AddCapability(spv::Capability::DenormFlushToZero);
+ ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormFlushToZero, 32U);
+ } else {
+ // Drivers will most likely flush denorms by default, no need to warn
+ }
+ } else if (info.uses_fp32_denorms_preserve) {
+ if (profile.support_fp32_denorm_preserve) {
+ ctx.AddCapability(spv::Capability::DenormPreserve);
+ ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 32U);
+ } else {
+ // LOG_WARNING(HW_GPU, "Fp32 denorm preserve used in shader without host support");
+ }
+ }
+ if (!profile.support_separate_denorm_behavior) {
+ // No separate denorm behavior
+ return;
+ }
+ if (info.uses_fp16_denorms_flush && info.uses_fp16_denorms_preserve) {
+ // LOG_ERROR(HW_GPU, "Fp16 denorm flush and preserve on the same shader");
+ } else if (info.uses_fp16_denorms_flush) {
+ if (profile.support_fp16_denorm_flush) {
+ ctx.AddCapability(spv::Capability::DenormFlushToZero);
+ ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U);
+ } else {
+ // Same as fp32, no need to warn as most drivers will flush by default
+ }
+ } else if (info.uses_fp32_denorms_preserve) {
+ if (profile.support_fp16_denorm_preserve) {
+ ctx.AddCapability(spv::Capability::DenormPreserve);
+ ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U);
+ } else {
+ // LOG_WARNING(HW_GPU, "Fp16 denorm preserve used in shader without host support");
+ }
+ }
+}
} // Anonymous namespace
-std::vector<u32> EmitSPIRV(Environment& env, IR::Program& program) {
+std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) {
EmitContext ctx{program};
const Id void_function{ctx.TypeFunction(ctx.void_id)};
// FIXME: Forward declare functions (needs sirit support)
@@ -131,10 +181,11 @@ std::vector<u32> EmitSPIRV(Environment& env, IR::Program& program) {
ctx.OpFunctionEnd();
}
boost::container::small_vector<Id, 32> interfaces;
- if (program.info.uses_workgroup_id) {
+ const Info& info{program.info};
+ if (info.uses_workgroup_id) {
interfaces.push_back(ctx.workgroup_id);
}
- if (program.info.uses_local_invocation_id) {
+ if (info.uses_local_invocation_id) {
interfaces.push_back(ctx.local_invocation_id);
}
const std::span interfaces_span(interfaces.data(), interfaces.size());
@@ -144,6 +195,8 @@ std::vector<u32> EmitSPIRV(Environment& env, IR::Program& program) {
ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1],
workgroup_size[2]);
+ SetupDenormControl(profile, program, ctx, func);
+
return ctx.Assemble();
}
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.h b/src/shader_recompiler/backend/spirv/emit_spirv.h
index 2b59c0b726..de624a1510 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv.h
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.h
@@ -11,10 +11,12 @@
#include "shader_recompiler/environment.h"
#include "shader_recompiler/frontend/ir/microinstruction.h"
#include "shader_recompiler/frontend/ir/program.h"
+#include "shader_recompiler/profile.h"
namespace Shader::Backend::SPIRV {
-[[nodiscard]] std::vector<u32> EmitSPIRV(Environment& env, IR::Program& program);
+[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env,
+ IR::Program& program);
// Microinstruction emitters
Id EmitPhi(EmitContext& ctx, IR::Inst* inst);
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 9ef1805310..c9687de377 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp
@@ -13,7 +13,10 @@ Id Decorate(EmitContext& ctx, IR::Inst* inst, Id op) {
ctx.Decorate(op, spv::Decoration::NoContraction);
}
switch (flags.rounding) {
+ case IR::FpRounding::DontCare:
+ break;
case IR::FpRounding::RN:
+ ctx.Decorate(op, spv::Decoration::FPRoundingMode, spv::FPRoundingMode::RTE);
break;
case IR::FpRounding::RM:
ctx.Decorate(op, spv::Decoration::FPRoundingMode, spv::FPRoundingMode::RTN);
@@ -25,9 +28,6 @@ Id Decorate(EmitContext& ctx, IR::Inst* inst, Id op) {
ctx.Decorate(op, spv::Decoration::FPRoundingMode, spv::FPRoundingMode::RTZ);
break;
}
- if (flags.fmz_mode != IR::FmzMode::FTZ) {
- throw NotImplementedException("Denorm management not implemented");
- }
return op;
}
--
cgit v1.2.3-70-g09d2