summaryrefslogtreecommitdiffstats
path: root/src
diff options
context:
space:
mode:
authorFernandoS27 <fsahmkow27@gmail.com>2021-03-27 22:30:24 +0100
committerameerj <52414509+ameerj@users.noreply.github.com>2021-07-23 03:51:25 +0200
commit34aba9627a8fad20b3b173180e2f3d679dd32293 (patch)
treea4f2faec67a793e8b44493532a683908dcefb4d8 /src
parentshader: Fix alignment checks on RZ (diff)
downloadyuzu-34aba9627a8fad20b3b173180e2f3d679dd32293.tar
yuzu-34aba9627a8fad20b3b173180e2f3d679dd32293.tar.gz
yuzu-34aba9627a8fad20b3b173180e2f3d679dd32293.tar.bz2
yuzu-34aba9627a8fad20b3b173180e2f3d679dd32293.tar.lz
yuzu-34aba9627a8fad20b3b173180e2f3d679dd32293.tar.xz
yuzu-34aba9627a8fad20b3b173180e2f3d679dd32293.tar.zst
yuzu-34aba9627a8fad20b3b173180e2f3d679dd32293.zip
Diffstat (limited to 'src')
-rw-r--r--src/shader_recompiler/CMakeLists.txt4
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv.h3
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp10
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp4
-rw-r--r--src/shader_recompiler/environment.h2
-rw-r--r--src/shader_recompiler/frontend/ir/ir_emitter.cpp12
-rw-r--r--src/shader_recompiler/frontend/ir/ir_emitter.h4
-rw-r--r--src/shader_recompiler/frontend/ir/microinstruction.cpp1
-rw-r--r--src/shader_recompiler/frontend/ir/opcodes.inc3
-rw-r--r--src/shader_recompiler/frontend/maxwell/control_flow.cpp58
-rw-r--r--src/shader_recompiler/frontend/maxwell/control_flow.h7
-rw-r--r--src/shader_recompiler/frontend/maxwell/indirect_branch_table_track.cpp108
-rw-r--r--src/shader_recompiler/frontend/maxwell/indirect_branch_table_track.h28
-rw-r--r--src/shader_recompiler/frontend/maxwell/instruction.h1
-rw-r--r--src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp57
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/branch_indirect.cpp36
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.cpp29
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.h39
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp8
-rw-r--r--src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp21
-rw-r--r--src/video_core/renderer_vulkan/vk_pipeline_cache.cpp50
21 files changed, 437 insertions, 48 deletions
diff --git a/src/shader_recompiler/CMakeLists.txt b/src/shader_recompiler/CMakeLists.txt
index 003cbefb1..44ab929b7 100644
--- a/src/shader_recompiler/CMakeLists.txt
+++ b/src/shader_recompiler/CMakeLists.txt
@@ -52,6 +52,8 @@ add_library(shader_recompiler STATIC
frontend/maxwell/control_flow.h
frontend/maxwell/decode.cpp
frontend/maxwell/decode.h
+ frontend/maxwell/indirect_branch_table_track.cpp
+ frontend/maxwell/indirect_branch_table_track.h
frontend/maxwell/instruction.h
frontend/maxwell/location.h
frontend/maxwell/maxwell.inc
@@ -63,6 +65,7 @@ add_library(shader_recompiler STATIC
frontend/maxwell/structured_control_flow.h
frontend/maxwell/translate/impl/bitfield_extract.cpp
frontend/maxwell/translate/impl/bitfield_insert.cpp
+ frontend/maxwell/translate/impl/branch_indirect.cpp
frontend/maxwell/translate/impl/common_encoding.h
frontend/maxwell/translate/impl/common_funcs.cpp
frontend/maxwell/translate/impl/common_funcs.h
@@ -110,6 +113,7 @@ add_library(shader_recompiler STATIC
frontend/maxwell/translate/impl/integer_short_multiply_add.cpp
frontend/maxwell/translate/impl/integer_to_integer_conversion.cpp
frontend/maxwell/translate/impl/load_constant.cpp
+ frontend/maxwell/translate/impl/load_constant.h
frontend/maxwell/translate/impl/load_effective_address.cpp
frontend/maxwell/translate/impl/load_store_attribute.cpp
frontend/maxwell/translate/impl/load_store_local_shared.cpp
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.h b/src/shader_recompiler/backend/spirv/emit_spirv.h
index 204c5f9e0..02648d769 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv.h
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.h
@@ -26,6 +26,7 @@ void EmitBranchConditional(EmitContext& ctx, Id condition, Id true_label, Id fal
void EmitLoopMerge(EmitContext& ctx, Id merge_label, Id continue_label);
void EmitSelectionMerge(EmitContext& ctx, Id merge_label);
void EmitReturn(EmitContext& ctx);
+void EmitUnreachable(EmitContext& ctx);
void EmitDemoteToHelperInvocation(EmitContext& ctx, Id continue_label);
void EmitPrologue(EmitContext& ctx);
void EmitEpilogue(EmitContext& ctx);
@@ -35,6 +36,8 @@ void EmitGetPred(EmitContext& ctx);
void EmitSetPred(EmitContext& ctx);
void EmitSetGotoVariable(EmitContext& ctx);
void EmitGetGotoVariable(EmitContext& ctx);
+void EmitSetIndirectBranchVariable(EmitContext& ctx);
+void EmitGetIndirectBranchVariable(EmitContext& ctx);
Id EmitGetCbufU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
Id EmitGetCbufS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
Id EmitGetCbufU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
index 52dcef8a4..4a267b16c 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
@@ -6,8 +6,6 @@
#include "shader_recompiler/backend/spirv/emit_spirv.h"
-#pragma optimize("", off)
-
namespace Shader::Backend::SPIRV {
namespace {
struct AttrInfo {
@@ -74,6 +72,14 @@ void EmitGetGotoVariable(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
+void EmitSetIndirectBranchVariable(EmitContext&) {
+ throw NotImplementedException("SPIR-V Instruction");
+}
+
+void EmitGetIndirectBranchVariable(EmitContext&) {
+ throw NotImplementedException("SPIR-V Instruction");
+}
+
static Id GetCbuf(EmitContext& ctx, Id result_type, Id UniformDefinitions::*member_ptr,
u32 element_size, const IR::Value& binding, const IR::Value& offset) {
if (!binding.IsImmediate()) {
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp
index 6b81f0169..335603f88 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp
@@ -26,6 +26,10 @@ void EmitReturn(EmitContext& ctx) {
ctx.OpReturn();
}
+void EmitUnreachable(EmitContext& ctx) {
+ ctx.OpUnreachable();
+}
+
void EmitDemoteToHelperInvocation(EmitContext& ctx, Id continue_label) {
ctx.OpDemoteToHelperInvocationEXT();
ctx.OpBranch(continue_label);
diff --git a/src/shader_recompiler/environment.h b/src/shader_recompiler/environment.h
index 9415d02f6..1c50ae51e 100644
--- a/src/shader_recompiler/environment.h
+++ b/src/shader_recompiler/environment.h
@@ -15,6 +15,8 @@ public:
[[nodiscard]] virtual u64 ReadInstruction(u32 address) = 0;
+ [[nodiscard]] virtual u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) = 0;
+
[[nodiscard]] virtual TextureType ReadTextureType(u32 cbuf_index, u32 cbuf_offset) = 0;
[[nodiscard]] virtual u32 TextureBoundBuffer() const = 0;
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.cpp b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
index 9b898e4e1..552472487 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.cpp
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
@@ -87,6 +87,10 @@ void IREmitter::Return() {
Inst(Opcode::Return);
}
+void IREmitter::Unreachable() {
+ Inst(Opcode::Unreachable);
+}
+
void IREmitter::DemoteToHelperInvocation(Block* continue_label) {
block->SetBranch(continue_label);
continue_label->AddImmediatePredecessor(block);
@@ -126,6 +130,14 @@ void IREmitter::SetGotoVariable(u32 id, const U1& value) {
Inst(Opcode::SetGotoVariable, id, value);
}
+U32 IREmitter::GetIndirectBranchVariable() {
+ return Inst<U32>(Opcode::GetIndirectBranchVariable);
+}
+
+void IREmitter::SetIndirectBranchVariable(const U32& value) {
+ Inst(Opcode::SetIndirectBranchVariable, value);
+}
+
void IREmitter::SetPred(IR::Pred pred, const U1& value) {
Inst(Opcode::SetPred, pred, value);
}
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.h b/src/shader_recompiler/frontend/ir/ir_emitter.h
index 269f367a4..17bc32fc8 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.h
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.h
@@ -37,6 +37,7 @@ public:
void LoopMerge(Block* merge_block, Block* continue_target);
void SelectionMerge(Block* merge_block);
void Return();
+ void Unreachable();
void DemoteToHelperInvocation(Block* continue_label);
void Prologue();
@@ -51,6 +52,9 @@ public:
[[nodiscard]] U1 GetGotoVariable(u32 id);
void SetGotoVariable(u32 id, const U1& value);
+ [[nodiscard]] U32 GetIndirectBranchVariable();
+ void SetIndirectBranchVariable(const U32& value);
+
[[nodiscard]] U32 GetCbuf(const U32& binding, const U32& byte_offset);
[[nodiscard]] UAny GetCbuf(const U32& binding, const U32& byte_offset, size_t bitsize,
bool is_signed);
diff --git a/src/shader_recompiler/frontend/ir/microinstruction.cpp b/src/shader_recompiler/frontend/ir/microinstruction.cpp
index 52a5e5034..c3ba6b522 100644
--- a/src/shader_recompiler/frontend/ir/microinstruction.cpp
+++ b/src/shader_recompiler/frontend/ir/microinstruction.cpp
@@ -55,6 +55,7 @@ bool Inst::MayHaveSideEffects() const noexcept {
case Opcode::LoopMerge:
case Opcode::SelectionMerge:
case Opcode::Return:
+ case Opcode::Unreachable:
case Opcode::DemoteToHelperInvocation:
case Opcode::Prologue:
case Opcode::Epilogue:
diff --git a/src/shader_recompiler/frontend/ir/opcodes.inc b/src/shader_recompiler/frontend/ir/opcodes.inc
index 9b050995b..fb79e3d8d 100644
--- a/src/shader_recompiler/frontend/ir/opcodes.inc
+++ b/src/shader_recompiler/frontend/ir/opcodes.inc
@@ -13,6 +13,7 @@ OPCODE(BranchConditional, Void, U1,
OPCODE(LoopMerge, Void, Label, Label, )
OPCODE(SelectionMerge, Void, Label, )
OPCODE(Return, Void, )
+OPCODE(Unreachable, Void, )
OPCODE(DemoteToHelperInvocation, Void, Label, )
// Special operations
@@ -26,6 +27,8 @@ OPCODE(GetPred, U1, Pred
OPCODE(SetPred, Void, Pred, U1, )
OPCODE(GetGotoVariable, U1, U32, )
OPCODE(SetGotoVariable, Void, U32, U1, )
+OPCODE(GetIndirectBranchVariable, U32, )
+OPCODE(SetIndirectBranchVariable, Void, U32, )
OPCODE(GetCbufU8, U32, U32, U32, )
OPCODE(GetCbufS8, U32, U32, U32, )
OPCODE(GetCbufU16, U32, U32, U32, )
diff --git a/src/shader_recompiler/frontend/maxwell/control_flow.cpp b/src/shader_recompiler/frontend/maxwell/control_flow.cpp
index 4f6707fae..1e9b8e426 100644
--- a/src/shader_recompiler/frontend/maxwell/control_flow.cpp
+++ b/src/shader_recompiler/frontend/maxwell/control_flow.cpp
@@ -14,6 +14,7 @@
#include "shader_recompiler/exception.h"
#include "shader_recompiler/frontend/maxwell/control_flow.h"
#include "shader_recompiler/frontend/maxwell/decode.h"
+#include "shader_recompiler/frontend/maxwell/indirect_branch_table_track.h"
#include "shader_recompiler/frontend/maxwell/location.h"
namespace Shader::Maxwell::Flow {
@@ -252,9 +253,7 @@ CFG::AnalysisState CFG::AnalyzeInst(Block* block, FunctionId function_id, Locati
const Opcode opcode{Decode(inst.raw)};
switch (opcode) {
case Opcode::BRA:
- case Opcode::BRX:
case Opcode::JMP:
- case Opcode::JMX:
case Opcode::RET:
if (!AnalyzeBranch(block, function_id, pc, inst, opcode)) {
return AnalysisState::Continue;
@@ -264,10 +263,6 @@ CFG::AnalysisState CFG::AnalyzeInst(Block* block, FunctionId function_id, Locati
case Opcode::JMP:
AnalyzeBRA(block, function_id, pc, inst, IsAbsoluteJump(opcode));
break;
- case Opcode::BRX:
- case Opcode::JMX:
- AnalyzeBRX(block, pc, inst, IsAbsoluteJump(opcode));
- break;
case Opcode::RET:
block->end_class = EndClass::Return;
break;
@@ -302,6 +297,9 @@ CFG::AnalysisState CFG::AnalyzeInst(Block* block, FunctionId function_id, Locati
case Opcode::SSY:
block->stack.Push(OpcodeToken(opcode), BranchOffset(pc, inst));
return AnalysisState::Continue;
+ case Opcode::BRX:
+ case Opcode::JMX:
+ return AnalyzeBRX(block, pc, inst, IsAbsoluteJump(opcode), function_id);
case Opcode::EXIT:
return AnalyzeEXIT(block, function_id, pc, inst);
case Opcode::PRET:
@@ -407,8 +405,46 @@ void CFG::AnalyzeBRA(Block* block, FunctionId function_id, Location pc, Instruct
block->branch_true = AddLabel(block, block->stack, bra_pc, function_id);
}
-void CFG::AnalyzeBRX(Block*, Location, Instruction, bool is_absolute) {
- throw NotImplementedException("{}", is_absolute ? "JMX" : "BRX");
+CFG::AnalysisState CFG::AnalyzeBRX(Block* block, Location pc, Instruction inst, bool is_absolute,
+ FunctionId function_id) {
+ const std::optional brx_table{TrackIndirectBranchTable(env, pc, block->begin)};
+ if (!brx_table) {
+ TrackIndirectBranchTable(env, pc, block->begin);
+ throw NotImplementedException("Failed to track indirect branch");
+ }
+ const IR::FlowTest flow_test{inst.branch.flow_test};
+ const Predicate pred{inst.Pred()};
+ if (flow_test != IR::FlowTest::T || pred != Predicate{true}) {
+ throw NotImplementedException("Conditional indirect branch");
+ }
+ std::vector<u32> targets;
+ targets.reserve(brx_table->num_entries);
+ for (u32 i = 0; i < brx_table->num_entries; ++i) {
+ u32 target{env.ReadCbufValue(brx_table->cbuf_index, brx_table->cbuf_offset + i * 4)};
+ if (!is_absolute) {
+ target += pc.Offset();
+ }
+ target += brx_table->branch_offset;
+ target += 8;
+ targets.push_back(target);
+ }
+ std::ranges::sort(targets);
+ targets.erase(std::unique(targets.begin(), targets.end()), targets.end());
+
+ block->indirect_branches.reserve(targets.size());
+ for (const u32 target : targets) {
+ Block* const branch{AddLabel(block, block->stack, target, function_id)};
+ block->indirect_branches.push_back(branch);
+ }
+ block->cond = IR::Condition{true};
+ block->end = pc + 1;
+ block->end_class = EndClass::IndirectBranch;
+ block->branch_reg = brx_table->branch_reg;
+ block->branch_offset = brx_table->branch_offset + 8;
+ if (!is_absolute) {
+ block->branch_offset += pc.Offset();
+ }
+ return AnalysisState::Branch;
}
CFG::AnalysisState CFG::AnalyzeEXIT(Block* block, FunctionId function_id, Location pc,
@@ -449,7 +485,6 @@ Block* CFG::AddLabel(Block* block, Stack stack, Location pc, FunctionId function
// Block already exists and it has been visited
return &*it;
}
- // TODO: FIX DANGLING BLOCKS
Block* const new_block{block_pool.Create(Block{
.begin{pc},
.end{pc},
@@ -494,6 +529,11 @@ std::string CFG::Dot() const {
add_branch(block.branch_false, false);
}
break;
+ case EndClass::IndirectBranch:
+ for (Block* const branch : block.indirect_branches) {
+ add_branch(branch, false);
+ }
+ break;
case EndClass::Call:
dot += fmt::format("\t\t{}->N{};\n", name, node_uid);
dot += fmt::format("\t\tN{}->{};\n", node_uid, NameOf(*block.return_block));
diff --git a/src/shader_recompiler/frontend/maxwell/control_flow.h b/src/shader_recompiler/frontend/maxwell/control_flow.h
index 22f134194..1e05fcb97 100644
--- a/src/shader_recompiler/frontend/maxwell/control_flow.h
+++ b/src/shader_recompiler/frontend/maxwell/control_flow.h
@@ -26,6 +26,7 @@ using FunctionId = size_t;
enum class EndClass {
Branch,
+ IndirectBranch,
Call,
Exit,
Return,
@@ -76,11 +77,14 @@ struct Block : boost::intrusive::set_base_hook<
union {
Block* branch_true;
FunctionId function_call;
+ IR::Reg branch_reg;
};
union {
Block* branch_false;
Block* return_block;
+ s32 branch_offset;
};
+ std::vector<Block*> indirect_branches;
};
struct Label {
@@ -139,7 +143,8 @@ private:
void AnalyzeBRA(Block* block, FunctionId function_id, Location pc, Instruction inst,
bool is_absolute);
- void AnalyzeBRX(Block* block, Location pc, Instruction inst, bool is_absolute);
+ AnalysisState AnalyzeBRX(Block* block, Location pc, Instruction inst, bool is_absolute,
+ FunctionId function_id);
AnalysisState AnalyzeEXIT(Block* block, FunctionId function_id, Location pc, Instruction inst);
/// Return the branch target block id
diff --git a/src/shader_recompiler/frontend/maxwell/indirect_branch_table_track.cpp b/src/shader_recompiler/frontend/maxwell/indirect_branch_table_track.cpp
new file mode 100644
index 000000000..96453509d
--- /dev/null
+++ b/src/shader_recompiler/frontend/maxwell/indirect_branch_table_track.cpp
@@ -0,0 +1,108 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include <optional>
+
+#include "common/common_types.h"
+#include "shader_recompiler/exception.h"
+#include "shader_recompiler/frontend/maxwell/decode.h"
+#include "shader_recompiler/frontend/maxwell/indirect_branch_table_track.h"
+#include "shader_recompiler/frontend/maxwell/opcodes.h"
+#include "shader_recompiler/frontend/maxwell/translate/impl/load_constant.h"
+
+namespace Shader::Maxwell {
+namespace {
+union Encoding {
+ u64 raw;
+ BitField<0, 8, IR::Reg> dest_reg;
+ BitField<8, 8, IR::Reg> src_reg;
+ BitField<20, 19, u64> immediate;
+ BitField<56, 1, u64> is_negative;
+ BitField<20, 24, s64> brx_offset;
+};
+
+template <typename Callable>
+std::optional<u64> Track(Environment& env, Location block_begin, Location& pos, Callable&& func) {
+ while (pos >= block_begin) {
+ const u64 insn{env.ReadInstruction(pos.Offset())};
+ --pos;
+ if (func(insn, Decode(insn))) {
+ return insn;
+ }
+ }
+ return std::nullopt;
+}
+
+std::optional<u64> TrackLDC(Environment& env, Location block_begin, Location& pos,
+ IR::Reg brx_reg) {
+ return Track(env, block_begin, pos, [brx_reg](u64 insn, Opcode opcode) {
+ const LDC::Encoding ldc{insn};
+ return opcode == Opcode::LDC && ldc.dest_reg == brx_reg && ldc.size == LDC::Size::B32 &&
+ ldc.mode == LDC::Mode::Default;
+ });
+}
+
+std::optional<u64> TrackSHL(Environment& env, Location block_begin, Location& pos,
+ IR::Reg ldc_reg) {
+ return Track(env, block_begin, pos, [ldc_reg](u64 insn, Opcode opcode) {
+ const Encoding shl{insn};
+ return opcode == Opcode::SHL_imm && shl.dest_reg == ldc_reg;
+ });
+}
+
+std::optional<u64> TrackIMNMX(Environment& env, Location block_begin, Location& pos,
+ IR::Reg shl_reg) {
+ return Track(env, block_begin, pos, [shl_reg](u64 insn, Opcode opcode) {
+ const Encoding imnmx{insn};
+ return opcode == Opcode::IMNMX_imm && imnmx.dest_reg == shl_reg;
+ });
+}
+} // Anonymous namespace
+
+std::optional<IndirectBranchTableInfo> TrackIndirectBranchTable(Environment& env, Location brx_pos,
+ Location block_begin) {
+ const u64 brx_insn{env.ReadInstruction(brx_pos.Offset())};
+ const Opcode brx_opcode{Decode(brx_insn)};
+ if (brx_opcode != Opcode::BRX && brx_opcode != Opcode::JMX) {
+ throw LogicError("Tracked instruction is not BRX or JMX");
+ }
+ const IR::Reg brx_reg{Encoding{brx_insn}.src_reg};
+ const s32 brx_offset{static_cast<s32>(Encoding{brx_insn}.brx_offset)};
+
+ Location pos{brx_pos};
+ const std::optional<u64> ldc_insn{TrackLDC(env, block_begin, pos, brx_reg)};
+ if (!ldc_insn) {
+ return std::nullopt;
+ }
+ const LDC::Encoding ldc{*ldc_insn};
+ const u32 cbuf_index{static_cast<u32>(ldc.index)};
+ const u32 cbuf_offset{static_cast<u32>(static_cast<s32>(ldc.offset.Value()))};
+ const IR::Reg ldc_reg{ldc.src_reg};
+
+ const std::optional<u64> shl_insn{TrackSHL(env, block_begin, pos, ldc_reg)};
+ if (!shl_insn) {
+ return std::nullopt;
+ }
+ const Encoding shl{*shl_insn};
+ const IR::Reg shl_reg{shl.src_reg};
+
+ const std::optional<u64> imnmx_insn{TrackIMNMX(env, block_begin, pos, shl_reg)};
+ if (!imnmx_insn) {
+ return std::nullopt;
+ }
+ const Encoding imnmx{*imnmx_insn};
+ if (imnmx.is_negative != 0) {
+ return std::nullopt;
+ }
+ const u32 imnmx_immediate{static_cast<u32>(imnmx.immediate.Value())};
+ return IndirectBranchTableInfo{
+ .cbuf_index{cbuf_index},
+ .cbuf_offset{cbuf_offset},
+ .num_entries{imnmx_immediate + 1},
+ .branch_offset{brx_offset},
+ .branch_reg{brx_reg},
+ };
+}
+
+} // namespace Shader::Maxwell
diff --git a/src/shader_recompiler/frontend/maxwell/indirect_branch_table_track.h b/src/shader_recompiler/frontend/maxwell/indirect_branch_table_track.h
new file mode 100644
index 000000000..eee5102fa
--- /dev/null
+++ b/src/shader_recompiler/frontend/maxwell/indirect_branch_table_track.h
@@ -0,0 +1,28 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#pragma once
+
+#include <optional>
+
+#include "common/bit_field.h"
+#include "common/common_types.h"
+#include "shader_recompiler/environment.h"
+#include "shader_recompiler/frontend/ir/reg.h"
+#include "shader_recompiler/frontend/maxwell/location.h"
+
+namespace Shader::Maxwell {
+
+struct IndirectBranchTableInfo {
+ u32 cbuf_index{};
+ u32 cbuf_offset{};
+ u32 num_entries{};
+ s32 branch_offset{};
+ IR::Reg branch_reg{};
+};
+
+std::optional<IndirectBranchTableInfo> TrackIndirectBranchTable(Environment& env, Location brx_pos,
+ Location block_begin);
+
+} // namespace Shader::Maxwell
diff --git a/src/shader_recompiler/frontend/maxwell/instruction.h b/src/shader_recompiler/frontend/maxwell/instruction.h
index 57fd531f2..743d68d61 100644
--- a/src/shader_recompiler/frontend/maxwell/instruction.h
+++ b/src/shader_recompiler/frontend/maxwell/instruction.h
@@ -7,6 +7,7 @@
#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/ir/flow_test.h"
+#include "shader_recompiler/frontend/ir/reg.h"
namespace Shader::Maxwell {
diff --git a/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp b/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
index 9d4688390..a6e55f61e 100644
--- a/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
+++ b/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
@@ -17,6 +17,7 @@
#include "shader_recompiler/environment.h"
#include "shader_recompiler/frontend/ir/basic_block.h"
#include "shader_recompiler/frontend/ir/ir_emitter.h"
+#include "shader_recompiler/frontend/maxwell/decode.h"
#include "shader_recompiler/frontend/maxwell/structured_control_flow.h"
#include "shader_recompiler/frontend/maxwell/translate/translate.h"
#include "shader_recompiler/object_pool.h"
@@ -46,12 +47,15 @@ enum class StatementType {
Break,
Return,
Kill,
+ Unreachable,
Function,
Identity,
Not,
Or,
SetVariable,
+ SetIndirectBranchVariable,
Variable,
+ IndirectBranchCond,
};
bool HasChildren(StatementType type) {
@@ -72,12 +76,15 @@ struct Loop {};
struct Break {};
struct Return {};
struct Kill {};
+struct Unreachable {};
struct FunctionTag {};
struct Identity {};
struct Not {};
struct Or {};
struct SetVariable {};
+struct SetIndirectBranchVariable {};
struct Variable {};
+struct IndirectBranchCond {};
#ifdef _MSC_VER
#pragma warning(push)
@@ -96,6 +103,7 @@ struct Statement : ListBaseHook {
: cond{cond_}, up{up_}, type{StatementType::Break} {}
Statement(Return) : type{StatementType::Return} {}
Statement(Kill) : type{StatementType::Kill} {}
+ Statement(Unreachable) : type{StatementType::Unreachable} {}
Statement(FunctionTag) : children{}, type{StatementType::Function} {}
Statement(Identity, IR::Condition cond_) : guest_cond{cond_}, type{StatementType::Identity} {}
Statement(Not, Statement* op_) : op{op_}, type{StatementType::Not} {}
@@ -103,7 +111,12 @@ struct Statement : ListBaseHook {
: op_a{op_a_}, op_b{op_b_}, type{StatementType::Or} {}
Statement(SetVariable, u32 id_, Statement* op_, Statement* up_)
: op{op_}, id{id_}, up{up_}, type{StatementType::SetVariable} {}
+ Statement(SetIndirectBranchVariable, IR::Reg branch_reg_, s32 branch_offset_)
+ : branch_offset{branch_offset_},
+ branch_reg{branch_reg_}, type{StatementType::SetIndirectBranchVariable} {}
Statement(Variable, u32 id_) : id{id_}, type{StatementType::Variable} {}
+ Statement(IndirectBranchCond, u32 location_)
+ : location{location_}, type{StatementType::IndirectBranchCond} {}
~Statement() {
if (HasChildren(type)) {
@@ -118,11 +131,14 @@ struct Statement : ListBaseHook {
IR::Condition guest_cond;
Statement* op;
Statement* op_a;
+ u32 location;
+ s32 branch_offset;
};
union {
Statement* cond;
Statement* op_b;
u32 id;
+ IR::Reg branch_reg;
};
Statement* up{};
StatementType type;
@@ -141,6 +157,8 @@ std::string DumpExpr(const Statement* stmt) {
return fmt::format("{} || {}", DumpExpr(stmt->op_a), DumpExpr(stmt->op_b));
case StatementType::Variable:
return fmt::format("goto_L{}", stmt->id);
+ case StatementType::IndirectBranchCond:
+ return fmt::format("(indirect_branch == {:x})", stmt->location);
default:
return "<invalid type>";
}
@@ -182,14 +200,22 @@ std::string DumpTree(const Tree& tree, u32 indentation = 0) {
case StatementType::Kill:
ret += fmt::format("{} kill;\n", indent);
break;
+ case StatementType::Unreachable:
+ ret += fmt::format("{} unreachable;\n", indent);
+ break;
case StatementType::SetVariable:
ret += fmt::format("{} goto_L{} = {};\n", indent, stmt->id, DumpExpr(stmt->op));
break;
+ case StatementType::SetIndirectBranchVariable:
+ ret += fmt::format("{} indirect_branch = {} + {};\n", indent, stmt->branch_reg,
+ stmt->branch_offset);
+ break;
case StatementType::Function:
case StatementType::Identity:
case StatementType::Not:
case StatementType::Or:
case StatementType::Variable:
+ case StatementType::IndirectBranchCond:
throw LogicError("Statement can't be printed");
}
}
@@ -417,6 +443,17 @@ private:
}
break;
}
+ case Flow::EndClass::IndirectBranch:
+ root.insert(ip, *pool.Create(SetIndirectBranchVariable{}, block.branch_reg,
+ block.branch_offset));
+ for (Flow::Block* const branch : block.indirect_branches) {
+ const Node indirect_label{local_labels.at(branch)};
+ Statement* cond{pool.Create(IndirectBranchCond{}, branch->begin.Offset())};
+ Statement* goto_stmt{pool.Create(Goto{}, cond, indirect_label, &root_stmt)};
+ gotos.push_back(root.insert(ip, *goto_stmt));
+ }
+ root.insert(ip, *pool.Create(Unreachable{}));
+ break;
case Flow::EndClass::Call: {
Flow::Function& call{cfg.Functions()[block.function_call]};
const Node call_return_label{local_labels.at(block.return_block)};
@@ -623,6 +660,8 @@ IR::Block* TryFindForwardBlock(const Statement& stmt) {
return ir.LogicalOr(VisitExpr(ir, *stmt.op_a), VisitExpr(ir, *stmt.op_b));
case StatementType::Variable:
return ir.GetGotoVariable(stmt.id);
+ case StatementType::IndirectBranchCond:
+ return ir.IEqual(ir.GetIndirectBranchVariable(), ir.Imm32(stmt.location));
default:
throw NotImplementedException("Statement type {}", stmt.type);
}
@@ -670,6 +709,15 @@ private:
ir.SetGotoVariable(stmt.id, VisitExpr(ir, *stmt.op));
break;
}
+ case StatementType::SetIndirectBranchVariable: {
+ if (!current_block) {
+ current_block = MergeBlock(parent, stmt);
+ }
+ IR::IREmitter ir{*current_block};
+ IR::U32 address{ir.IAdd(ir.GetReg(stmt.branch_reg), ir.Imm32(stmt.branch_offset))};
+ ir.SetIndirectBranchVariable(address);
+ break;
+ }
case StatementType::If: {
if (!current_block) {
current_block = block_pool.Create(inst_pool);
@@ -756,6 +804,15 @@ private:
current_block = demote_block;
break;
}
+ case StatementType::Unreachable: {
+ if (!current_block) {
+ current_block = block_pool.Create(inst_pool);
+ block_list.push_back(current_block);
+ }
+ IR::IREmitter{*current_block}.Unreachable();
+ current_block = nullptr;
+ break;
+ }
default:
throw NotImplementedException("Statement type {}", stmt.type);
}
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/branch_indirect.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/branch_indirect.cpp
new file mode 100644
index 000000000..371c0e0f7
--- /dev/null
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/branch_indirect.cpp
@@ -0,0 +1,36 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include "common/bit_field.h"
+#include "common/common_types.h"
+#include "shader_recompiler/exception.h"
+#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
+
+namespace Shader::Maxwell {
+namespace {
+void Check(u64 insn) {
+ union {
+ u64 raw;
+ BitField<5, 1, u64> cbuf_mode;
+ BitField<6, 1, u64> lmt;
+ } const encoding{insn};
+
+ if (encoding.cbuf_mode != 0) {
+ throw NotImplementedException("Constant buffer mode");
+ }
+ if (encoding.lmt != 0) {
+ throw NotImplementedException("LMT");
+ }
+}
+} // Anonymous namespace
+
+void TranslatorVisitor::BRX(u64 insn) {
+ Check(insn);
+}
+
+void TranslatorVisitor::JMX(u64 insn) {
+ Check(insn);
+}
+
+} // namespace Shader::Maxwell
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.cpp
index 39becf93c..49ccb7d62 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.cpp
@@ -5,25 +5,11 @@
#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
+#include "shader_recompiler/frontend/maxwell/translate/impl/load_constant.h"
namespace Shader::Maxwell {
+using namespace LDC;
namespace {
-enum class Mode : u64 {
- Default,
- IL,
- IS,
- ISL,
-};
-
-enum class Size : u64 {
- U8,
- S8,
- U16,
- S16,
- B32,
- B64,
-};
-
std::pair<IR::U32, IR::U32> Slot(IR::IREmitter& ir, Mode mode, const IR::U32& imm_index,
const IR::U32& reg, const IR::U32& imm) {
switch (mode) {
@@ -37,16 +23,7 @@ std::pair<IR::U32, IR::U32> Slot(IR::IREmitter& ir, Mode mode, const IR::U32& im
} // Anonymous namespace
void TranslatorVisitor::LDC(u64 insn) {
- union {
- u64 raw;
- BitField<0, 8, IR::Reg> dest_reg;
- BitField<8, 8, IR::Reg> src_reg;
- BitField<20, 16, s64> offset;
- BitField<36, 5, u64> index;
- BitField<44, 2, Mode> mode;
- BitField<48, 3, Size> size;
- } const ldc{insn};
-
+ const Encoding ldc{insn};
const IR::U32 imm_index{ir.Imm32(static_cast<u32>(ldc.index))};
const IR::U32 reg{X(ldc.src_reg)};
const IR::U32 imm{ir.Imm32(static_cast<s32>(ldc.offset))};
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.h b/src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.h
new file mode 100644
index 000000000..3074ea0e3
--- /dev/null
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/load_constant.h
@@ -0,0 +1,39 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#pragma once
+
+#include "common/bit_field.h"
+#include "common/common_types.h"
+#include "shader_recompiler/frontend/ir/reg.h"
+
+namespace Shader::Maxwell::LDC {
+
+enum class Mode : u64 {
+ Default,
+ IL,
+ IS,
+ ISL,
+};
+
+enum class Size : u64 {
+ U8,
+ S8,
+ U16,
+ S16,
+ B32,
+ B64,
+};
+
+union Encoding {
+ u64 raw;
+ BitField<0, 8, IR::Reg> dest_reg;
+ BitField<8, 8, IR::Reg> src_reg;
+ BitField<20, 16, s64> offset;
+ BitField<36, 5, u64> index;
+ BitField<44, 2, Mode> mode;
+ BitField<48, 3, Size> size;
+};
+
+} // namespace Shader::Maxwell::LDC
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp
index b62d8ee2a..a0057a473 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp
@@ -53,10 +53,6 @@ void TranslatorVisitor::BRK(u64) {
ThrowNotImplemented(Opcode::BRK);
}
-void TranslatorVisitor::BRX(u64) {
- ThrowNotImplemented(Opcode::BRX);
-}
-
void TranslatorVisitor::CAL() {
// CAL is a no-op
}
@@ -181,10 +177,6 @@ void TranslatorVisitor::JMP(u64) {
ThrowNotImplemented(Opcode::JMP);
}
-void TranslatorVisitor::JMX(u64) {
- ThrowNotImplemented(Opcode::JMX);
-}
-
void TranslatorVisitor::KIL() {
// KIL is a no-op
}
diff --git a/src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp b/src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp
index bab7ca186..259233746 100644
--- a/src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp
+++ b/src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp
@@ -48,8 +48,12 @@ struct GotoVariable : FlagTag {
u32 index;
};
+struct IndirectBranchVariable {
+ auto operator<=>(const IndirectBranchVariable&) const noexcept = default;
+};
+
using Variant = std::variant<IR::Reg, IR::Pred, ZeroFlagTag, SignFlagTag, CarryFlagTag,
- OverflowFlagTag, GotoVariable>;
+ OverflowFlagTag, GotoVariable, IndirectBranchVariable>;
using ValueMap = boost::container::flat_map<IR::Block*, IR::Value, std::less<IR::Block*>>;
struct DefTable {
@@ -65,6 +69,10 @@ struct DefTable {
return goto_vars[goto_variable.index];
}
+ [[nodiscard]] ValueMap& operator[](IndirectBranchVariable) {
+ return indirect_branch_var;
+ }
+
[[nodiscard]] ValueMap& operator[](ZeroFlagTag) noexcept {
return zero_flag;
}
@@ -84,6 +92,7 @@ struct DefTable {
std::array<ValueMap, IR::NUM_USER_REGS> regs;
std::array<ValueMap, IR::NUM_USER_PREDS> preds;
boost::container::flat_map<u32, ValueMap> goto_vars;
+ ValueMap indirect_branch_var;
ValueMap zero_flag;
ValueMap sign_flag;
ValueMap carry_flag;
@@ -102,6 +111,10 @@ IR::Opcode UndefOpcode(const FlagTag&) noexcept {
return IR::Opcode::UndefU1;
}
+IR::Opcode UndefOpcode(IndirectBranchVariable) noexcept {
+ return IR::Opcode::UndefU32;
+}
+
[[nodiscard]] bool IsPhi(const IR::Inst& inst) noexcept {
return inst.Opcode() == IR::Opcode::Phi;
}
@@ -219,6 +232,9 @@ void VisitInst(Pass& pass, IR::Block* block, IR::Inst& inst) {
case IR::Opcode::SetGotoVariable:
pass.WriteVariable(GotoVariable{inst.Arg(0).U32()}, block, inst.Arg(1));
break;
+ case IR::Opcode::SetIndirectBranchVariable:
+ pass.WriteVariable(IndirectBranchVariable{}, block, inst.Arg(0));
+ break;
case IR::Opcode::SetZFlag:
pass.WriteVariable(ZeroFlagTag{}, block, inst.Arg(0));
break;
@@ -244,6 +260,9 @@ void VisitInst(Pass& pass, IR::Block* block, IR::Inst& inst) {
case IR::Opcode::GetGotoVariable:
inst.ReplaceUsesWith(pass.ReadVariable(GotoVariable{inst.Arg(0).U32()}, block));
break;
+ case IR::Opcode::GetIndirectBranchVariable:
+ inst.ReplaceUsesWith(pass.ReadVariable(IndirectBranchVariable{}, block));
+ break;
case IR::Opcode::GetZFlag:
inst.ReplaceUsesWith(pass.ReadVariable(ZeroFlagTag{}, block));
break;
diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
index 8b2816c13..6cde01491 100644
--- a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
+++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
@@ -47,7 +47,7 @@ auto MakeSpan(Container& container) {
}
u64 MakeCbufKey(u32 index, u32 offset) {
- return (static_cast<u64>(index) << 32) | static_cast<u64>(offset);
+ return (static_cast<u64>(index) << 32) | offset;
}
class GenericEnvironment : public Shader::Environment {
@@ -114,11 +114,13 @@ public:
gpu_memory->ReadBlock(program_base + read_lowest, data.get(), code_size);
const u64 num_texture_types{static_cast<u64>(texture_types.size())};
+ const u64 num_cbuf_values{static_cast<u64>(cbuf_values.size())};
const u32 local_memory_size{LocalMemorySize()};
const u32 texture_bound{TextureBoundBuffer()};
file.write(reinterpret_cast<const char*>(&code_size), sizeof(code_size))
.write(reinterpret_cast<const char*>(&num_texture_types), sizeof(num_texture_types))
+ .write(reinterpret_cast<const char*>(&num_cbuf_values), sizeof(num_cbuf_values))
.write(reinterpret_cast<const char*>(&local_memory_size), sizeof(local_memory_size))
.write(reinterpret_cast<const char*>(&texture_bound), sizeof(texture_bound))
.write(reinterpret_cast<const char*>(&start_address), sizeof(start_address))
@@ -130,6 +132,10 @@ public:
file.write(reinterpret_cast<const char*>(&key), sizeof(key))
.write(reinterpret_cast<const char*>(&type), sizeof(type));
}
+ for (const auto [key, type] : cbuf_values) {
+ file.write(reinterpret_cast<const char*>(&key), sizeof(key))
+ .write(reinterpret_cast<const char*>(&type), sizeof(type));
+ }
if (stage == Shader::Stage::Compute) {
const std::array<u32, 3> workgroup_size{WorkgroupSize()};
const u32 shared_memory_size{SharedMemorySize()};
@@ -212,6 +218,7 @@ protected:
std::vector<u64> code;
std::unordered_map<u64, Shader::TextureType> texture_types;
+ std::unordered_map<u64, u32> cbuf_values;
u32 read_lowest = std::numeric_limits<u32>::max();
u32 read_highest = 0;
@@ -267,6 +274,17 @@ public:
~GraphicsEnvironment() override = default;
+ u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override {
+ const auto& cbuf{maxwell3d->state.shader_stages[stage_index].const_buffers[cbuf_index]};
+ ASSERT(cbuf.enabled);
+ u32 value{};
+ if (cbuf_offset < cbuf.size) {
+ value = gpu_memory->Read<u32>(cbuf.address + cbuf_offset);
+ }
+ cbuf_values.emplace(MakeCbufKey(cbuf_index, cbuf_offset), value);
+ return value;
+ }
+
Shader::TextureType ReadTextureType(u32 cbuf_index, u32 cbuf_offset) override {
const auto& regs{maxwell3d->regs};
const auto& cbuf{maxwell3d->state.shader_stages[stage_index].const_buffers[cbuf_index]};
@@ -312,6 +330,18 @@ public:
~ComputeEnvironment() override = default;
+ u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override {
+ const auto& qmd{kepler_compute->launch_description};
+ ASSERT(((qmd.const_buffer_enable_mask.Value() >> cbuf_index) & 1) != 0);
+ const auto& cbuf{qmd.const_buffer_config[cbuf_index]};
+ u32 value{};
+ if (cbuf_offset < cbuf.size) {
+ value = gpu_memory->Read<u32>(cbuf.Address() + cbuf_offset);
+ }
+ cbuf_values.emplace(MakeCbufKey(cbuf_index, cbuf_offset), value);
+ return value;
+ }
+
Shader::TextureType ReadTextureType(u32 cbuf_index, u32 cbuf_offset) override {
const auto& regs{kepler_compute->regs};
const auto& qmd{kepler_compute->launch_description};
@@ -386,8 +416,10 @@ public:
void Deserialize(std::ifstream& file) {
u64 code_size{};
u64 num_texture_types{};
+ u64 num_cbuf_values{};
file.read(reinterpret_cast<char*>(&code_size), sizeof(code_size))
.read(reinterpret_cast<char*>(&num_texture_types), sizeof(num_texture_types))
+ .read(reinterpret_cast<char*>(&num_cbuf_values), sizeof(num_cbuf_values))
.read(reinterpret_cast<char*>(&local_memory_size), sizeof(local_memory_size))
.read(reinterpret_cast<char*>(&texture_bound), sizeof(texture_bound))
.read(reinterpret_cast<char*>(&start_address), sizeof(start_address))
@@ -403,6 +435,13 @@ public:
.read(reinterpret_cast<char*>(&type), sizeof(type));
texture_types.emplace(key, type);
}
+ for (size_t i = 0; i < num_cbuf_values; ++i) {
+ u64 key;
+ u32 value;
+ file.read(reinterpret_cast<char*>(&key), sizeof(key))
+ .read(reinterpret_cast<char*>(&value), sizeof(value));
+ cbuf_values.emplace(key, value);
+ }
if (stage == Shader::Stage::Compute) {
file.read(reinterpret_cast<char*>(&workgroup_size), sizeof(workgroup_size))
.read(reinterpret_cast<char*>(&shared_memory_size), sizeof(shared_memory_size));
@@ -418,6 +457,14 @@ public:
return code[(address - read_lowest) / sizeof(u64)];
}
+ u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override {
+ const auto it{cbuf_values.find(MakeCbufKey(cbuf_index, cbuf_offset))};
+ if (it == cbuf_values.end()) {
+ throw Shader::LogicError("Uncached read texture type");
+ }
+ return it->second;
+ }
+
Shader::TextureType ReadTextureType(u32 cbuf_index, u32 cbuf_offset) override {
const auto it{texture_types.find(MakeCbufKey(cbuf_index, cbuf_offset))};
if (it == texture_types.end()) {
@@ -445,6 +492,7 @@ public:
private:
std::unique_ptr<u64[]> code;
std::unordered_map<u64, Shader::TextureType> texture_types;
+ std::unordered_map<u64, u32> cbuf_values;
std::array<u32, 3> workgroup_size{};
u32 local_memory_size{};
u32 shared_memory_size{};