From b10cf64c486d8730fcfeb53a333814915b3b5fbe Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Wed, 5 May 2021 02:19:08 -0300
Subject: glasm: Add GLASM backend infrastructure

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 95 ++++++++++++++++++++++
 1 file changed, 95 insertions(+)
 create mode 100644 src/shader_recompiler/backend/glasm/emit_glasm.cpp

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
new file mode 100644
index 000000000..59d7c0f96
--- /dev/null
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -0,0 +1,95 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include <string>
+#include <tuple>
+
+#include "shader_recompiler/backend/bindings.h"
+#include "shader_recompiler/backend/glasm/emit_context.h"
+#include "shader_recompiler/backend/glasm/emit_glasm.h"
+#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/frontend/ir/program.h"
+#include "shader_recompiler/profile.h"
+
+namespace Shader::Backend::GLASM {
+namespace {
+template <class Func>
+struct FuncTraits {};
+
+template <class ReturnType_, class... Args>
+struct FuncTraits<ReturnType_ (*)(Args...)> {
+    using ReturnType = ReturnType_;
+
+    static constexpr size_t NUM_ARGS = sizeof...(Args);
+
+    template <size_t I>
+    using ArgType = std::tuple_element_t<I, std::tuple<Args...>>;
+};
+
+template <typename ArgType>
+auto Arg(EmitContext& ctx, const IR::Value& arg) {
+    if constexpr (std::is_same_v<ArgType, std::string_view>) {
+        return ctx.reg_alloc.Consume(arg);
+    } else if constexpr (std::is_same_v<ArgType, const IR::Value&>) {
+        return arg;
+    } else if constexpr (std::is_same_v<ArgType, u32>) {
+        return arg.U32();
+    } else if constexpr (std::is_same_v<ArgType, IR::Block*>) {
+        return arg.Label();
+    } else if constexpr (std::is_same_v<ArgType, IR::Attribute>) {
+        return arg.Attribute();
+    } else if constexpr (std::is_same_v<ArgType, IR::Patch>) {
+        return arg.Patch();
+    } else if constexpr (std::is_same_v<ArgType, IR::Reg>) {
+        return arg.Reg();
+    }
+}
+
+template <auto func, bool is_first_arg_inst, size_t... I>
+void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) {
+    using Traits = FuncTraits<decltype(func)>;
+    if constexpr (is_first_arg_inst) {
+        func(ctx, inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
+    } else {
+        func(ctx, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...);
+    }
+}
+
+template <auto func>
+void Invoke(EmitContext& ctx, IR::Inst* inst) {
+    using Traits = FuncTraits<decltype(func)>;
+    static_assert(Traits::NUM_ARGS >= 1, "Insufficient arguments");
+    if constexpr (Traits::NUM_ARGS == 1) {
+        Invoke<func, false>(ctx, inst, std::make_index_sequence<0>{});
+    } else {
+        using FirstArgType = typename Traits::template ArgType<1>;
+        static constexpr bool is_first_arg_inst = std::is_same_v<FirstArgType, IR::Inst*>;
+        using Indices = std::make_index_sequence<Traits::NUM_ARGS - (is_first_arg_inst ? 2 : 1)>;
+        Invoke<func, is_first_arg_inst>(ctx, inst, Indices{});
+    }
+}
+
+void EmitInst(EmitContext& ctx, IR::Inst* inst) {
+    switch (inst->GetOpcode()) {
+#define OPCODE(name, result_type, ...)                                                             \
+    case IR::Opcode::name:                                                                         \
+        return Invoke<&Emit##name>(ctx, inst);
+#include "shader_recompiler/frontend/ir/opcodes.inc"
+#undef OPCODE
+    }
+    throw LogicError("Invalid opcode {}", inst->GetOpcode());
+}
+} // Anonymous namespace
+
+std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
+    EmitContext ctx;
+    for (IR::Block* const block : program.blocks) {
+        for (IR::Inst& inst : block->Instructions()) {
+            EmitInst(ctx, &inst);
+        }
+    }
+    return ctx.code;
+}
+
+} // namespace Shader::Backend::GLASM
-- 
cgit v1.2.3


From 6fd190d1ae4275a06ed2e488401e1d63912954be Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sat, 8 May 2021 16:28:52 -0300
Subject: glasm: Implement basic GLASM instructions

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 66 +++++++++++++++++++++-
 1 file changed, 63 insertions(+), 3 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 59d7c0f96..65600f58c 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -50,7 +50,7 @@ template <auto func, bool is_first_arg_inst, size_t... I>
 void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) {
     using Traits = FuncTraits<decltype(func)>;
     if constexpr (is_first_arg_inst) {
-        func(ctx, inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
+        func(ctx, *inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
     } else {
         func(ctx, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...);
     }
@@ -64,7 +64,7 @@ void Invoke(EmitContext& ctx, IR::Inst* inst) {
         Invoke<func, false>(ctx, inst, std::make_index_sequence<0>{});
     } else {
         using FirstArgType = typename Traits::template ArgType<1>;
-        static constexpr bool is_first_arg_inst = std::is_same_v<FirstArgType, IR::Inst*>;
+        static constexpr bool is_first_arg_inst = std::is_same_v<FirstArgType, IR::Inst&>;
         using Indices = std::make_index_sequence<Traits::NUM_ARGS - (is_first_arg_inst ? 2 : 1)>;
         Invoke<func, is_first_arg_inst>(ctx, inst, Indices{});
     }
@@ -80,16 +80,76 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) {
     }
     throw LogicError("Invalid opcode {}", inst->GetOpcode());
 }
+
+void Identity(IR::Inst& inst, const IR::Value& value) {
+    if (value.IsImmediate()) {
+        return;
+    }
+    IR::Inst* const value_inst{value.InstRecursive()};
+    if (inst.GetOpcode() == IR::Opcode::Identity) {
+        value_inst->DestructiveAddUsage(inst.UseCount());
+        value_inst->DestructiveRemoveUsage();
+    }
+    inst.SetDefinition(value_inst->Definition<Id>());
+}
 } // Anonymous namespace
 
 std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
-    EmitContext ctx;
+    EmitContext ctx{program};
     for (IR::Block* const block : program.blocks) {
         for (IR::Inst& inst : block->Instructions()) {
             EmitInst(ctx, &inst);
         }
     }
+    std::string header = "!!NVcp5.0\n"
+                         "OPTION NV_internal;";
+    switch (program.stage) {
+    case Stage::Compute:
+        header += fmt::format("GROUP_SIZE {} {} {};", program.workgroup_size[0],
+                              program.workgroup_size[1], program.workgroup_size[2]);
+        break;
+    default:
+        break;
+    }
+    header += "TEMP ";
+    for (size_t index = 0; index < ctx.reg_alloc.NumUsedRegisters(); ++index) {
+        header += fmt::format("R{},", index);
+    }
+    header += "RC;";
+    if (!program.info.storage_buffers_descriptors.empty()) {
+        header += "LONG TEMP LC;";
+    }
+    ctx.code.insert(0, header);
+    ctx.code += "END";
     return ctx.code;
 }
 
+void EmitIdentity(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+    Identity(inst, value);
+}
+
+void EmitBitCastU16F16(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+    Identity(inst, value);
+}
+
+void EmitBitCastU32F32(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+    Identity(inst, value);
+}
+
+void EmitBitCastU64F64(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+    Identity(inst, value);
+}
+
+void EmitBitCastF16U16(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+    Identity(inst, value);
+}
+
+void EmitBitCastF32U32(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+    Identity(inst, value);
+}
+
+void EmitBitCastF64U64(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+    Identity(inst, value);
+}
+
 } // namespace Shader::Backend::GLASM
-- 
cgit v1.2.3


From 2b04b4d27fc38d9865cef5bf8eabb335bc29eb83 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sat, 8 May 2021 16:43:26 -0300
Subject: glasm: Remove unused argument in identity instructions on GLASM

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 65600f58c..7ec880c81 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -124,31 +124,31 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     return ctx.code;
 }
 
