summaryrefslogtreecommitdiffstats
path: root/src/shader_recompiler
diff options
context:
space:
mode:
Diffstat (limited to 'src/shader_recompiler')
-rw-r--r--src/shader_recompiler/CMakeLists.txt5
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv.cpp63
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv.h4
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp6
-rw-r--r--src/shader_recompiler/frontend/ir/ir_emitter.cpp32
-rw-r--r--src/shader_recompiler/frontend/ir/ir_emitter.h8
-rw-r--r--src/shader_recompiler/frontend/ir/modifiers.h23
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/floating_point_conversion_integer.cpp19
-rw-r--r--src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp71
-rw-r--r--src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp1
-rw-r--r--src/shader_recompiler/main.cpp13
-rw-r--r--src/shader_recompiler/profile.h9
-rw-r--r--src/shader_recompiler/recompiler.cpp5
-rw-r--r--src/shader_recompiler/recompiler.h4
-rw-r--r--src/shader_recompiler/shader_info.h7
15 files changed, 210 insertions, 60 deletions
diff --git a/src/shader_recompiler/CMakeLists.txt b/src/shader_recompiler/CMakeLists.txt
index 6047f3ebe..fbd4ec6dc 100644
--- a/src/shader_recompiler/CMakeLists.txt
+++ b/src/shader_recompiler/CMakeLists.txt
@@ -32,6 +32,7 @@ add_library(shader_recompiler STATIC
frontend/ir/ir_emitter.h
frontend/ir/microinstruction.cpp
frontend/ir/microinstruction.h
+ frontend/ir/modifiers.h
frontend/ir/opcodes.cpp
frontend/ir/opcodes.h
frontend/ir/opcodes.inc
@@ -94,9 +95,7 @@ add_library(shader_recompiler STATIC
shader_info.h
)
-target_include_directories(shader_recompiler PRIVATE sirit)
-target_link_libraries(shader_recompiler PRIVATE fmt::fmt sirit)
-target_link_libraries(shader_recompiler INTERFACE fmt::fmt sirit)
+target_link_libraries(shader_recompiler PUBLIC fmt::fmt sirit)
add_executable(shader_util main.cpp)
target_link_libraries(shader_util PRIVATE shader_recompiler)
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
index 4ce07c281..2519e446a 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 2b59c0b72..de624a151 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 9ef180531..c9687de37 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;
}
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.cpp b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
index 559ab9cca..8f120a2f6 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.cpp
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
@@ -558,53 +558,53 @@ F16F32F64 IREmitter::FPSaturate(const F16F32F64& value) {
}
}
-F16F32F64 IREmitter::FPRoundEven(const F16F32F64& value) {
+F16F32F64 IREmitter::FPRoundEven(const F16F32F64& value, FpControl control) {
switch (value.Type()) {
case Type::F16:
- return Inst<F16>(Opcode::FPRoundEven16, value);
+ return Inst<F16>(Opcode::FPRoundEven16, Flags{control}, value);
case Type::F32:
- return Inst<F32>(Opcode::FPRoundEven32, value);
+ return Inst<F32>(Opcode::FPRoundEven32, Flags{control}, value);
case Type::F64:
- return Inst<F64>(Opcode::FPRoundEven64, value);
+ return Inst<F64>(Opcode::FPRoundEven64, Flags{control}, value);
default:
ThrowInvalidType(value.Type());
}
}
-F16F32F64 IREmitter::FPFloor(const F16F32F64& value) {
+F16F32F64 IREmitter::FPFloor(const F16F32F64& value, FpControl control) {
switch (value.Type()) {
case Type::F16:
- return Inst<F16>(Opcode::FPFloor16, value);
+ return Inst<F16>(Opcode::FPFloor16, Flags{control}, value);
case Type::F32:
- return Inst<F32>(Opcode::FPFloor32, value);
+ return Inst<F32>(Opcode::FPFloor32, Flags{control}, value);
case Type::F64:
- return Inst<F64>(Opcode::FPFloor64, value);
+ return Inst<F64>(Opcode::FPFloor64, Flags{control}, value);
default:
ThrowInvalidType(value.Type());
}
}
-F16F32F64 IREmitter::FPCeil(const F16F32F64& value) {
+F16F32F64 IREmitter::FPCeil(const F16F32F64& value, FpControl control) {
switch (value.Type()) {
case Type::F16:
- return Inst<F16>(Opcode::FPCeil16, value);
+ return Inst<F16>(Opcode::FPCeil16, Flags{control}, value);
case Type::F32:
- return Inst<F32>(Opcode::FPCeil32, value);
+ return Inst<F32>(Opcode::FPCeil32, Flags{control}, value);
case Type::F64:
- return Inst<F64>(Opcode::FPCeil64, value);
+ return Inst<F64>(Opcode::FPCeil64, Flags{control}, value);
default:
ThrowInvalidType(value.Type());
}
}
-F16F32F64 IREmitter::FPTrunc(const F16F32F64& value) {
+F16F32F64 IREmitter::FPTrunc(const F16F32F64& value, FpControl control) {
switch (value.Type()) {
case Type::F16:
- return Inst<F16>(Opcode::FPTrunc16, value);
+ return Inst<F16>(Opcode::FPTrunc16, Flags{control}, value);
case Type::F32:
- return Inst<F32>(Opcode::FPTrunc32, value);
+ return Inst<F32>(Opcode::FPTrunc32, Flags{control}, value);
case Type::F64:
- return Inst<F64>(Opcode::FPTrunc64, value);
+ return Inst<F64>(Opcode::FPTrunc64, Flags{control}, value);
default:
ThrowInvalidType(value.Type());
}
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.h b/src/shader_recompiler/frontend/ir/ir_emitter.h
index 24b012a39..959f4f9da 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.h
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.h
@@ -129,10 +129,10 @@ public:
[[nodiscard]] F32 FPSinNotReduced(const F32& value);
[[nodiscard]] F32 FPSqrt(const F32& value);
[[nodiscard]] F16F32F64 FPSaturate(const F16F32F64& value);
- [[nodiscard]] F16F32F64 FPRoundEven(const F16F32F64& value);
- [[nodiscard]] F16F32F64 FPFloor(const F16F32F64& value);
- [[nodiscard]] F16F32F64 FPCeil(const F16F32F64& value);
- [[nodiscard]] F16F32F64 FPTrunc(const F16F32F64& value);
+ [[nodiscard]] F16F32F64 FPRoundEven(const F16F32F64& value, FpControl control = {});
+ [[nodiscard]] F16F32F64 FPFloor(const F16F32F64& value, FpControl control = {});
+ [[nodiscard]] F16F32F64 FPCeil(const F16F32F64& value, FpControl control = {});
+ [[nodiscard]] F16F32F64 FPTrunc(const F16F32F64& value, FpControl control = {});
[[nodiscard]] U32U64 IAdd(const U32U64& a, const U32U64& b);
[[nodiscard]] U32U64 ISub(const U32U64& a, const U32U64& b);
diff --git a/src/shader_recompiler/frontend/ir/modifiers.h b/src/shader_recompiler/frontend/ir/modifiers.h
index c288eede0..44652eae7 100644
--- a/src/shader_recompiler/frontend/ir/modifiers.h
+++ b/src/shader_recompiler/frontend/ir/modifiers.h
@@ -4,25 +4,30 @@
#pragma once
+#include "common/common_types.h"
+
namespace Shader::IR {
enum class FmzMode : u8 {
- None, // Denorms are not flushed, NAN is propagated (nouveau)
- FTZ, // Flush denorms to zero, NAN is propagated (D3D11, NVN, GL, VK)
- FMZ, // Flush denorms to zero, x * 0 == 0 (D3D9)
+ DontCare, // Not specified for this instruction
+ FTZ, // Flush denorms to zero, NAN is propagated (D3D11, NVN, GL, VK)
+ FMZ, // Flush denorms to zero, x * 0 == 0 (D3D9)
+ None, // Denorms are not flushed, NAN is propagated (nouveau)
};
enum class FpRounding : u8 {
- RN, // Round to nearest even,
- RM, // Round towards negative infinity
- RP, // Round towards positive infinity
- RZ, // Round towards zero
+ DontCare, // Not specified for this instruction
+ RN, // Round to nearest even,
+ RM, // Round towards negative infinity
+ RP, // Round towards positive infinity
+ RZ, // Round towards zero
};
struct FpControl {
bool no_contraction{false};
- FpRounding rounding{FpRounding::RN};
- FmzMode fmz_mode{FmzMode::FTZ};
+ FpRounding rounding{FpRounding::DontCare};
+ FmzMode fmz_mode{FmzMode::DontCare};
};
static_assert(sizeof(FpControl) <= sizeof(u32));
+
} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/floating_point_conversion_integer.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/floating_point_conversion_integer.cpp
index ae2d37405..4d82a0009 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/floating_point_conversion_integer.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/floating_point_conversion_integer.cpp
@@ -81,17 +81,28 @@ void TranslateF2I(TranslatorVisitor& v, u64 insn, const IR::F16F32F64& src_a) {
// F2I is used to convert from a floating point value to an integer
const F2I f2i{insn};
+ const bool denorm_cares{f2i.src_format != SrcFormat::F16 && f2i.src_format != SrcFormat::F64 &&
+ f2i.dest_format != DestFormat::I64};
+ IR::FmzMode fmz_mode{IR::FmzMode::DontCare};
+ if (denorm_cares) {
+ fmz_mode = f2i.ftz != 0 ? IR::FmzMode::FTZ : IR::FmzMode::None;
+ }
+ const IR::FpControl fp_control{
+ .no_contraction{true},
+ .rounding{IR::FpRounding::DontCare},
+ .fmz_mode{fmz_mode},
+ };
const IR::F16F32F64 op_a{v.ir.FPAbsNeg(src_a, f2i.abs != 0, f2i.neg != 0)};
const IR::F16F32F64 rounded_value{[&] {
switch (f2i.rounding) {
case Rounding::Round:
- return v.ir.FPRoundEven(op_a);
+ return v.ir.FPRoundEven(op_a, fp_control);
case Rounding::Floor:
- return v.ir.FPFloor(op_a);
+ return v.ir.FPFloor(op_a, fp_control);
case Rounding::Ceil:
- return v.ir.FPCeil(op_a);
+ return v.ir.FPCeil(op_a, fp_control);
case Rounding::Trunc:
- return v.ir.FPTrunc(op_a);
+ return v.ir.FPTrunc(op_a, fp_control);
default:
throw NotImplementedException("Invalid F2I rounding {}", f2i.rounding.Value());
}
diff --git a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
index f7f102f53..6662ef4cd 100644
--- a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
+++ b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
@@ -2,23 +2,28 @@
// Licensed under GPLv2 or any later version
// Refer to the license.txt file included.
+#include "shader_recompiler/frontend/ir/microinstruction.h"
+#include "shader_recompiler/frontend/ir/modifiers.h"
#include "shader_recompiler/frontend/ir/program.h"
#include "shader_recompiler/shader_info.h"
namespace Shader::Optimization {
namespace {
-void AddConstantBufferDescriptor(Info& info, u32 index) {
- auto& descriptor{info.constant_buffers.at(index)};
- if (descriptor) {
+void AddConstantBufferDescriptor(Info& info, u32 index, u32 count) {
+ if (count != 1) {
+ throw NotImplementedException("Constant buffer descriptor indexing");
+ }
+ if ((info.constant_buffer_mask & (1U << index)) != 0) {
return;
}
- descriptor = &info.constant_buffer_descriptors.emplace_back(Info::ConstantBufferDescriptor{
+ info.constant_buffer_mask |= 1U << index;
+ info.constant_buffer_descriptors.push_back({
.index{index},
.count{1},
});
}
-void Visit(Info& info, IR::Inst& inst) {
+void VisitUsages(Info& info, IR::Inst& inst) {
switch (inst.Opcode()) {
case IR::Opcode::WorkgroupId:
info.uses_workgroup_id = true;
@@ -72,7 +77,7 @@ void Visit(Info& info, IR::Inst& inst) {
break;
case IR::Opcode::GetCbuf:
if (const IR::Value index{inst.Arg(0)}; index.IsImmediate()) {
- AddConstantBufferDescriptor(info, index.U32());
+ AddConstantBufferDescriptor(info, index.U32(), 1);
} else {
throw NotImplementedException("Constant buffer with non-immediate index");
}
@@ -81,6 +86,60 @@ void Visit(Info& info, IR::Inst& inst) {
break;
}
}
+
+void VisitFpModifiers(Info& info, IR::Inst& inst) {
+ switch (inst.Opcode()) {
+ 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: {
+ const auto control{inst.Flags<IR::FpControl>()};
+ switch (control.fmz_mode) {
+ case IR::FmzMode::DontCare:
+ break;
+ case IR::FmzMode::FTZ:
+ case IR::FmzMode::FMZ:
+ info.uses_fp16_denorms_flush = true;
+ break;
+ case IR::FmzMode::None:
+ info.uses_fp16_denorms_preserve = true;
+ break;
+ }
+ break;
+ }
+ 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: {
+ const auto control{inst.Flags<IR::FpControl>()};
+ switch (control.fmz_mode) {
+ case IR::FmzMode::DontCare:
+ break;
+ case IR::FmzMode::FTZ:
+ case IR::FmzMode::FMZ:
+ info.uses_fp32_denorms_flush = true;
+ break;
+ case IR::FmzMode::None:
+ info.uses_fp32_denorms_preserve = true;
+ break;
+ }
+ break;
+ }
+ default:
+ break;
+ }
+}
+
+void Visit(Info& info, IR::Inst& inst) {
+ VisitUsages(info, inst);
+ VisitFpModifiers(info, inst);
+}
} // Anonymous namespace
void CollectShaderInfoPass(IR::Program& program) {
diff --git a/src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp b/src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp
index bf230a850..03bd547b7 100644
--- a/src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp
+++ b/src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp
@@ -351,7 +351,6 @@ void GlobalMemoryToStorageBufferPass(IR::Program& program) {
.cbuf_offset{storage_buffer.offset},
.count{1},
});
- info.storage_buffers[storage_index] = &info.storage_buffers_descriptors.back();
++storage_index;
}
for (const StorageInst& storage_inst : to_replace) {
diff --git a/src/shader_recompiler/main.cpp b/src/shader_recompiler/main.cpp
index abd44e323..72565f477 100644
--- a/src/shader_recompiler/main.cpp
+++ b/src/shader_recompiler/main.cpp
@@ -60,6 +60,17 @@ void RunDatabase() {
fmt::print(stdout, "{} ms", duration_cast<milliseconds>(t - t0).count() / double(N));
}
+static constexpr Profile PROFILE{
+ .unified_descriptor_binding = true,
+ .support_float_controls = true,
+ .support_separate_denorm_behavior = true,
+ .support_separate_rounding_mode = true,
+ .support_fp16_denorm_preserve = true,
+ .support_fp32_denorm_preserve = true,
+ .support_fp16_denorm_flush = true,
+ .support_fp32_denorm_flush = true,
+};
+
int main() {
// RunDatabase();
@@ -76,7 +87,7 @@ int main() {
fmt::print(stdout, "{}\n", cfg.Dot());
IR::Program program{TranslateProgram(inst_pool, block_pool, env, cfg)};
fmt::print(stdout, "{}\n", IR::DumpProgram(program));
- const std::vector<u32> spirv{Backend::SPIRV::EmitSPIRV(env, program)};
+ const std::vector<u32> spirv{Backend::SPIRV::EmitSPIRV(PROFILE, env, program)};
std::FILE* const file{std::fopen("D:\\shader.spv", "wb")};
std::fwrite(spirv.data(), spirv.size(), sizeof(u32), file);
std::fclose(file);
diff --git a/src/shader_recompiler/profile.h b/src/shader_recompiler/profile.h
index c96d783b7..9881bebab 100644
--- a/src/shader_recompiler/profile.h
+++ b/src/shader_recompiler/profile.h
@@ -7,7 +7,14 @@
namespace Shader {
struct Profile {
- bool unified_descriptor_binding;
+ bool unified_descriptor_binding{};
+ bool support_float_controls{};
+ bool support_separate_denorm_behavior{};
+ bool support_separate_rounding_mode{};
+ bool support_fp16_denorm_preserve{};
+ bool support_fp32_denorm_preserve{};
+ bool support_fp16_denorm_flush{};
+ bool support_fp32_denorm_flush{};
};
} // namespace Shader
diff --git a/src/shader_recompiler/recompiler.cpp b/src/shader_recompiler/recompiler.cpp
index b25081e39..527e19c27 100644
--- a/src/shader_recompiler/recompiler.cpp
+++ b/src/shader_recompiler/recompiler.cpp
@@ -14,14 +14,15 @@
namespace Shader {
-std::pair<Info, std::vector<u32>> RecompileSPIRV(Environment& env, u32 start_address) {
+std::pair<Info, std::vector<u32>> RecompileSPIRV(const Profile& profile, Environment& env,
+ u32 start_address) {
ObjectPool<Maxwell::Flow::Block> flow_block_pool;
ObjectPool<IR::Inst> inst_pool;
ObjectPool<IR::Block> block_pool;
Maxwell::Flow::CFG cfg{env, flow_block_pool, start_address};
IR::Program program{Maxwell::TranslateProgram(inst_pool, block_pool, env, cfg)};
- return {std::move(program.info), Backend::SPIRV::EmitSPIRV(env, program)};
+ return {std::move(program.info), Backend::SPIRV::EmitSPIRV(profile, env, program)};
}
} // namespace Shader
diff --git a/src/shader_recompiler/recompiler.h b/src/shader_recompiler/recompiler.h
index 4cb973878..2529463ae 100644
--- a/src/shader_recompiler/recompiler.h
+++ b/src/shader_recompiler/recompiler.h
@@ -9,10 +9,12 @@
#include "common/common_types.h"
#include "shader_recompiler/environment.h"
+#include "shader_recompiler/profile.h"
#include "shader_recompiler/shader_info.h"
namespace Shader {
-[[nodiscard]] std::pair<Info, std::vector<u32>> RecompileSPIRV(Environment& env, u32 start_address);
+[[nodiscard]] std::pair<Info, std::vector<u32>> RecompileSPIRV(const Profile& profile,
+ Environment& env, u32 start_address);
} // namespace Shader
diff --git a/src/shader_recompiler/shader_info.h b/src/shader_recompiler/shader_info.h
index f49a79368..8766bf13e 100644
--- a/src/shader_recompiler/shader_info.h
+++ b/src/shader_recompiler/shader_info.h
@@ -31,14 +31,15 @@ struct Info {
bool uses_local_invocation_id{};
bool uses_fp16{};
bool uses_fp64{};
+ bool uses_fp16_denorms_flush{};
+ bool uses_fp16_denorms_preserve{};
+ bool uses_fp32_denorms_flush{};
+ bool uses_fp32_denorms_preserve{};
u32 constant_buffer_mask{};
- std::array<ConstantBufferDescriptor*, MAX_CBUFS> constant_buffers{};
boost::container::static_vector<ConstantBufferDescriptor, MAX_CBUFS>
constant_buffer_descriptors;
-
- std::array<StorageBufferDescriptor*, MAX_SSBOS> storage_buffers{};
boost::container::static_vector<StorageBufferDescriptor, MAX_SSBOS> storage_buffers_descriptors;
};
generated by cgit v1.2.3 (git 2.25.1) at 2024-11-13 04:12:02 +0100