shader: Implement BRX

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

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);

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()) {

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);

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;

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);
 }

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);

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:

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,                                                            )

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) {
+CFG::AnalysisState CFG::AnalyzeBRX(Block* block, Location pc, Instruction inst, bool is_absolute,
-    throw NotImplementedException("{}", is_absolute ? "JMX" : "BRX");
+                                   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));

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

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

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

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 {
 

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);
             }

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

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 {
+    const Encoding ldc{insn};
-        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 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))};

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

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
 }

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;

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{};