-void EmitIdentity(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+void EmitIdentity(EmitContext&, IR::Inst& inst, const IR::Value& value) {
     Identity(inst, value);
 }
 
-void EmitBitCastU16F16(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+void EmitBitCastU16F16(EmitContext&, IR::Inst& inst, const IR::Value& value) {
     Identity(inst, value);
 }
 
-void EmitBitCastU32F32(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+void EmitBitCastU32F32(EmitContext&, IR::Inst& inst, const IR::Value& value) {
     Identity(inst, value);
 }
 
-void EmitBitCastU64F64(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+void EmitBitCastU64F64(EmitContext&, IR::Inst& inst, const IR::Value& value) {
     Identity(inst, value);
 }
 
-void EmitBitCastF16U16(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+void EmitBitCastF16U16(EmitContext&, IR::Inst& inst, const IR::Value& value) {
     Identity(inst, value);
 }
 
-void EmitBitCastF32U32(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+void EmitBitCastF32U32(EmitContext&, IR::Inst& inst, const IR::Value& value) {
     Identity(inst, value);
 }
 
-void EmitBitCastF64U64(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) {
+void EmitBitCastF64U64(EmitContext&, IR::Inst& inst, const IR::Value& value) {
     Identity(inst, value);
 }
 
-- 
cgit v1.2.3


From 1c9307969c4e3f6206947f1360acae33f95a4a86 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sun, 9 May 2021 03:11:34 -0300
Subject: glasm: Make GLASM aware of types

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 95 ++++++++++++++++++----
 1 file changed, 77 insertions(+), 18 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 7ec880c81..8981cf300 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -27,22 +27,80 @@ struct FuncTraits<ReturnType_ (*)(Args...)> {
     using ArgType = std::tuple_element_t<I, std::tuple<Args...>>;
 };
 
+template <typename T>
+struct Identity {
+    Identity(const T& data_) : data{data_} {}
+
+    const T& Extract() {
+        return data;
+    }
+
+    T data;
+};
+
+template <bool scalar>
+struct RegWrapper {
+    RegWrapper(EmitContext& ctx, Value value)
+        : reg_alloc{ctx.reg_alloc}, allocated{value.type != Type::Register} {
+        reg = allocated ? reg_alloc.AllocReg() : Register{value};
+        switch (value.type) {
+        case Type::Register:
+            break;
+        case Type::U32:
+            ctx.Add("MOV.U {}.x,{};", reg, value.imm_u32);
+            break;
+        case Type::S32:
+            ctx.Add("MOV.S {}.x,{};", reg, value.imm_s32);
+            break;
+        case Type::F32:
+            ctx.Add("MOV.F {}.x,{};", reg, value.imm_f32);
+            break;
+        }
+    }
+    ~RegWrapper() {
+        if (allocated) {
+            reg_alloc.FreeReg(reg);
+        }
+    }
+
+    auto Extract() {
+        return std::conditional_t<scalar, ScalarRegister, Register>{Value{reg}};
+    }
+
+    RegAlloc& reg_alloc;
+    Register reg{};
+    bool allocated{};
+};
+
 template <typename ArgType>
 auto Arg(EmitContext& ctx, const IR::Value& arg) {
-    if constexpr (std::is_same_v<ArgType, std::string_view>) {
-        return ctx.reg_alloc.Consume(arg);
+    if constexpr (std::is_same_v<ArgType, Register>) {
+        return RegWrapper<false>{ctx, ctx.reg_alloc.Consume(arg)};
+    } else if constexpr (std::is_same_v<ArgType, ScalarRegister>) {
+        return RegWrapper<true>{ctx, ctx.reg_alloc.Consume(arg)};
+    } else if constexpr (std::is_base_of_v<Value, ArgType>) {
+        return Identity{ArgType{ctx.reg_alloc.Consume(arg)}};
     } else if constexpr (std::is_same_v<ArgType, const IR::Value&>) {
-        return arg;
+        return Identity{arg};
     } else if constexpr (std::is_same_v<ArgType, u32>) {
-        return arg.U32();
+        return Identity{arg.U32()};
     } else if constexpr (std::is_same_v<ArgType, IR::Block*>) {
-        return arg.Label();
+        return Identity{arg.Label()};
     } else if constexpr (std::is_same_v<ArgType, IR::Attribute>) {
-        return arg.Attribute();
+        return Identity{arg.Attribute()};
     } else if constexpr (std::is_same_v<ArgType, IR::Patch>) {
-        return arg.Patch();
+        return Identity{arg.Patch()};
     } else if constexpr (std::is_same_v<ArgType, IR::Reg>) {
-        return arg.Reg();
+        return Identity{arg.Reg()};
+    }
+}
+
+template <auto func, bool is_first_arg_inst, typename... Args>
+void InvokeCall(EmitContext& ctx, IR::Inst* inst, Args&&... args) {
+    if constexpr (is_first_arg_inst) {
+        func(ctx, *inst, std::forward<Args>(args.Extract())...);
+    } else {
+        func(ctx, std::forward<Args>(args.Extract())...);
     }
 }
 
@@ -50,9 +108,10 @@ template <auto func, bool is_first_arg_inst, size_t... I>
 void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) {
     using Traits = FuncTraits<decltype(func)>;
     if constexpr (is_first_arg_inst) {
-        func(ctx, *inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
+        func(ctx, *inst,
+             Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I)).Extract()...);
     } else {
-        func(ctx, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...);
+        func(ctx, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I)).Extract()...);
     }
 }
 
@@ -81,7 +140,7 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) {
     throw LogicError("Invalid opcode {}", inst->GetOpcode());
 }
 
-void Identity(IR::Inst& inst, const IR::Value& value) {
+void Alias(IR::Inst& inst, const IR::Value& value) {
     if (value.IsImmediate()) {
         return;
     }
@@ -125,31 +184,31 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
 }
 
 void EmitIdentity(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Identity(inst, value);
+    Alias(inst, value);
 }
 
 void EmitBitCastU16F16(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Identity(inst, value);
+    Alias(inst, value);
 }
 
 void EmitBitCastU32F32(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Identity(inst, value);
+    Alias(inst, value);
 }
 
 void EmitBitCastU64F64(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Identity(inst, value);
+    Alias(inst, value);
 }
 
 void EmitBitCastF16U16(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Identity(inst, value);
+    Alias(inst, value);
 }
 
 void EmitBitCastF32U32(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Identity(inst, value);
+    Alias(inst, value);
 }
 
 void EmitBitCastF64U64(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Identity(inst, value);
+    Alias(inst, value);
 }
 
 } // namespace Shader::Backend::GLASM
-- 
cgit v1.2.3


From 9f851e3832fb85c20f406eacfadd12a8bb7d982d Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sun, 9 May 2021 04:18:37 -0300
Subject: glasm: Implement GLASM fp16 packing and move bitwise insns

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 40 ----------------------
 1 file changed, 40 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 8981cf300..842ec157d 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -139,18 +139,6 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) {
     }
     throw LogicError("Invalid opcode {}", inst->GetOpcode());
 }
-
-void Alias(IR::Inst& inst, const IR::Value& value) {
-    if (value.IsImmediate()) {
-        return;
-    }
-    IR::Inst* const value_inst{value.InstRecursive()};
-    if (inst.GetOpcode() == IR::Opcode::Identity) {
-        value_inst->DestructiveAddUsage(inst.UseCount());
-        value_inst->DestructiveRemoveUsage();
-    }
-    inst.SetDefinition(value_inst->Definition<Id>());
-}
 } // Anonymous namespace
 
 std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
@@ -183,32 +171,4 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     return ctx.code;
 }
 
-void EmitIdentity(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Alias(inst, value);
-}
-
-void EmitBitCastU16F16(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Alias(inst, value);
-}
-
-void EmitBitCastU32F32(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Alias(inst, value);
-}
-
-void EmitBitCastU64F64(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Alias(inst, value);
-}
-
-void EmitBitCastF16U16(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Alias(inst, value);
-}
-
-void EmitBitCastF32U32(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Alias(inst, value);
-}
-
-void EmitBitCastF64U64(EmitContext&, IR::Inst& inst, const IR::Value& value) {
-    Alias(inst, value);
-}
-
 } // namespace Shader::Backend::GLASM
-- 
cgit v1.2.3


From 4502595bc2518eecf934110e9393b11bf0c2f75a Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sun, 9 May 2021 18:03:01 -0300
Subject: glasm: Initial GLASM fp64 support

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 17 +++++++++++++----
 1 file changed, 13 insertions(+), 4 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 842ec157d..9db6eb4a0 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -42,7 +42,11 @@ template <bool scalar>
 struct RegWrapper {
     RegWrapper(EmitContext& ctx, Value value)
         : reg_alloc{ctx.reg_alloc}, allocated{value.type != Type::Register} {
-        reg = allocated ? reg_alloc.AllocReg() : Register{value};
+        if (allocated) {
+            reg = value.type == Type::F64 ? reg_alloc.AllocLongReg() : reg_alloc.AllocReg();
+        } else {
+            reg = Register{value};
+        }
         switch (value.type) {
         case Type::Register:
             break;
@@ -55,6 +59,9 @@ struct RegWrapper {
         case Type::F32:
             ctx.Add("MOV.F {}.x,{};", reg, value.imm_f32);
             break;
+        case Type::F64:
+            ctx.Add("MOV.F64 {}.x,{};", reg, value.imm_f64);
+            break;
         }
     }
     ~RegWrapper() {
@@ -162,10 +169,12 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     for (size_t index = 0; index < ctx.reg_alloc.NumUsedRegisters(); ++index) {
         header += fmt::format("R{},", index);
     }
-    header += "RC;";
-    if (!program.info.storage_buffers_descriptors.empty()) {
-        header += "LONG TEMP LC;";
+    header += "RC;"
+              "LONG TEMP ";
+    for (size_t index = 0; index < ctx.reg_alloc.NumUsedLongRegisters(); ++index) {
+        header += fmt::format("D{},", index);
     }
+    header += "DC;";
     ctx.code.insert(0, header);
     ctx.code += "END";
     return ctx.code;
-- 
cgit v1.2.3


From ad61b47f80b96436ef675abcf1123668d9c1180d Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sun, 9 May 2021 22:43:29 -0300
Subject: glasm: Add conversion instructions to GLASM

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 +++
 1 file changed, 3 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 9db6eb4a0..0e4b189c9 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -59,6 +59,9 @@ struct RegWrapper {
         case Type::F32:
             ctx.Add("MOV.F {}.x,{};", reg, value.imm_f32);
             break;
+        case Type::U64:
+            ctx.Add("MOV.U64 {}.x,{};", reg, value.imm_u64);
+            break;
         case Type::F64:
             ctx.Add("MOV.F64 {}.x,{};", reg, value.imm_f64);
             break;
-- 
cgit v1.2.3


From 80813b1d144a7f0f11047e7348620b720def93a9 Mon Sep 17 00:00:00 2001
From: ameerj <52414509+ameerj@users.noreply.github.com>
Date: Sun, 9 May 2021 22:01:03 -0400
Subject: glasm: Implement storage atomic ops

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 13 +++++++++++++
 1 file changed, 13 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 0e4b189c9..e6e065e7f 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -149,6 +149,18 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) {
     }
     throw LogicError("Invalid opcode {}", inst->GetOpcode());
 }
+
+void SetupOptions(std::string& header, Info info) {
+    if (info.uses_int64_bit_atomics) {
+        header += "OPTION NV_shader_atomic_int64;";
+    }
+    if (info.uses_atomic_f32_add) {
+        header += "OPTION NV_shader_atomic_float;";
+    }
+    if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) {
+        header += "OPTION NV_shader_atomic_fp16_vector;";
+    }
+}
 } // Anonymous namespace
 
 std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
@@ -160,6 +172,7 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     }
     std::string header = "!!NVcp5.0\n"
                          "OPTION NV_internal;";
