ShaderCompiler: Inline driver specific constants.

This commit is contained in:
Fernando Sahmkow 2022-12-27 21:39:46 -05:00
parent b62ffb612d
commit a045e860dd
5 changed files with 39 additions and 3 deletions

View file

@ -57,11 +57,16 @@ public:
return start_address; return start_address;
} }
[[nodiscard]] bool IsPropietaryDriver() const noexcept {
return is_propietary_driver;
}
protected: protected:
ProgramHeader sph{}; ProgramHeader sph{};
std::array<u32, 8> gp_passthrough_mask{}; std::array<u32, 8> gp_passthrough_mask{};
Stage stage{}; Stage stage{};
u32 start_address{}; u32 start_address{};
bool is_propietary_driver{};
}; };
} // namespace Shader } // namespace Shader

View file

@ -677,6 +677,30 @@ void FoldConstBuffer(Environment& env, IR::Block& block, IR::Inst& inst) {
} }
} }
void FoldDriverConstBuffer(Environment& env, IR::Block& block, IR::Inst& inst, u32 which_bank,
u32 offset_start = 0, u32 offset_end = std::numeric_limits<u16>::max()) {
const IR::Value bank{inst.Arg(0)};
const IR::Value offset{inst.Arg(1)};
if (!bank.IsImmediate() || !offset.IsImmediate()) {
return;
}
const auto bank_value = bank.U32();
if (bank_value != which_bank) {
return;
}
const auto offset_value = offset.U32();
if (offset_value < offset_start || offset_value >= offset_end) {
return;
}
IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(inst)};
if (inst.GetOpcode() == IR::Opcode::GetCbufU32) {
inst.ReplaceUsesWith(IR::Value{env.ReadCbufValue(bank_value, offset_value)});
} else {
inst.ReplaceUsesWith(
IR::Value{Common::BitCast<f32>(env.ReadCbufValue(bank_value, offset_value))});
}
}
void ConstantPropagation(Environment& env, IR::Block& block, IR::Inst& inst) { void ConstantPropagation(Environment& env, IR::Block& block, IR::Inst& inst) {
switch (inst.GetOpcode()) { switch (inst.GetOpcode()) {
case IR::Opcode::GetRegister: case IR::Opcode::GetRegister:
@ -825,13 +849,17 @@ void ConstantPropagation(Environment& env, IR::Block& block, IR::Inst& inst) {
case IR::Opcode::GetCbufF32: case IR::Opcode::GetCbufF32:
case IR::Opcode::GetCbufU32: case IR::Opcode::GetCbufU32:
if (env.HasHLEMacroState()) { if (env.HasHLEMacroState()) {
return FoldConstBuffer(env, block, inst); FoldConstBuffer(env, block, inst);
}
if (env.IsPropietaryDriver()) {
FoldDriverConstBuffer(env, block, inst, 1);
} }
break; break;
default: default:
break; break;
} }
} }
} // Anonymous namespace } // Anonymous namespace
void ConstantPropagationPass(Environment& env, IR::Program& program) { void ConstantPropagationPass(Environment& env, IR::Program& program) {

View file

@ -51,7 +51,7 @@ using VideoCommon::LoadPipelines;
using VideoCommon::SerializePipeline; using VideoCommon::SerializePipeline;
using Context = ShaderContext::Context; using Context = ShaderContext::Context;
constexpr u32 CACHE_VERSION = 8; constexpr u32 CACHE_VERSION = 9;
template <typename Container> template <typename Container>
auto MakeSpan(Container& container) { auto MakeSpan(Container& container) {

View file

@ -54,7 +54,7 @@ using VideoCommon::FileEnvironment;
using VideoCommon::GenericEnvironment; using VideoCommon::GenericEnvironment;
using VideoCommon::GraphicsEnvironment; using VideoCommon::GraphicsEnvironment;
constexpr u32 CACHE_VERSION = 9; constexpr u32 CACHE_VERSION = 10;
template <typename Container> template <typename Container>
auto MakeSpan(Container& container) { auto MakeSpan(Container& container) {

View file

@ -325,6 +325,7 @@ GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_,
ASSERT(local_size <= std::numeric_limits<u32>::max()); ASSERT(local_size <= std::numeric_limits<u32>::max());
local_memory_size = static_cast<u32>(local_size) + sph.common3.shader_local_memory_crs_size; local_memory_size = static_cast<u32>(local_size) + sph.common3.shader_local_memory_crs_size;
texture_bound = maxwell3d->regs.bindless_texture_const_buffer_slot; texture_bound = maxwell3d->regs.bindless_texture_const_buffer_slot;
is_propietary_driver = texture_bound == 2;
has_hle_engine_state = has_hle_engine_state =
maxwell3d->engine_state == Tegra::Engines::Maxwell3D::EngineHint::OnHLEMacro; maxwell3d->engine_state == Tegra::Engines::Maxwell3D::EngineHint::OnHLEMacro;
} }
@ -399,6 +400,7 @@ ComputeEnvironment::ComputeEnvironment(Tegra::Engines::KeplerCompute& kepler_com
stage = Shader::Stage::Compute; stage = Shader::Stage::Compute;
local_memory_size = qmd.local_pos_alloc + qmd.local_crs_alloc; local_memory_size = qmd.local_pos_alloc + qmd.local_crs_alloc;
texture_bound = kepler_compute->regs.tex_cb_index; texture_bound = kepler_compute->regs.tex_cb_index;
is_propietary_driver = texture_bound == 2;
shared_memory_size = qmd.shared_alloc; shared_memory_size = qmd.shared_alloc;
workgroup_size = {qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z}; workgroup_size = {qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z};
} }
@ -498,6 +500,7 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
file.read(reinterpret_cast<char*>(&gp_passthrough_mask), sizeof(gp_passthrough_mask)); file.read(reinterpret_cast<char*>(&gp_passthrough_mask), sizeof(gp_passthrough_mask));
} }
} }
is_propietary_driver = texture_bound == 2;
} }
void FileEnvironment::Dump(u64 hash) { void FileEnvironment::Dump(u64 hash) {