+    SetupOptions(header, program.info);
     switch (program.stage) {
     case Stage::Compute:
         header += fmt::format("GROUP_SIZE {} {} {};", program.workgroup_size[0],
-- 
cgit v1.2.3


From 8eb72ff0dc3eb428c28b578ffb3912c1bd1c42dd Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 10 May 2021 01:33:24 -0300
Subject: glasm: Fix moving U64 immediates to registers in GLASM

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index e6e065e7f..8e5d575a9 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -43,7 +43,8 @@ struct RegWrapper {
     RegWrapper(EmitContext& ctx, Value value)
         : reg_alloc{ctx.reg_alloc}, allocated{value.type != Type::Register} {
         if (allocated) {
-            reg = value.type == Type::F64 ? reg_alloc.AllocLongReg() : reg_alloc.AllocReg();
+            const bool is_long{value.type == Type::F64 || value.type == Type::U64};
+            reg = is_long ? reg_alloc.AllocLongReg() : reg_alloc.AllocReg();
         } else {
             reg = Register{value};
         }
-- 
cgit v1.2.3


From deda89372f78dc78b37e941bf86e3026708e3ea2 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 10 May 2021 03:47:31 -0300
Subject: glasm: Fix register allocation when moving immediate on GLASM

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 50 +++++++++++++++++-----
 1 file changed, 39 insertions(+), 11 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 8e5d575a9..ad27b8b06 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -39,14 +39,16 @@ struct Identity {
 };
 
 template <bool scalar>
-struct RegWrapper {
-    RegWrapper(EmitContext& ctx, Value value)
-        : reg_alloc{ctx.reg_alloc}, allocated{value.type != Type::Register} {
-        if (allocated) {
+class RegWrapper {
+public:
+    RegWrapper(EmitContext& ctx, const IR::Value& ir_value) : reg_alloc{ctx.reg_alloc} {
+        const Value value{reg_alloc.Peek(ir_value)};
+        if (value.type == Type::Register) {
+            inst = ir_value.InstRecursive();
+            reg = Register{value};
+        } else {
             const bool is_long{value.type == Type::F64 || value.type == Type::U64};
             reg = is_long ? reg_alloc.AllocLongReg() : reg_alloc.AllocReg();
-        } else {
-            reg = Register{value};
         }
         switch (value.type) {
         case Type::Register:
@@ -68,8 +70,11 @@ struct RegWrapper {
             break;
         }
     }
+
     ~RegWrapper() {
-        if (allocated) {
+        if (inst) {
+            reg_alloc.Unref(*inst);
+        } else {
             reg_alloc.FreeReg(reg);
         }
     }
@@ -78,19 +83,42 @@ struct RegWrapper {
         return std::conditional_t<scalar, ScalarRegister, Register>{Value{reg}};
     }
 
+private:
     RegAlloc& reg_alloc;
+    IR::Inst* inst{};
     Register reg{};
-    bool allocated{};
+};
+
+template <typename ArgType>
+class ValueWrapper {
+public:
+    ValueWrapper(EmitContext& ctx, const IR::Value& ir_value_)
+        : reg_alloc{ctx.reg_alloc}, ir_value{ir_value_}, value{reg_alloc.Peek(ir_value)} {}
+
+    ~ValueWrapper() {
+        if (!ir_value.IsImmediate()) {
+            reg_alloc.Unref(*ir_value.InstRecursive());
+        }
+    }
+
+    ArgType Extract() {
+        return value;
+    }
+
+private:
+    RegAlloc& reg_alloc;
+    const IR::Value& ir_value;
+    ArgType value;
 };
 
 template <typename ArgType>
 auto Arg(EmitContext& ctx, const IR::Value& arg) {
     if constexpr (std::is_same_v<ArgType, Register>) {
-        return RegWrapper<false>{ctx, ctx.reg_alloc.Consume(arg)};
+        return RegWrapper<false>{ctx, arg};
     } else if constexpr (std::is_same_v<ArgType, ScalarRegister>) {
-        return RegWrapper<true>{ctx, ctx.reg_alloc.Consume(arg)};
+        return RegWrapper<true>{ctx, arg};
     } else if constexpr (std::is_base_of_v<Value, ArgType>) {
-        return Identity{ArgType{ctx.reg_alloc.Consume(arg)}};
+        return ValueWrapper<ArgType>{ctx, arg};
     } else if constexpr (std::is_same_v<ArgType, const IR::Value&>) {
         return Identity{arg};
     } else if constexpr (std::is_same_v<ArgType, u32>) {
-- 
cgit v1.2.3


From c4fd6b55bc9acd06b2fc89f84fd175d78e14110a Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 10 May 2021 18:21:28 -0300
Subject: glasm: Implement shuffle and vote instructions on GLASM

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 6 ++++++
 1 file changed, 6 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index ad27b8b06..8b42cbf79 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -189,6 +189,12 @@ void SetupOptions(std::string& header, Info info) {
     if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) {
         header += "OPTION NV_shader_atomic_fp16_vector;";
     }
+    if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask) {
+        header += "OPTION NV_shader_thread_group;";
+    }
+    if (info.uses_subgroup_shuffles) {
+        header += "OPTION NV_shader_thread_shuffle;";
+    }
 }
 } // Anonymous namespace
 
-- 
cgit v1.2.3


From c917290497b313abe2f9ad6983050703615b1888 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 10 May 2021 19:20:44 -0300
Subject: glasm: Enable unintentionally disabled register aliasing on GLASM

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 27 +++++++++-------------
 1 file changed, 11 insertions(+), 16 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 8b42cbf79..c90b80e48 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -29,9 +29,9 @@ struct FuncTraits<ReturnType_ (*)(Args...)> {
 
 template <typename T>
 struct Identity {
-    Identity(const T& data_) : data{data_} {}
+    Identity(T data_) : data{data_} {}
 
-    const T& Extract() {
+    T Extract() {
         return data;
     }
 
@@ -71,15 +71,12 @@ public:
         }
     }
 
-    ~RegWrapper() {
+    auto Extract() {
         if (inst) {
             reg_alloc.Unref(*inst);
         } else {
             reg_alloc.FreeReg(reg);
         }
-    }
-
-    auto Extract() {
         return std::conditional_t<scalar, ScalarRegister, Register>{Value{reg}};
     }
 
@@ -95,13 +92,10 @@ public:
     ValueWrapper(EmitContext& ctx, const IR::Value& ir_value_)
         : reg_alloc{ctx.reg_alloc}, ir_value{ir_value_}, value{reg_alloc.Peek(ir_value)} {}
 
-    ~ValueWrapper() {
+    ArgType Extract() {
         if (!ir_value.IsImmediate()) {
             reg_alloc.Unref(*ir_value.InstRecursive());
         }
-    }
-
-    ArgType Extract() {
         return value;
     }
 
@@ -120,7 +114,7 @@ auto Arg(EmitContext& ctx, const IR::Value& arg) {
     } else if constexpr (std::is_base_of_v<Value, ArgType>) {
         return ValueWrapper<ArgType>{ctx, arg};
     } else if constexpr (std::is_same_v<ArgType, const IR::Value&>) {
-        return Identity{arg};
+        return Identity<const IR::Value&>{arg};
     } else if constexpr (std::is_same_v<ArgType, u32>) {
         return Identity{arg.U32()};
     } else if constexpr (std::is_same_v<ArgType, IR::Block*>) {
@@ -137,9 +131,9 @@ auto Arg(EmitContext& ctx, const IR::Value& arg) {
 template <auto func, bool is_first_arg_inst, typename... Args>
 void InvokeCall(EmitContext& ctx, IR::Inst* inst, Args&&... args) {
     if constexpr (is_first_arg_inst) {
-        func(ctx, *inst, std::forward<Args>(args.Extract())...);
+        func(ctx, *inst, args.Extract()...);
     } else {
-        func(ctx, std::forward<Args>(args.Extract())...);
+        func(ctx, args.Extract()...);
     }
 }
 
@@ -147,10 +141,11 @@ template <auto func, bool is_first_arg_inst, size_t... I>
 void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) {
     using Traits = FuncTraits<decltype(func)>;
     if constexpr (is_first_arg_inst) {
-        func(ctx, *inst,
-             Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I)).Extract()...);
+        InvokeCall<func, is_first_arg_inst>(
+            ctx, inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
     } else {
-        func(ctx, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I)).Extract()...);
+        InvokeCall<func, is_first_arg_inst>(
+            ctx, inst, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...);
     }
 }
 
-- 
cgit v1.2.3


From 8c81a20ace8c65d0a9d58b9cf333049a2bc0383a Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 10 May 2021 19:32:10 -0300
Subject: glasm: Ensure reg alloc order across compilers on GLASM

Use a struct constructor to serialize register allocation arguments to
ensure registers are allocated in the same order regardless of the
compiler used.

The A and B functions can be called in any order when passed as
arguments to "foo":

  foo(A(), B())

But the order is guaranteed for curly-braced constructor calls in
classes:

  Foo{A(), B()}

Use this to get consistent behavior.
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 25 ++++++++++++----------
 1 file changed, 14 insertions(+), 11 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index c90b80e48..047b2f89c 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -128,24 +128,27 @@ auto Arg(EmitContext& ctx, const IR::Value& arg) {
     }
 }
 
-template <auto func, bool is_first_arg_inst, typename... Args>
-void InvokeCall(EmitContext& ctx, IR::Inst* inst, Args&&... args) {
-    if constexpr (is_first_arg_inst) {
-        func(ctx, *inst, args.Extract()...);
-    } else {
-        func(ctx, args.Extract()...);
+template <auto func, bool is_first_arg_inst>
+struct InvokeCall {
+    template <typename... Args>
+    InvokeCall(EmitContext& ctx, IR::Inst* inst, Args&&... args) {
+        if constexpr (is_first_arg_inst) {
+            func(ctx, *inst, args.Extract()...);
+        } else {
+            func(ctx, args.Extract()...);
+        }
     }
-}
+};
 
 template <auto func, bool is_first_arg_inst, size_t... I>
 void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) {
     using Traits = FuncTraits<decltype(func)>;
     if constexpr (is_first_arg_inst) {
-        InvokeCall<func, is_first_arg_inst>(
-            ctx, inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
+        InvokeCall<func, is_first_arg_inst>{
+            ctx, inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...};
     } else {
-        InvokeCall<func, is_first_arg_inst>(
-            ctx, inst, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...);
+        InvokeCall<func, is_first_arg_inst>{
+            ctx, inst, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...};
     }
 }
 
-- 
cgit v1.2.3


From 7ff5851608031baca2adceb9f72e7c75eda9b3a9 Mon Sep 17 00:00:00 2001
From: ameerj <52414509+ameerj@users.noreply.github.com>
Date: Mon, 10 May 2021 22:35:16 -0400
Subject: glasm: Implement Storage atomics

StorageAtomicExchangeU64 is failing test seemingly due to failure storing 64-bit
result into the register
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 7 +++++++
 1 file changed, 7 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 047b2f89c..056d8cbf8 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -193,6 +193,9 @@ void SetupOptions(std::string& header, Info info) {
     if (info.uses_subgroup_shuffles) {
         header += "OPTION NV_shader_thread_shuffle;";
     }
+    // TODO: Track the shared atomic ops
+    header +=
+        "OPTION NV_shader_storage_buffer;OPTION NV_gpu_program_fp64;OPTION NV_bindless_texture;";
 }
 } // Anonymous namespace
 
@@ -214,6 +217,10 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     default:
         break;
     }
+    if (program.shared_memory_size > 0) {
+        header += fmt::format("SHARED_MEMORY {};", program.shared_memory_size);
+        header += fmt::format("SHARED shared_mem[]={{program.sharedmem}};");
+    }
     header += "TEMP ";
     for (size_t index = 0; index < ctx.reg_alloc.NumUsedRegisters(); ++index) {
         header += fmt::format("R{},", index);
-- 
cgit v1.2.3


From d54d7de40e7295827b0e4e4026441b53d3fc9569 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Fri, 14 May 2021 00:40:54 -0300
Subject: glasm: Rework control flow introducing a syntax list

This commit regresses VertexA shaders, their transformation pass has to
be adapted to the new control flow.
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 41 ++++++++++++++++++----
 1 file changed, 34 insertions(+), 7 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 056d8cbf8..51ca83d18 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -117,8 +117,6 @@ auto Arg(EmitContext& ctx, const IR::Value& arg) {
         return Identity<const IR::Value&>{arg};
     } else if constexpr (std::is_same_v<ArgType, u32>) {
         return Identity{arg.U32()};
-    } else if constexpr (std::is_same_v<ArgType, IR::Block*>) {
-        return Identity{arg.Label()};
     } else if constexpr (std::is_same_v<ArgType, IR::Attribute>) {
         return Identity{arg.Attribute()};
     } else if constexpr (std::is_same_v<ArgType, IR::Patch>) {
@@ -177,6 +175,39 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) {
     throw LogicError("Invalid opcode {}", inst->GetOpcode());
 }
 
+void EmitCode(EmitContext& ctx, const IR::Program& program) {
+    const auto eval{
+        [&](const IR::U1& cond) { return ScalarS32{ctx.reg_alloc.Consume(IR::Value{cond})}; }};
+    for (const IR::AbstractSyntaxNode& node : program.syntax_list) {
+        switch (node.type) {
+        case IR::AbstractSyntaxNode::Type::Block:
+            for (IR::Inst& inst : node.block->Instructions()) {
+                EmitInst(ctx, &inst);
+            }
+            break;
+        case IR::AbstractSyntaxNode::Type::If:
+            ctx.Add("MOV.S.CC RC,{};IF NE.x;", eval(node.if_node.cond));
+            break;
+        case IR::AbstractSyntaxNode::Type::EndIf:
+            ctx.Add("ENDIF;");
+            break;
+        case IR::AbstractSyntaxNode::Type::Loop:
+            ctx.Add("REP;");
+            break;
+        case IR::AbstractSyntaxNode::Type::Repeat:
+            ctx.Add("MOV.S.CC RC,{};BRK NE.x;ENDREP;", eval(node.repeat.cond));
+            break;
+        case IR::AbstractSyntaxNode::Type::Break:
+            ctx.Add("MOV.S.CC RC,{};BRK NE.x;", eval(node.repeat.cond));
+            break;
+        case IR::AbstractSyntaxNode::Type::Return:
+        case IR::AbstractSyntaxNode::Type::Unreachable:
+            ctx.Add("RET;");
+            break;
+        }
+    }
+}
+
 void SetupOptions(std::string& header, Info info) {
     if (info.uses_int64_bit_atomics) {
         header += "OPTION NV_shader_atomic_int64;";
@@ -201,11 +232,7 @@ void SetupOptions(std::string& header, Info info) {
 
 std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     EmitContext ctx{program};
-    for (IR::Block* const block : program.blocks) {
-        for (IR::Inst& inst : block->Instructions()) {
-            EmitInst(ctx, &inst);
-        }
-    }
+    EmitCode(ctx, program);
     std::string header = "!!NVcp5.0\n"
                          "OPTION NV_internal;";
     SetupOptions(header, program.info);
-- 
cgit v1.2.3


From d4385c34e3aee6718502a1c5bc814535a657dc4f Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Fri, 14 May 2021 02:09:33 -0300
Subject: glasm: Declare NV_shader_thread_group when needed

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 51ca83d18..fa48ba25c 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -218,15 +218,16 @@ void SetupOptions(std::string& header, Info info) {
     if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) {
         header += "OPTION NV_shader_atomic_fp16_vector;";
     }
-    if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask) {
+    if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote) {
         header += "OPTION NV_shader_thread_group;";
     }
     if (info.uses_subgroup_shuffles) {
         header += "OPTION NV_shader_thread_shuffle;";
     }
     // TODO: Track the shared atomic ops
-    header +=
-        "OPTION NV_shader_storage_buffer;OPTION NV_gpu_program_fp64;OPTION NV_bindless_texture;";
+    header += "OPTION NV_shader_storage_buffer;"
+              "OPTION NV_gpu_program_fp64;"
+              "OPTION NV_bindless_texture;";
 }
 } // Anonymous namespace
 
-- 
cgit v1.2.3


From bf5e48ffe4bd48ea681f2a01c8919c97125e88df Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Fri, 14 May 2021 04:48:46 -0300
Subject: glasm: Initial implementation of phi nodes on GLASM

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 59 ++++++++++++++++++++--
 1 file changed, 56 insertions(+), 3 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index fa48ba25c..775dd9e7e 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -2,6 +2,7 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
+#include <ranges>
 #include <string>
 #include <tuple>
 
@@ -9,6 +10,7 @@
 #include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/frontend/ir/ir_emitter.h"
 #include "shader_recompiler/frontend/ir/program.h"
 #include "shader_recompiler/profile.h"
 
@@ -175,6 +177,34 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) {
     throw LogicError("Invalid opcode {}", inst->GetOpcode());
 }
 
+void Precolor(EmitContext& ctx, const IR::Program& program) {
+    for (IR::Block* const block : program.blocks) {
+        for (IR::Inst& phi : block->Instructions() | std::views::take_while(IR::IsPhi)) {
+            switch (phi.Arg(0).Type()) {
+            case IR::Type::U1:
+            case IR::Type::U32:
+            case IR::Type::F32:
+                ctx.reg_alloc.Define(phi);
+                break;
+            case IR::Type::U64:
+            case IR::Type::F64:
+                ctx.reg_alloc.LongDefine(phi);
+                break;
+            default:
+                throw NotImplementedException("Phi node type {}", phi.Type());
+            }
+            const size_t num_args{phi.NumArgs()};
+            for (size_t i = 0; i < num_args; ++i) {
+                IR::IREmitter{*phi.PhiBlock(i)}.PhiMove(phi, phi.Arg(i));
+            }
+            // Add reference to the phi node on the phi predecessor to avoid overwritting it
+            for (size_t i = 0; i < num_args; ++i) {
+                IR::IREmitter{*phi.PhiBlock(i)}.DummyReference(IR::Value{&phi});
+            }
+        }
+    }
+}
+
 void EmitCode(EmitContext& ctx, const IR::Program& program) {
     const auto eval{
         [&](const IR::U1& cond) { return ScalarS32{ctx.reg_alloc.Consume(IR::Value{cond})}; }};
@@ -186,7 +216,9 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
             }
             break;
         case IR::AbstractSyntaxNode::Type::If:
-            ctx.Add("MOV.S.CC RC,{};IF NE.x;", eval(node.if_node.cond));
+            ctx.Add("MOV.S.CC RC,{};"
+                    "IF NE.x;",
+                    eval(node.if_node.cond));
             break;
         case IR::AbstractSyntaxNode::Type::EndIf:
             ctx.Add("ENDIF;");
@@ -195,10 +227,30 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
             ctx.Add("REP;");
             break;
         case IR::AbstractSyntaxNode::Type::Repeat:
-            ctx.Add("MOV.S.CC RC,{};BRK NE.x;ENDREP;", eval(node.repeat.cond));
+            if (node.repeat.cond.IsImmediate()) {
+                if (node.repeat.cond.U1()) {
+                    ctx.Add("ENDREP;");
+                } else {
+                    ctx.Add("BRK;"
+                            "ENDREP;");
+                }
+            } else {
+                ctx.Add("MOV.S.CC RC,{};"
+                        "BRK (EQ.x);"
+                        "ENDREP;",
+                        eval(node.repeat.cond));
+            }
             break;
         case IR::AbstractSyntaxNode::Type::Break:
-            ctx.Add("MOV.S.CC RC,{};BRK NE.x;", eval(node.repeat.cond));
+            if (node.break_node.cond.IsImmediate()) {
+                if (node.break_node.cond.U1()) {
+                    ctx.Add("BRK;");
+                }
+            } else {
+                ctx.Add("MOV.S.CC RC,{};"
+                        "BRK (NE.x);",
+                        eval(node.break_node.cond));
+            }
             break;
         case IR::AbstractSyntaxNode::Type::Return:
         case IR::AbstractSyntaxNode::Type::Unreachable:
@@ -233,6 +285,7 @@ void SetupOptions(std::string& header, Info info) {
 
 std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     EmitContext ctx{program};
+    Precolor(ctx, program);
     EmitCode(ctx, program);
     std::string header = "!!NVcp5.0\n"
                          "OPTION NV_internal;";
-- 
cgit v1.2.3


From 057dee48562b0cce69b1fa8bdb02bc0367852b4d Mon Sep 17 00:00:00 2001
From: ameerj <52414509+ameerj@users.noreply.github.com>
Date: Fri, 14 May 2021 21:18:53 -0400
Subject: glasm: Implement local memory for glasm

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 +++
 1 file changed, 3 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 775dd9e7e..0b70bf3f6 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -306,6 +306,9 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     for (size_t index = 0; index < ctx.reg_alloc.NumUsedRegisters(); ++index) {
         header += fmt::format("R{},", index);
     }
+    if (program.local_memory_size > 0) {
+        header += fmt::format("lmem[{}],", program.local_memory_size);
+    }
     header += "RC;"
               "LONG TEMP ";
     for (size_t index = 0; index < ctx.reg_alloc.NumUsedLongRegisters(); ++index) {
-- 
cgit v1.2.3


From 3764750339fa60f2d79cf3abe1b91ca42ba61401 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sat, 15 May 2021 18:14:29 -0300
Subject: glasm: Add graphics specific shader declarations to GLASM

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 32 ++++++++++++++++++----
 1 file changed, 26 insertions(+), 6 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 0b70bf3f6..ab6790ce8 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -261,6 +261,12 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
 }
 
 void SetupOptions(std::string& header, Info info) {
+    // TODO: Track the shared atomic ops
+    header += "OPTION NV_internal;"
+              "OPTION NV_shader_storage_buffer;"
+              "OPTION NV_gpu_program_fp64;"
+              "OPTION NV_bindless_texture;"
+              "OPTION ARB_derivative_control;";
     if (info.uses_int64_bit_atomics) {
         header += "OPTION NV_shader_atomic_int64;";
     }
@@ -276,10 +282,25 @@ void SetupOptions(std::string& header, Info info) {
     if (info.uses_subgroup_shuffles) {
         header += "OPTION NV_shader_thread_shuffle;";
     }
-    // TODO: Track the shared atomic ops
-    header += "OPTION NV_shader_storage_buffer;"
-              "OPTION NV_gpu_program_fp64;"
-              "OPTION NV_bindless_texture;";
+}
+
+std::string_view StageHeader(Stage stage) {
+    switch (stage) {
+    case Stage::VertexA:
+    case Stage::VertexB:
+        return "!!NVvp5.0\n";
+    case Stage::TessellationControl:
+        return "!!NVtcs5.0\n";
+    case Stage::TessellationEval:
+        return "!!NVtes5.0\n";
+    case Stage::Geometry:
+        return "!!NVgp5.0\n";
+    case Stage::Fragment:
+        return "!!NVfp5.0\n";
+    case Stage::Compute:
+        return "!!NVcp5.0\n";
+    }
+    throw InvalidArgument("Invalid stage {}", stage);
 }
 } // Anonymous namespace
 
@@ -287,8 +308,7 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
     EmitContext ctx{program};
     Precolor(ctx, program);
     EmitCode(ctx, program);
-    std::string header = "!!NVcp5.0\n"
-                         "OPTION NV_internal;";
+    std::string header{StageHeader(program.stage)};
     SetupOptions(header, program.info);
     switch (program.stage) {
     case Stage::Compute:
-- 
cgit v1.2.3


From 31d402ee74d7f7045aec7e748fdee489a434db6b Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sat, 15 May 2021 18:15:13 -0300
Subject: glasm: Add Void type to GLASM values

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 1 +
 1 file changed, 1 insertion(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index ab6790ce8..e5c96eb7f 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -54,6 +54,7 @@ public:
         }
         switch (value.type) {
         case Type::Register:
+        case Type::Void:
             break;
         case Type::U32:
             ctx.Add("MOV.U {}.x,{};", reg, value.imm_u32);
-- 
cgit v1.2.3


From f7a2340205b4fa2db32403f20d7b7afe32b15f33 Mon Sep 17 00:00:00 2001
From: lat9nq <22451773+lat9nq@users.noreply.github.com>
Date: Sun, 16 May 2021 17:06:13 -0400
Subject: shader_recompiler: GCC fixes

Fixes members of unnamed union not being accessible, and one function
without a declaration.
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 16 ++++++++--------
 1 file changed, 8 insertions(+), 8 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index e5c96eb7f..0a76423f4 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -212,14 +212,14 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
     for (const IR::AbstractSyntaxNode& node : program.syntax_list) {
         switch (node.type) {
         case IR::AbstractSyntaxNode::Type::Block:
-            for (IR::Inst& inst : node.block->Instructions()) {
+            for (IR::Inst& inst : node.data.block->Instructions()) {
                 EmitInst(ctx, &inst);
             }
             break;
         case IR::AbstractSyntaxNode::Type::If:
             ctx.Add("MOV.S.CC RC,{};"
                     "IF NE.x;",
-                    eval(node.if_node.cond));
+                    eval(node.data.if_node.cond));
             break;
         case IR::AbstractSyntaxNode::Type::EndIf:
             ctx.Add("ENDIF;");
@@ -228,8 +228,8 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
             ctx.Add("REP;");
             break;
         case IR::AbstractSyntaxNode::Type::Repeat:
-            if (node.repeat.cond.IsImmediate()) {
-                if (node.repeat.cond.U1()) {
+            if (node.data.repeat.cond.IsImmediate()) {
+                if (node.data.repeat.cond.U1()) {
                     ctx.Add("ENDREP;");
                 } else {
                     ctx.Add("BRK;"
@@ -239,18 +239,18 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
                 ctx.Add("MOV.S.CC RC,{};"
                         "BRK (EQ.x);"
                         "ENDREP;",
-                        eval(node.repeat.cond));
+                        eval(node.data.repeat.cond));
             }
             break;
         case IR::AbstractSyntaxNode::Type::Break:
-            if (node.break_node.cond.IsImmediate()) {
-                if (node.break_node.cond.U1()) {
+            if (node.data.break_node.cond.IsImmediate()) {
+                if (node.data.break_node.cond.U1()) {
                     ctx.Add("BRK;");
                 }
             } else {
                 ctx.Add("MOV.S.CC RC,{};"
                         "BRK (NE.x);",
-                        eval(node.break_node.cond));
+                        eval(node.data.break_node.cond));
             }
             break;
         case IR::AbstractSyntaxNode::Type::Return:
-- 
cgit v1.2.3


From db2f0f410810d3d8310a6a476a8bcfd5e509869e Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sun, 16 May 2021 17:52:30 -0300
Subject: emit_glasm: Enable ARB_draw_buffers when needed

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 4 ++++
 1 file changed, 4 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 0a76423f4..0c591f73c 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -283,6 +283,10 @@ void SetupOptions(std::string& header, Info info) {
     if (info.uses_subgroup_shuffles) {
         header += "OPTION NV_shader_thread_shuffle;";
     }
+    const auto non_zero_frag_colors{info.stores_frag_color | std::views::drop(1)};
+    if (std::ranges::find(non_zero_frag_colors, true) != non_zero_frag_colors.end()) {
+        header += "OPTION ARB_draw_buffers;";
+    }
 }
 
 std::string_view StageHeader(Stage stage) {
-- 
cgit v1.2.3


From ec6fc5fe78c9038fc9ad7259b7b3a7be751ecef6 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 17 May 2021 02:52:01 -0300
Subject: glasm: Implement TEX and TEXS instructions

Remove lod clamp from texture instructions with lod, as this is not
needed (nor supported).
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 +++
 1 file changed, 3 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 0c591f73c..d7a08e4b3 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -283,6 +283,9 @@ void SetupOptions(std::string& header, Info info) {
     if (info.uses_subgroup_shuffles) {
         header += "OPTION NV_shader_thread_shuffle;";
     }
+    if (info.uses_sparse_residency) {
+        header += "OPTION EXT_sparse_texture2;";
+    }
     const auto non_zero_frag_colors{info.stores_frag_color | std::views::drop(1)};
     if (std::ranges::find(non_zero_frag_colors, true) != non_zero_frag_colors.end()) {
         header += "OPTION ARB_draw_buffers;";
-- 
cgit v1.2.3


From 9bb3e008c9f4bbdd35c095b506c3a3312d17e383 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Tue, 18 May 2021 02:04:22 -0300
Subject: shader: Read branch conditions from an instruction

Fixes the identity removal pass.
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index d7a08e4b3..a893fa3fb 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -200,7 +200,7 @@ void Precolor(EmitContext& ctx, const IR::Program& program) {
             }
             // Add reference to the phi node on the phi predecessor to avoid overwritting it
             for (size_t i = 0; i < num_args; ++i) {
-                IR::IREmitter{*phi.PhiBlock(i)}.DummyReference(IR::Value{&phi});
+                IR::IREmitter{*phi.PhiBlock(i)}.Reference(IR::Value{&phi});
             }
         }
     }
-- 
cgit v1.2.3


From 8b7d5912d61d56f65fb7e3a03bba544a4c40bfa6 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Tue, 18 May 2021 21:04:09 -0300
Subject: glasm: Support textures used in more than one stage

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index a893fa3fb..edf6f5e13 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -312,8 +312,8 @@ std::string_view StageHeader(Stage stage) {
 }
 } // Anonymous namespace
 
-std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) {
-    EmitContext ctx{program};
+std::string EmitGLASM(const Profile&, IR::Program& program, Bindings& bindings) {
+    EmitContext ctx{program, bindings};
     Precolor(ctx, program);
     EmitCode(ctx, program);
     std::string header{StageHeader(program.stage)};
-- 
cgit v1.2.3


From accad56ee7cc9f77886d164701a35f1e89a3519b Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Wed, 19 May 2021 16:32:03 -0300
Subject: glasm: Implement stores to gl_ViewportIndex

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 16 ++++++++++++----
 1 file changed, 12 insertions(+), 4 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index edf6f5e13..9dc0cacbe 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -261,7 +261,10 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
     }
 }
 
-void SetupOptions(std::string& header, Info info) {
+void SetupOptions(const IR::Program& program, const Profile& profile, std::string& header) {
+    const Info& info{program.info};
+    const Stage stage{program.stage};
+
     // TODO: Track the shared atomic ops
     header += "OPTION NV_internal;"
               "OPTION NV_shader_storage_buffer;"
@@ -286,6 +289,11 @@ void SetupOptions(std::string& header, Info info) {
     if (info.uses_sparse_residency) {
         header += "OPTION EXT_sparse_texture2;";
     }
+    if ((info.stores_viewport_index || info.stores_layer) && stage != Stage::Geometry) {
+        if (profile.support_viewport_index_layer_non_geometry) {
+            header += "OPTION NV_viewport_array2;";
+        }
+    }
     const auto non_zero_frag_colors{info.stores_frag_color | std::views::drop(1)};
     if (std::ranges::find(non_zero_frag_colors, true) != non_zero_frag_colors.end()) {
         header += "OPTION ARB_draw_buffers;";
@@ -312,12 +320,12 @@ std::string_view StageHeader(Stage stage) {
 }
 } // Anonymous namespace
 
-std::string EmitGLASM(const Profile&, IR::Program& program, Bindings& bindings) {
-    EmitContext ctx{program, bindings};
+std::string EmitGLASM(const Profile& profile, IR::Program& program, Bindings& bindings) {
+    EmitContext ctx{program, bindings, profile};
     Precolor(ctx, program);
     EmitCode(ctx, program);
     std::string header{StageHeader(program.stage)};
-    SetupOptions(header, program.info);
+    SetupOptions(program, profile, header);
     switch (program.stage) {
     case Stage::Compute:
         header += fmt::format("GROUP_SIZE {} {} {};", program.workgroup_size[0],
-- 
cgit v1.2.3


From dadd192b30d547dfa078057796a5ae16820eb4be Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Thu, 20 May 2021 02:18:52 -0300
Subject: glasm: Implement ImageRead

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 9dc0cacbe..3910d00ee 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -270,7 +270,8 @@ void SetupOptions(const IR::Program& program, const Profile& profile, std::strin
               "OPTION NV_shader_storage_buffer;"
               "OPTION NV_gpu_program_fp64;"
               "OPTION NV_bindless_texture;"
-              "OPTION ARB_derivative_control;";
+              "OPTION ARB_derivative_control;"
+              "OPTION EXT_shader_image_load_formatted;";
     if (info.uses_int64_bit_atomics) {
         header += "OPTION NV_shader_atomic_int64;";
     }
-- 
cgit v1.2.3


From fad139a3e6f8273acb5b14296ba8fcbd0946fe76 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Thu, 20 May 2021 17:27:39 -0300
Subject: glasm: Declare geometry program headers

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 35 ++++++++++++++++++++++
 1 file changed, 35 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 3910d00ee..b6b8d504e 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -319,6 +319,34 @@ std::string_view StageHeader(Stage stage) {
     }
     throw InvalidArgument("Invalid stage {}", stage);
 }
+
+std::string_view InputPrimitive(InputTopology topology) {
+    switch (topology) {
+    case InputTopology::Points:
+        return "POINTS";
+    case InputTopology::Lines:
+        return "LINES";
+    case InputTopology::LinesAdjacency:
+        return "LINESS_ADJACENCY";
+    case InputTopology::Triangles:
+        return "TRIANGLES";
+    case InputTopology::TrianglesAdjacency:
+        return "TRIANGLES_ADJACENCY";
+    }
+    throw InvalidArgument("Invalid input topology {}", topology);
+}
+
+std::string_view OutputPrimitive(OutputTopology topology) {
+    switch (topology) {
+    case OutputTopology::PointList:
+        return "POINTS";
+    case OutputTopology::LineStrip:
+        return "LINE_STRIP";
+    case OutputTopology::TriangleStrip:
+        return "TRIANGLE_STRIP";
+    }
+    throw InvalidArgument("Invalid output topology {}", topology);
+}
 } // Anonymous namespace
 
 std::string EmitGLASM(const Profile& profile, IR::Program& program, Bindings& bindings) {
@@ -328,6 +356,13 @@ std::string EmitGLASM(const Profile& profile, IR::Program& program, Bindings& bi
     std::string header{StageHeader(program.stage)};
     SetupOptions(program, profile, header);
     switch (program.stage) {
+    case Stage::Geometry:
+        header += fmt::format("PRIMITIVE_IN {};"
+                              "PRIMITIVE_OUT {};"
+                              "VERTICES_OUT {};",
+                              InputPrimitive(profile.input_topology),
+                              OutputPrimitive(program.output_topology), program.output_vertices);
+        break;
     case Stage::Compute:
         header += fmt::format("GROUP_SIZE {} {} {};", program.workgroup_size[0],
                               program.workgroup_size[1], program.workgroup_size[2]);
-- 
cgit v1.2.3


From 9ec2303ad6a399cea9e66fa522f65671046f1879 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Thu, 20 May 2021 21:18:39 -0300
Subject: glasm: Add tessellation shader declarations

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 35 ++++++++++++++++++++++
 1 file changed, 35 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index b6b8d504e..476cdda54 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -347,6 +347,30 @@ std::string_view OutputPrimitive(OutputTopology topology) {
     }
     throw InvalidArgument("Invalid output topology {}", topology);
 }
+
+std::string_view GetTessMode(TessPrimitive primitive) {
+    switch (primitive) {
+    case TessPrimitive::Triangles:
+        return "TRIANGLES";
+    case TessPrimitive::Quads:
+        return "QUADS";
+    case TessPrimitive::Isolines:
+        return "ISOLINES";
+    }
+    throw InvalidArgument("Invalid tessellation primitive {}", primitive);
+}
+
+std::string_view GetTessSpacing(TessSpacing spacing) {
+    switch (spacing) {
+    case TessSpacing::Equal:
+        return "EQUAL";
+    case TessSpacing::FractionalOdd:
+        return "FRACTIONAL_ODD";
+    case TessSpacing::FractionalEven:
+        return "FRACTIONAL_EVEN";
+    }
+    throw InvalidArgument("Invalid tessellation spacing {}", spacing);
+}
 } // Anonymous namespace
 
 std::string EmitGLASM(const Profile& profile, IR::Program& program, Bindings& bindings) {
@@ -356,6 +380,17 @@ std::string EmitGLASM(const Profile& profile, IR::Program& program, Bindings& bi
     std::string header{StageHeader(program.stage)};
     SetupOptions(program, profile, header);
     switch (program.stage) {
+    case Stage::TessellationControl:
+        header += fmt::format("VERTICES_OUT {};", program.invocations);
+        break;
+    case Stage::TessellationEval:
+        header +=
+            fmt::format("TESS_MODE {};"
+                        "TESS_SPACING {};"
+                        "TESS_VERTEX_ORDER {};",
+                        GetTessMode(profile.tess_primitive), GetTessSpacing(profile.tess_spacing),
+                        profile.tess_clockwise ? "CW" : "CCW");
+        break;
     case Stage::Geometry:
         header += fmt::format("PRIMITIVE_IN {};"
                               "PRIMITIVE_OUT {};"
-- 
cgit v1.2.3


From 48d4e263264e9ae0214ad6f0064e8e32aba17fc4 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Thu, 20 May 2021 21:39:45 -0300
Subject: glasm: Fix tessellation headers

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 476cdda54..4fc7d2f2f 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -307,9 +307,9 @@ std::string_view StageHeader(Stage stage) {
     case Stage::VertexB:
         return "!!NVvp5.0\n";
     case Stage::TessellationControl:
-        return "!!NVtcs5.0\n";
+        return "!!NVtcp5.0\n";
     case Stage::TessellationEval:
-        return "!!NVtes5.0\n";
+        return "!!NVtep5.0\n";
     case Stage::Geometry:
         return "!!NVgp5.0\n";
     case Stage::Fragment:
-- 
cgit v1.2.3


From 36d040da7059e438fa35f1a5de5d5aed4cef5ca4 Mon Sep 17 00:00:00 2001
From: ameerj <52414509+ameerj@users.noreply.github.com>
Date: Wed, 19 May 2021 01:00:51 -0400
Subject: glasm: Implement FSWZADD

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 16 +++++++++++++++-
 1 file changed, 15 insertions(+), 1 deletion(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 4fc7d2f2f..f110fd7f8 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -281,7 +281,8 @@ void SetupOptions(const IR::Program& program, const Profile& profile, std::strin
     if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) {
         header += "OPTION NV_shader_atomic_fp16_vector;";
     }
-    if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote) {
+    if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote ||
+        info.uses_fswzadd) {
         header += "OPTION NV_shader_thread_group;";
     }
     if (info.uses_subgroup_shuffles) {
@@ -416,12 +417,25 @@ std::string EmitGLASM(const Profile& profile, IR::Program& program, Bindings& bi
     if (program.local_memory_size > 0) {
         header += fmt::format("lmem[{}],", program.local_memory_size);
     }
+    if (program.info.uses_fswzadd) {
+        header += "FSWZA[4],FSWZB[4],";
+    }
     header += "RC;"
               "LONG TEMP ";
     for (size_t index = 0; index < ctx.reg_alloc.NumUsedLongRegisters(); ++index) {
         header += fmt::format("D{},", index);
     }
     header += "DC;";
+    if (program.info.uses_fswzadd) {
+        header += "MOV.F FSWZA[0],-1;"
+                  "MOV.F FSWZA[1],1;"
+                  "MOV.F FSWZA[2],-1;"
+                  "MOV.F FSWZA[3],0;"
+                  "MOV.F FSWZB[0],-1;"
+                  "MOV.F FSWZB[1],-1;"
+                  "MOV.F FSWZB[2],1;"
+                  "MOV.F FSWZB[3],-1;";
+    }
     ctx.code.insert(0, header);
     ctx.code += "END";
     return ctx.code;
-- 
cgit v1.2.3


From 9e7b6622c25aa858b96bf0f1c7f94223a2f449a2 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Fri, 21 May 2021 02:12:32 -0300
Subject: shader: Split profile and runtime information in separate structs

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 19 ++++++++++---------
 1 file changed, 10 insertions(+), 9 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index f110fd7f8..edff04a44 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -374,8 +374,9 @@ std::string_view GetTessSpacing(TessSpacing spacing) {
 }
 } // Anonymous namespace
 
-std::string EmitGLASM(const Profile& profile, IR::Program& program, Bindings& bindings) {
-    EmitContext ctx{program, bindings, profile};
+std::string EmitGLASM(const Profile& profile, const RuntimeInfo& runtime_info, IR::Program& program,
+                      Bindings& bindings) {
+    EmitContext ctx{program, bindings, profile, runtime_info};
     Precolor(ctx, program);
     EmitCode(ctx, program);
     std::string header{StageHeader(program.stage)};
@@ -385,18 +386,18 @@ std::string EmitGLASM(const Profile& profile, IR::Program& program, Bindings& bi
         header += fmt::format("VERTICES_OUT {};", program.invocations);
         break;
     case Stage::TessellationEval:
-        header +=
-            fmt::format("TESS_MODE {};"
-                        "TESS_SPACING {};"
-                        "TESS_VERTEX_ORDER {};",
-                        GetTessMode(profile.tess_primitive), GetTessSpacing(profile.tess_spacing),
-                        profile.tess_clockwise ? "CW" : "CCW");
+        header += fmt::format("TESS_MODE {};"
+                              "TESS_SPACING {};"
+                              "TESS_VERTEX_ORDER {};",
+                              GetTessMode(runtime_info.tess_primitive),
+                              GetTessSpacing(runtime_info.tess_spacing),
+                              runtime_info.tess_clockwise ? "CW" : "CCW");
         break;
     case Stage::Geometry:
         header += fmt::format("PRIMITIVE_IN {};"
                               "PRIMITIVE_OUT {};"
                               "VERTICES_OUT {};",
-                              InputPrimitive(profile.input_topology),
+                              InputPrimitive(runtime_info.input_topology),
                               OutputPrimitive(program.output_topology), program.output_vertices);
         break;
     case Stage::Compute:
-- 
cgit v1.2.3


From 84feabac881443d27f84f8fec5eba6dc3b13b620 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Fri, 21 May 2021 18:27:37 -0300
Subject: glasm: Implement forced early Z

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 8 ++++++--
 1 file changed, 6 insertions(+), 2 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index edff04a44..0c2bbf284 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -261,7 +261,8 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
     }
 }
 
-void SetupOptions(const IR::Program& program, const Profile& profile, std::string& header) {
+void SetupOptions(const IR::Program& program, const Profile& profile,
+                  const RuntimeInfo& runtime_info, std::string& header) {
     const Info& info{program.info};
     const Stage stage{program.stage};
 
@@ -296,6 +297,9 @@ void SetupOptions(const IR::Program& program, const Profile& profile, std::strin
             header += "OPTION NV_viewport_array2;";
         }
     }
+    if (stage == Stage::Fragment && runtime_info.force_early_z != 0) {
+        header += "OPTION NV_early_fragment_tests;";
+    }
     const auto non_zero_frag_colors{info.stores_frag_color | std::views::drop(1)};
     if (std::ranges::find(non_zero_frag_colors, true) != non_zero_frag_colors.end()) {
         header += "OPTION ARB_draw_buffers;";
@@ -380,7 +384,7 @@ std::string EmitGLASM(const Profile& profile, const RuntimeInfo& runtime_info, I
     Precolor(ctx, program);
     EmitCode(ctx, program);
     std::string header{StageHeader(program.stage)};
-    SetupOptions(program, profile, header);
+    SetupOptions(program, profile, runtime_info, header);
     switch (program.stage) {
     case Stage::TessellationControl:
         header += fmt::format("VERTICES_OUT {};", program.invocations);
-- 
cgit v1.2.3


From c31521512fd49603ea42c93e2a6eac5d7985cd78 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Fri, 21 May 2021 20:46:40 -0300
Subject: gl_shader_cache,glasm: Conditionally use typeless image reads
 extension

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 0c2bbf284..8718cc7ec 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -271,8 +271,7 @@ void SetupOptions(const IR::Program& program, const Profile& profile,
               "OPTION NV_shader_storage_buffer;"
               "OPTION NV_gpu_program_fp64;"
               "OPTION NV_bindless_texture;"
-              "OPTION ARB_derivative_control;"
-              "OPTION EXT_shader_image_load_formatted;";
+              "OPTION ARB_derivative_control;";
     if (info.uses_int64_bit_atomics) {
         header += "OPTION NV_shader_atomic_int64;";
     }
@@ -297,6 +296,9 @@ void SetupOptions(const IR::Program& program, const Profile& profile,
             header += "OPTION NV_viewport_array2;";
         }
     }
+    if (info.uses_typeless_image_reads && profile.support_typeless_image_loads) {
+        header += "OPTION EXT_shader_image_load_formatted;";
+    }
     if (stage == Stage::Fragment && runtime_info.force_early_z != 0) {
         header += "OPTION NV_early_fragment_tests;";
     }
-- 
cgit v1.2.3


From 9fbfe7d676790dea160368eda6492e8feb6e2f4a Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 24 May 2021 20:59:49 -0300
Subject: glasm: Fix usage counting on phi nodes

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 12 +++++++++---
 1 file changed, 9 insertions(+), 3 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 8718cc7ec..2ce839059 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -178,6 +178,10 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) {
     throw LogicError("Invalid opcode {}", inst->GetOpcode());
 }
 
+bool IsReference(IR::Inst& inst) {
+    return inst.GetOpcode() == IR::Opcode::Reference;
+}
+
 void Precolor(EmitContext& ctx, const IR::Program& program) {
     for (IR::Block* const block : program.blocks) {
         for (IR::Inst& phi : block->Instructions() | std::views::take_while(IR::IsPhi)) {
@@ -194,11 +198,13 @@ void Precolor(EmitContext& ctx, const IR::Program& program) {
             default:
                 throw NotImplementedException("Phi node type {}", phi.Type());
             }
+            // Insert phi moves before references to avoid overwritting them
             const size_t num_args{phi.NumArgs()};
             for (size_t i = 0; i < num_args; ++i) {
-                IR::IREmitter{*phi.PhiBlock(i)}.PhiMove(phi, phi.Arg(i));
+                IR::Block& phi_block{*phi.PhiBlock(i)};
+                auto it{std::find_if_not(phi_block.rbegin(), phi_block.rend(), IsReference).base()};
+                IR::IREmitter{phi_block, it}.PhiMove(phi, phi.Arg(i));
             }
-            // Add reference to the phi node on the phi predecessor to avoid overwritting it
             for (size_t i = 0; i < num_args; ++i) {
                 IR::IREmitter{*phi.PhiBlock(i)}.Reference(IR::Value{&phi});
             }
@@ -237,7 +243,7 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
                 }
             } else {
                 ctx.Add("MOV.S.CC RC,{};"
-                        "BRK (EQ.x);"
+                        "BRK(EQ.x);"
                         "ENDREP;",
                         eval(node.data.repeat.cond));
             }
-- 
cgit v1.2.3


From ca05a13c62ad7693f8be924c168e400e8139b0d2 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Tue, 25 May 2021 02:22:21 -0300
Subject: glasm: Catch more register leaks

Add support for null registers. These are used when an instruction has
no usages.

This comes handy when an instruction is only used for its CC value, with
the caveat of having to invalidate all pseudo-instructions before
defining the instruction itself in the register allocator. This commits
changes this.

Workaround a bug on Nvidia's condition codes conditional execution using
branches.
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 8 +++++++-
 1 file changed, 7 insertions(+), 1 deletion(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 2ce839059..4aa3682c2 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -203,7 +203,13 @@ void Precolor(EmitContext& ctx, const IR::Program& program) {
             for (size_t i = 0; i < num_args; ++i) {
                 IR::Block& phi_block{*phi.PhiBlock(i)};
                 auto it{std::find_if_not(phi_block.rbegin(), phi_block.rend(), IsReference).base()};
-                IR::IREmitter{phi_block, it}.PhiMove(phi, phi.Arg(i));
+                IR::IREmitter ir{phi_block, it};
+                const IR::Value arg{phi.Arg(i)};
+                if (arg.IsImmediate()) {
+                    ir.PhiMove(phi, arg);
+                } else {
+                    ir.PhiMove(phi, IR::Value{&RegAlloc::AliasInst(*arg.Inst())});
+                }
             }
             for (size_t i = 0; i < num_args; ++i) {
                 IR::IREmitter{*phi.PhiBlock(i)}.Reference(IR::Value{&phi});
-- 
cgit v1.2.3


From 379b305b4bc09799d53981fa6e5d9cbe6be99561 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Tue, 25 May 2021 02:34:08 -0300
Subject: glasm: Throw when there are register leaks

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 +++
 1 file changed, 3 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 4aa3682c2..0e9dc06a6 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -271,6 +271,9 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
             break;
         }
     }
+    if (!ctx.reg_alloc.IsEmpty()) {
+        throw LogicError("Register allocator is not empty");
+    }
 }
 
 void SetupOptions(const IR::Program& program, const Profile& profile,
-- 
cgit v1.2.3


From 75fd0079db9ac2f3bc6bcf182ed080a58538ed06 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Tue, 25 May 2021 02:46:51 -0300
Subject: glasm: Remove unnecessary value types

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 12 +-----------
 1 file changed, 1 insertion(+), 11 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 0e9dc06a6..5ffefaad2 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -49,8 +49,7 @@ public:
             inst = ir_value.InstRecursive();
             reg = Register{value};
         } else {
-            const bool is_long{value.type == Type::F64 || value.type == Type::U64};
-            reg = is_long ? reg_alloc.AllocLongReg() : reg_alloc.AllocReg();
+            reg = value.type == Type::U64 ? reg_alloc.AllocLongReg() : reg_alloc.AllocReg();
         }
         switch (value.type) {
         case Type::Register:
@@ -59,18 +58,9 @@ public:
         case Type::U32:
             ctx.Add("MOV.U {}.x,{};", reg, value.imm_u32);
             break;
-        case Type::S32:
-            ctx.Add("MOV.S {}.x,{};", reg, value.imm_s32);
-            break;
-        case Type::F32:
-            ctx.Add("MOV.F {}.x,{};", reg, value.imm_f32);
-            break;
         case Type::U64:
             ctx.Add("MOV.U64 {}.x,{};", reg, value.imm_u64);
             break;
-        case Type::F64:
-            ctx.Add("MOV.F64 {}.x,{};", reg, value.imm_f64);
-            break;
         }
     }
 
-- 
cgit v1.2.3


From 48aafe0961a2ddfb52b627c6ba6bce8276330550 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Tue, 25 May 2021 17:37:35 -0300
Subject: glasm: Release phi node registers after they are no longer needed

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 53 +++++++++-------------
 1 file changed, 22 insertions(+), 31 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 5ffefaad2..4f838b699 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -172,38 +172,29 @@ bool IsReference(IR::Inst& inst) {
     return inst.GetOpcode() == IR::Opcode::Reference;
 }
 
-void Precolor(EmitContext& ctx, const IR::Program& program) {
+void PrecolorInst(IR::Inst& phi) {
+    // Insert phi moves before references to avoid overwritting other phis
+    const size_t num_args{phi.NumArgs()};
+    for (size_t i = 0; i < num_args; ++i) {
+        IR::Block& phi_block{*phi.PhiBlock(i)};
+        auto it{std::find_if_not(phi_block.rbegin(), phi_block.rend(), IsReference).base()};
+        IR::IREmitter ir{phi_block, it};
+        const IR::Value arg{phi.Arg(i)};
+        if (arg.IsImmediate()) {
+            ir.PhiMove(phi, arg);
+        } else {
+            ir.PhiMove(phi, IR::Value{&RegAlloc::AliasInst(*arg.Inst())});
+        }
+    }
+    for (size_t i = 0; i < num_args; ++i) {
+        IR::IREmitter{*phi.PhiBlock(i)}.Reference(IR::Value{&phi});
+    }
+}
+
+void Precolor(const IR::Program& program) {
     for (IR::Block* const block : program.blocks) {
         for (IR::Inst& phi : block->Instructions() | std::views::take_while(IR::IsPhi)) {
-            switch (phi.Arg(0).Type()) {
-            case IR::Type::U1:
-            case IR::Type::U32:
-            case IR::Type::F32:
-                ctx.reg_alloc.Define(phi);
-                break;
-            case IR::Type::U64:
-            case IR::Type::F64:
-                ctx.reg_alloc.LongDefine(phi);
-                break;
-            default:
-                throw NotImplementedException("Phi node type {}", phi.Type());
-            }
-            // Insert phi moves before references to avoid overwritting them
-            const size_t num_args{phi.NumArgs()};
-            for (size_t i = 0; i < num_args; ++i) {
-                IR::Block& phi_block{*phi.PhiBlock(i)};
-                auto it{std::find_if_not(phi_block.rbegin(), phi_block.rend(), IsReference).base()};
-                IR::IREmitter ir{phi_block, it};
-                const IR::Value arg{phi.Arg(i)};
-                if (arg.IsImmediate()) {
-                    ir.PhiMove(phi, arg);
-                } else {
-                    ir.PhiMove(phi, IR::Value{&RegAlloc::AliasInst(*arg.Inst())});
-                }
-            }
-            for (size_t i = 0; i < num_args; ++i) {
-                IR::IREmitter{*phi.PhiBlock(i)}.Reference(IR::Value{&phi});
-            }
+            PrecolorInst(phi);
         }
     }
 }
@@ -388,7 +379,7 @@ std::string_view GetTessSpacing(TessSpacing spacing) {
 std::string EmitGLASM(const Profile& profile, const RuntimeInfo& runtime_info, IR::Program& program,
                       Bindings& bindings) {
     EmitContext ctx{program, bindings, profile, runtime_info};
-    Precolor(ctx, program);
+    Precolor(program);
     EmitCode(ctx, program);
     std::string header{StageHeader(program.stage)};
     SetupOptions(program, profile, runtime_info, header);
-- 
cgit v1.2.3


From f58f79c85dad7ad018a015cc6913f2789540ec22 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Wed, 26 May 2021 16:00:36 -0300
Subject: glasm: Implement Y direction

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 +++
 1 file changed, 3 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 4f838b699..2a0524609 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -439,6 +439,9 @@ std::string EmitGLASM(const Profile& profile, const RuntimeInfo& runtime_info, I
                   "MOV.F FSWZB[2],1;"
                   "MOV.F FSWZB[3],-1;";
     }
+    if (ctx.uses_y_direction) {
+        header += "PARAM y_direction[1]={state.material.front.ambient};";
+    }
     ctx.code.insert(0, header);
     ctx.code += "END";
     return ctx.code;
-- 
cgit v1.2.3


From b6c087496b14f1f5b253c3ecb82c00ded743418a Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Thu, 27 May 2021 17:54:08 -0300
Subject: glasm: Reduce reg allocation leaks from an exception to a log

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 2a0524609..e23208d2c 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -253,7 +253,7 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
         }
     }
     if (!ctx.reg_alloc.IsEmpty()) {
-        throw LogicError("Register allocator is not empty");
+        // LOG_WARNING ...;
     }
 }
 
-- 
cgit v1.2.3


From 916ca7432474e891864524dcbc6c879d5cdbfb72 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Sun, 30 May 2021 03:40:19 -0300
Subject: opengl: Declare fragment outputs even if they are not used

Fixes Ori and the Blind Forest's menu on GLASM. For some reason
(probably high level optimizations) it is not sanitized on SPIR-V for
OpenGL. Vulkan is unaffected by this change.
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index e23208d2c..70ca6f621 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -298,8 +298,7 @@ void SetupOptions(const IR::Program& program, const Profile& profile,
     if (stage == Stage::Fragment && runtime_info.force_early_z != 0) {
         header += "OPTION NV_early_fragment_tests;";
     }
-    const auto non_zero_frag_colors{info.stores_frag_color | std::views::drop(1)};
-    if (std::ranges::find(non_zero_frag_colors, true) != non_zero_frag_colors.end()) {
+    if (stage == Stage::Fragment) {
         header += "OPTION ARB_draw_buffers;";
     }
 }
-- 
cgit v1.2.3


From 79f2fe1a39120f498e915fa0c740b15dc0f09793 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Wed, 2 Jun 2021 03:02:33 -0300
Subject: glasm: Use ARB_derivative_control conditionally

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 70ca6f621..fc01797b6 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -265,9 +265,7 @@ void SetupOptions(const IR::Program& program, const Profile& profile,
     // TODO: Track the shared atomic ops
     header += "OPTION NV_internal;"
               "OPTION NV_shader_storage_buffer;"
-              "OPTION NV_gpu_program_fp64;"
-              "OPTION NV_bindless_texture;"
-              "OPTION ARB_derivative_control;";
+              "OPTION NV_gpu_program_fp64;";
     if (info.uses_int64_bit_atomics) {
         header += "OPTION NV_shader_atomic_int64;";
     }
@@ -295,6 +293,9 @@ void SetupOptions(const IR::Program& program, const Profile& profile,
     if (info.uses_typeless_image_reads && profile.support_typeless_image_loads) {
         header += "OPTION EXT_shader_image_load_formatted;";
     }
+    if (profile.support_derivative_control) {
+        header += "OPTION ARB_derivative_control;";
+    }
     if (stage == Stage::Fragment && runtime_info.force_early_z != 0) {
         header += "OPTION NV_early_fragment_tests;";
     }
-- 
cgit v1.2.3


From 61cd7dd30128633b656ce3264da74bef1ba00bb5 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 14 Jun 2021 02:27:49 -0300
Subject: shader: Add logging

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index fc01797b6..832b4fd40 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -253,7 +253,7 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
         }
     }
     if (!ctx.reg_alloc.IsEmpty()) {
-        // LOG_WARNING ...;
+        LOG_WARNING(Shader_GLASM, "Register leak after generating code");
     }
 }
 
-- 
cgit v1.2.3


From 0ffea97e2ea2c8f58928e13dc2488d620ea98ea8 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Wed, 16 Jun 2021 03:22:56 -0300
Subject: shader: Split profile and runtime info headers

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 1 +
 1 file changed, 1 insertion(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 832b4fd40..66e4aea04 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -13,6 +13,7 @@
 #include "shader_recompiler/frontend/ir/ir_emitter.h"
 #include "shader_recompiler/frontend/ir/program.h"
 #include "shader_recompiler/profile.h"
+#include "shader_recompiler/runtime_info.h"
 
 namespace Shader::Backend::GLASM {
 namespace {
-- 
cgit v1.2.3


From 3a2dd1b48310e2912e7f7f90da15bff555ef7256 Mon Sep 17 00:00:00 2001
From: ameerj <52414509+ameerj@users.noreply.github.com>
Date: Sun, 20 Jun 2021 01:40:21 -0400
Subject: glasm: Implement SetAttribute ViewportMask

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 66e4aea04..f39b02f77 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -286,7 +286,8 @@ void SetupOptions(const IR::Program& program, const Profile& profile,
     if (info.uses_sparse_residency) {
         header += "OPTION EXT_sparse_texture2;";
     }
-    if ((info.stores_viewport_index || info.stores_layer) && stage != Stage::Geometry) {
+    if (((info.stores_viewport_index || info.stores_layer) && stage != Stage::Geometry) ||
+        info.stores_viewport_mask) {
         if (profile.support_viewport_index_layer_non_geometry) {
             header += "OPTION NV_viewport_array2;";
         }
-- 
cgit v1.2.3


From 808ef97a086e7cc58a3ceded1de516ad6a6be5d3 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 21 Jun 2021 01:07:10 -0300
Subject: shader: Move loop safety tests to code emission

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 17 +++++++++++++++++
 1 file changed, 17 insertions(+)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index f39b02f77..79314f130 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -6,6 +6,8 @@
 #include <string>
 #include <tuple>
 
+#include "common/div_ceil.h"
+#include "common/settings.h"
 #include "shader_recompiler/backend/bindings.h"
 #include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm.h"
@@ -222,6 +224,14 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
             ctx.Add("REP;");
             break;
         case IR::AbstractSyntaxNode::Type::Repeat:
+            if (!Settings::values.disable_shader_loop_safety_checks) {
+                const u32 loop_index{ctx.num_safety_loop_vars++};
+                const u32 vector_index{loop_index / 4};
+                const char component{"xyzw"[loop_index % 4]};
+                ctx.Add("SUB.S.CC loop{}.{},loop{}.{},1;"
+                        "BRK(LT.{});",
+                        vector_index, component, vector_index, component, component);
+            }
             if (node.data.repeat.cond.IsImmediate()) {
                 if (node.data.repeat.cond.U1()) {
                     ctx.Add("ENDREP;");
@@ -425,6 +435,10 @@ std::string EmitGLASM(const Profile& profile, const RuntimeInfo& runtime_info, I
     if (program.info.uses_fswzadd) {
         header += "FSWZA[4],FSWZB[4],";
     }
+    const u32 num_safety_loop_vectors{Common::DivCeil(ctx.num_safety_loop_vars, 4u)};
+    for (u32 index = 0; index < num_safety_loop_vectors; ++index) {
+        header += fmt::format("loop{},", index);
+    }
     header += "RC;"
               "LONG TEMP ";
     for (size_t index = 0; index < ctx.reg_alloc.NumUsedLongRegisters(); ++index) {
@@ -441,6 +455,9 @@ std::string EmitGLASM(const Profile& profile, const RuntimeInfo& runtime_info, I
                   "MOV.F FSWZB[2],1;"
                   "MOV.F FSWZB[3],-1;";
     }
+    for (u32 index = 0; index < num_safety_loop_vectors; ++index) {
+        header += fmt::format("MOV.S loop{},{{0x2000,0x2000,0x2000,0x2000}};", index);
+    }
     if (ctx.uses_y_direction) {
         header += "PARAM y_direction[1]={state.material.front.ambient};";
     }
-- 
cgit v1.2.3


From 7dafa96ab59892b7f1fbffdb61e4326e6443955f Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Thu, 24 Jun 2021 02:41:09 -0300
Subject: shader: Rework varyings and implement passthrough geometry shaders

Put all varyings into a single std::bitset with helpers to access it.

Implement passthrough geometry shaders using host's.
---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 79314f130..2b96977b3 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -296,8 +296,10 @@ void SetupOptions(const IR::Program& program, const Profile& profile,
     if (info.uses_sparse_residency) {
         header += "OPTION EXT_sparse_texture2;";
     }
-    if (((info.stores_viewport_index || info.stores_layer) && stage != Stage::Geometry) ||
-        info.stores_viewport_mask) {
+    const bool stores_viewport_layer{info.stores[IR::Attribute::ViewportIndex] ||
+                                     info.stores[IR::Attribute::Layer]};
+    if ((stage != Stage::Geometry && stores_viewport_layer) ||
+        info.stores[IR::Attribute::ViewportMask]) {
         if (profile.support_viewport_index_layer_non_geometry) {
             header += "OPTION NV_viewport_array2;";
         }
-- 
cgit v1.2.3


From 8a3427a4c857aa08e365d1776d1f0d9f32639c9c Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Thu, 24 Jun 2021 17:40:24 -0300
Subject: glasm: Add passthrough geometry shader support

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 28 ++++++++++++++++++----
 1 file changed, 23 insertions(+), 5 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 2b96977b3..64787b353 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -304,6 +304,9 @@ void SetupOptions(const IR::Program& program, const Profile& profile,
             header += "OPTION NV_viewport_array2;";
         }
     }
+    if (program.is_geometry_passthrough && profile.support_geometry_shader_passthrough) {
+        header += "OPTION NV_geometry_shader_passthrough;";
+    }
     if (info.uses_typeless_image_reads && profile.support_typeless_image_loads) {
         header += "OPTION EXT_shader_image_load_formatted;";
     }
@@ -410,11 +413,26 @@ std::string EmitGLASM(const Profile& profile, const RuntimeInfo& runtime_info, I
                               runtime_info.tess_clockwise ? "CW" : "CCW");
         break;
     case Stage::Geometry:
-        header += fmt::format("PRIMITIVE_IN {};"
-                              "PRIMITIVE_OUT {};"
-                              "VERTICES_OUT {};",
-                              InputPrimitive(runtime_info.input_topology),
-                              OutputPrimitive(program.output_topology), program.output_vertices);
+        header += fmt::format("PRIMITIVE_IN {};", InputPrimitive(runtime_info.input_topology));
+        if (program.is_geometry_passthrough) {
+            if (profile.support_geometry_shader_passthrough) {
+                for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
+                    if (program.info.passthrough.Generic(index)) {
+                        header += fmt::format("PASSTHROUGH result.attrib[{}];", index);
+                    }
+                }
+                if (program.info.passthrough.AnyComponent(IR::Attribute::PositionX)) {
+                    header += "PASSTHROUGH result.position;";
+                }
+            } else {
+                LOG_WARNING(Shader_GLASM, "Passthrough geometry program used but not supported");
+            }
+        } else {
+            header +=
+                fmt::format("VERTICES_OUT {};"
+                            "PRIMITIVE_OUT {};",
+                            program.output_vertices, OutputPrimitive(program.output_topology));
+        }
         break;
     case Stage::Compute:
         header += fmt::format("GROUP_SIZE {} {} {};", program.workgroup_size[0],
-- 
cgit v1.2.3


From bf2956d77ab0ad06c4b5505cc9906e51e5878274 Mon Sep 17 00:00:00 2001
From: ReinUsesLisp <reinuseslisp@airmail.cc>
Date: Mon, 12 Jul 2021 05:22:01 -0300
Subject: shader: Avoid usage of C++20 ranges to build in clang

---
 src/shader_recompiler/backend/glasm/emit_glasm.cpp | 7 +++++--
 1 file changed, 5 insertions(+), 2 deletions(-)

(limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')

diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index 64787b353..a5e8c9b6e 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -2,7 +2,7 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
-#include <ranges>
+#include <algorithm>
 #include <string>
 #include <tuple>
 
@@ -196,7 +196,10 @@ void PrecolorInst(IR::Inst& phi) {
 
 void Precolor(const IR::Program& program) {
     for (IR::Block* const block : program.blocks) {
-        for (IR::Inst& phi : block->Instructions() | std::views::take_while(IR::IsPhi)) {
+        for (IR::Inst& phi : block->Instructions()) {
+            if (!IR::IsPhi(phi)) {
+                break;
+            }
             PrecolorInst(phi);
         }
     }
-- 
cgit v1.2.3