early-access version 2281

2021-12-07 02:20:09 +01:00
parent c2ae6d480a
commit c4fa174d53
591 changed files with 36978 additions and 18653 deletions
--- a/src/shader_recompiler/CMakeLists.txt
+++ b/src/shader_recompiler/CMakeLists.txt
@@ -1,7 +1,5 @@
 add_library(shader_recompiler STATIC
    backend/bindings.h
-    backend/glasm/emit_context.cpp
-    backend/glasm/emit_context.h
    backend/glasm/emit_glasm.cpp
    backend/glasm/emit_glasm.h
    backend/glasm/emit_glasm_barriers.cpp
@@ -22,10 +20,10 @@ add_library(shader_recompiler STATIC
    backend/glasm/emit_glasm_special.cpp
    backend/glasm/emit_glasm_undefined.cpp
    backend/glasm/emit_glasm_warp.cpp
+    backend/glasm/glasm_emit_context.cpp
+    backend/glasm/glasm_emit_context.h
    backend/glasm/reg_alloc.cpp
    backend/glasm/reg_alloc.h
-    backend/glsl/emit_context.cpp
-    backend/glsl/emit_context.h
    backend/glsl/emit_glsl.cpp
    backend/glsl/emit_glsl.h
    backend/glsl/emit_glsl_atomic.cpp
@@ -47,10 +45,10 @@ add_library(shader_recompiler STATIC
    backend/glsl/emit_glsl_special.cpp
    backend/glsl/emit_glsl_undefined.cpp
    backend/glsl/emit_glsl_warp.cpp
+    backend/glsl/glsl_emit_context.cpp
+    backend/glsl/glsl_emit_context.h
    backend/glsl/var_alloc.cpp
    backend/glsl/var_alloc.h
-    backend/spirv/emit_context.cpp
-    backend/spirv/emit_context.h
    backend/spirv/emit_spirv.cpp
    backend/spirv/emit_spirv.h
    backend/spirv/emit_spirv_atomic.cpp
@@ -72,6 +70,8 @@ add_library(shader_recompiler STATIC
    backend/spirv/emit_spirv_special.cpp
    backend/spirv/emit_spirv_undefined.cpp
    backend/spirv/emit_spirv_warp.cpp
+    backend/spirv/spirv_emit_context.cpp
+    backend/spirv/spirv_emit_context.h
    environment.h
    exception.h
    frontend/ir/abstract_syntax_list.h
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -9,9 +9,9 @@
 #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"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/ir_emitter.h"
 #include "shader_recompiler/frontend/ir/program.h"
 #include "shader_recompiler/profile.h"
--- a/src/shader_recompiler/backend/glasm/emit_glasm_barriers.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_barriers.cpp
@@ -0,0 +1,22 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
+
+namespace Shader::Backend::GLASM {
+
+void EmitBarrier(EmitContext& ctx) {
+    ctx.Add("BAR;");
+}
+
+void EmitWorkgroupMemoryBarrier(EmitContext& ctx) {
+    ctx.Add("MEMBAR.CTA;");
+}
+
+void EmitDeviceMemoryBarrier(EmitContext& ctx) {
+    ctx.Add("MEMBAR;");
+}
+
+} // namespace Shader::Backend::GLASM
--- a/src/shader_recompiler/backend/glasm/emit_glasm_bitwise_conversion.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_bitwise_conversion.cpp
@@ -2,8 +2,8 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLASM {
--- a/src/shader_recompiler/backend/glasm/emit_glasm_composite.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_composite.cpp
@@ -2,8 +2,8 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLASM {
--- a/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"
 #include "shader_recompiler/profile.h"
 #include "shader_recompiler/shader_info.h"
@@ -335,6 +335,35 @@ void EmitSetFragDepth(EmitContext& ctx, ScalarF32 value) {
    ctx.Add("MOV.F result.depth.z,{};", value);
 }

+void EmitWorkgroupId(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {},invocation.groupid;", inst);
+}
+
+void EmitLocalInvocationId(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {},invocation.localid;", inst);
+}
+
+void EmitInvocationId(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {}.x,primitive_invocation.x;", inst);
+}
+
+void EmitSampleId(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {}.x,fragment.sampleid.x;", inst);
+}
+
+void EmitIsHelperInvocation(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {}.x,fragment.helperthread.x;", inst);
+}
+
+void EmitYDirection(EmitContext& ctx, IR::Inst& inst) {
+    ctx.uses_y_direction = true;
+    ctx.Add("MOV.F {}.x,y_direction[0].w;", inst);
+}
+
+void EmitResolutionDownFactor(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.F {}.x,scaling[0].z;", inst);
+}
+
 void EmitLoadLocal(EmitContext& ctx, IR::Inst& inst, ScalarU32 word_offset) {
    ctx.Add("MOV.U {},lmem[{}].x;", inst, word_offset);
 }
--- a/src/shader_recompiler/backend/glasm/emit_glasm_control_flow.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_control_flow.cpp
@@ -0,0 +1,18 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
+
+namespace Shader::Backend::GLASM {
+
+void EmitJoin(EmitContext&) {
+    throw NotImplementedException("Join shouldn't be emitted");
+}
+
+void EmitDemoteToHelperInvocation(EmitContext& ctx) {
+    ctx.Add("KIL TR.x;");
+}
+
+} // namespace Shader::Backend::GLASM
--- a/src/shader_recompiler/backend/glasm/emit_glasm_convert.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_convert.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"
 #include "shader_recompiler/frontend/ir/value.h"

--- a/src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"
 #include "shader_recompiler/frontend/ir/value.h"

--- a/src/shader_recompiler/backend/glasm/emit_glasm_image.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_image.cpp
@@ -4,8 +4,8 @@

 #include <utility>

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"
 #include "shader_recompiler/frontend/ir/value.h"

--- a/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp
@@ -2,8 +2,8 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLASM {
--- a/src/shader_recompiler/backend/glasm/emit_glasm_logical.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_logical.cpp
@@ -0,0 +1,26 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
+
+namespace Shader::Backend::GLASM {
+
+void EmitLogicalOr(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) {
+    ctx.Add("OR.S {},{},{};", inst, a, b);
+}
+
+void EmitLogicalAnd(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) {
+    ctx.Add("AND.S {},{},{};", inst, a, b);
+}
+
+void EmitLogicalXor(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) {
+    ctx.Add("XOR.S {},{},{};", inst, a, b);
+}
+
+void EmitLogicalNot(EmitContext& ctx, IR::Inst& inst, ScalarS32 value) {
+    ctx.Add("SEQ.S {},{},0;", inst, value);
+}
+
+} // namespace Shader::Backend::GLASM
--- a/src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/program.h"
 #include "shader_recompiler/frontend/ir/value.h"
 #include "shader_recompiler/runtime_info.h"
--- a/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/program.h"
 #include "shader_recompiler/frontend/ir/value.h"

@@ -17,110 +17,6 @@ namespace Shader::Backend::GLASM {

 #define NotImplemented() throw NotImplementedException("GLASM instruction {}", __LINE__)

-static void DefinePhi(EmitContext& ctx, IR::Inst& phi) {
-    switch (phi.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());
-    }
-}
-
-void EmitPhi(EmitContext& ctx, IR::Inst& phi) {
-    const size_t num_args{phi.NumArgs()};
-    for (size_t i = 0; i < num_args; ++i) {
-        ctx.reg_alloc.Consume(phi.Arg(i));
-    }
-    if (!phi.Definition<Id>().is_valid) {
-        // The phi node wasn't forward defined
-        DefinePhi(ctx, phi);
-    }
-}
-
-void EmitVoid(EmitContext&) {}
-
-void EmitReference(EmitContext& ctx, const IR::Value& value) {
-    ctx.reg_alloc.Consume(value);
-}
-
-void EmitPhiMove(EmitContext& ctx, const IR::Value& phi_value, const IR::Value& value) {
-    IR::Inst& phi{RegAlloc::AliasInst(*phi_value.Inst())};
-    if (!phi.Definition<Id>().is_valid) {
-        // The phi node wasn't forward defined
-        DefinePhi(ctx, phi);
-    }
-    const Register phi_reg{ctx.reg_alloc.Consume(IR::Value{&phi})};
-    const Value eval_value{ctx.reg_alloc.Consume(value)};
-
-    if (phi_reg == eval_value) {
-        return;
-    }
-    switch (phi.Flags<IR::Type>()) {
-    case IR::Type::U1:
-    case IR::Type::U32:
-    case IR::Type::F32:
-        ctx.Add("MOV.S {}.x,{};", phi_reg, ScalarS32{eval_value});
-        break;
-    case IR::Type::U64:
-    case IR::Type::F64:
-        ctx.Add("MOV.U64 {}.x,{};", phi_reg, ScalarRegister{eval_value});
-        break;
-    default:
-        throw NotImplementedException("Phi node type {}", phi.Type());
-    }
-}
-
-void EmitJoin(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitDemoteToHelperInvocation(EmitContext& ctx) {
-    ctx.Add("KIL TR.x;");
-}
-
-void EmitBarrier(EmitContext& ctx) {
-    ctx.Add("BAR;");
-}
-
-void EmitWorkgroupMemoryBarrier(EmitContext& ctx) {
-    ctx.Add("MEMBAR.CTA;");
-}
-
-void EmitDeviceMemoryBarrier(EmitContext& ctx) {
-    ctx.Add("MEMBAR;");
-}
-
-void EmitPrologue(EmitContext& ctx) {
-    // TODO
-}
-
-void EmitEpilogue(EmitContext& ctx) {
-    // TODO
-}
-
-void EmitEmitVertex(EmitContext& ctx, ScalarS32 stream) {
-    if (stream.type == Type::U32 && stream.imm_u32 == 0) {
-        ctx.Add("EMIT;");
-    } else {
-        ctx.Add("EMITS {};", stream);
-    }
-}
-
-void EmitEndPrimitive(EmitContext& ctx, const IR::Value& stream) {
-    if (!stream.IsImmediate()) {
-        LOG_WARNING(Shader_GLASM, "Stream is not immediate");
-    }
-    ctx.reg_alloc.Consume(stream);
-    ctx.Add("ENDPRIM;");
-}
-
 void EmitGetRegister(EmitContext& ctx) {
    NotImplemented();
 }
@@ -185,55 +81,6 @@ void EmitSetOFlag(EmitContext& ctx) {
    NotImplemented();
 }

-void EmitWorkgroupId(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {},invocation.groupid;", inst);
-}
-
-void EmitLocalInvocationId(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {},invocation.localid;", inst);
-}
-
-void EmitInvocationId(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {}.x,primitive_invocation.x;", inst);
-}
-
-void EmitSampleId(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {}.x,fragment.sampleid.x;", inst);
-}
-
-void EmitIsHelperInvocation(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {}.x,fragment.helperthread.x;", inst);
-}
-
-void EmitYDirection(EmitContext& ctx, IR::Inst& inst) {
-    ctx.uses_y_direction = true;
-    ctx.Add("MOV.F {}.x,y_direction[0].w;", inst);
-}
-
-void EmitResolutionDownFactor(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.F {}.x,scaling[0].z;", inst);
-}
-
-void EmitUndefU1(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {}.x,0;", inst);
-}
-
-void EmitUndefU8(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {}.x,0;", inst);
-}
-
-void EmitUndefU16(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {}.x,0;", inst);
-}
-
-void EmitUndefU32(EmitContext& ctx, IR::Inst& inst) {
-    ctx.Add("MOV.S {}.x,0;", inst);
-}
-
-void EmitUndefU64(EmitContext& ctx, IR::Inst& inst) {
-    ctx.LongAdd("MOV.S64 {}.x,0;", inst);
-}
-
 void EmitGetZeroFromOp(EmitContext& ctx) {
    NotImplemented();
 }
@@ -258,20 +105,4 @@ void EmitGetInBoundsFromOp(EmitContext& ctx) {
    NotImplemented();
 }

-void EmitLogicalOr(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) {
-    ctx.Add("OR.S {},{},{};", inst, a, b);
-}
-
-void EmitLogicalAnd(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) {
-    ctx.Add("AND.S {},{},{};", inst, a, b);
-}
-
-void EmitLogicalXor(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) {
-    ctx.Add("XOR.S {},{},{};", inst, a, b);
-}
-
-void EmitLogicalNot(EmitContext& ctx, IR::Inst& inst, ScalarS32 value) {
-    ctx.Add("SEQ.S {},{},0;", inst, value);
-}
-
 } // namespace Shader::Backend::GLASM
--- a/src/shader_recompiler/backend/glasm/emit_glasm_select.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_select.cpp
@@ -3,8 +3,8 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLASM {
--- a/src/shader_recompiler/backend/glasm/emit_glasm_shared_memory.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_shared_memory.cpp
@@ -3,8 +3,8 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLASM {
--- a/src/shader_recompiler/backend/glasm/emit_glasm_special.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_special.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 "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
+#include "shader_recompiler/frontend/ir/value.h"
+
+namespace Shader::Backend::GLASM {
+
+static void DefinePhi(EmitContext& ctx, IR::Inst& phi) {
+    switch (phi.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());
+    }
+}
+
+void EmitPhi(EmitContext& ctx, IR::Inst& phi) {
+    const size_t num_args{phi.NumArgs()};
+    for (size_t i = 0; i < num_args; ++i) {
+        ctx.reg_alloc.Consume(phi.Arg(i));
+    }
+    if (!phi.Definition<Id>().is_valid) {
+        // The phi node wasn't forward defined
+        DefinePhi(ctx, phi);
+    }
+}
+
+void EmitVoid(EmitContext&) {}
+
+void EmitReference(EmitContext& ctx, const IR::Value& value) {
+    ctx.reg_alloc.Consume(value);
+}
+
+void EmitPhiMove(EmitContext& ctx, const IR::Value& phi_value, const IR::Value& value) {
+    IR::Inst& phi{RegAlloc::AliasInst(*phi_value.Inst())};
+    if (!phi.Definition<Id>().is_valid) {
+        // The phi node wasn't forward defined
+        DefinePhi(ctx, phi);
+    }
+    const Register phi_reg{ctx.reg_alloc.Consume(IR::Value{&phi})};
+    const Value eval_value{ctx.reg_alloc.Consume(value)};
+
+    if (phi_reg == eval_value) {
+        return;
+    }
+    switch (phi.Flags<IR::Type>()) {
+    case IR::Type::U1:
+    case IR::Type::U32:
+    case IR::Type::F32:
+        ctx.Add("MOV.S {}.x,{};", phi_reg, ScalarS32{eval_value});
+        break;
+    case IR::Type::U64:
+    case IR::Type::F64:
+        ctx.Add("MOV.U64 {}.x,{};", phi_reg, ScalarRegister{eval_value});
+        break;
+    default:
+        throw NotImplementedException("Phi node type {}", phi.Type());
+    }
+}
+
+void EmitPrologue(EmitContext&) {
+    // TODO
+}
+
+void EmitEpilogue(EmitContext&) {
+    // TODO
+}
+
+void EmitEmitVertex(EmitContext& ctx, ScalarS32 stream) {
+    if (stream.type == Type::U32 && stream.imm_u32 == 0) {
+        ctx.Add("EMIT;");
+    } else {
+        ctx.Add("EMITS {};", stream);
+    }
+}
+
+void EmitEndPrimitive(EmitContext& ctx, const IR::Value& stream) {
+    if (!stream.IsImmediate()) {
+        LOG_WARNING(Shader_GLASM, "Stream is not immediate");
+    }
+    ctx.reg_alloc.Consume(stream);
+    ctx.Add("ENDPRIM;");
+}
+
+} // namespace Shader::Backend::GLASM
--- a/src/shader_recompiler/backend/glasm/emit_glasm_undefined.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_undefined.cpp
@@ -0,0 +1,30 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
+
+namespace Shader::Backend::GLASM {
+
+void EmitUndefU1(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {}.x,0;", inst);
+}
+
+void EmitUndefU8(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {}.x,0;", inst);
+}
+
+void EmitUndefU16(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {}.x,0;", inst);
+}
+
+void EmitUndefU32(EmitContext& ctx, IR::Inst& inst) {
+    ctx.Add("MOV.S {}.x,0;", inst);
+}
+
+void EmitUndefU64(EmitContext& ctx, IR::Inst& inst) {
+    ctx.LongAdd("MOV.S64 {}.x,0;", inst);
+}
+
+} // namespace Shader::Backend::GLASM
--- a/src/shader_recompiler/backend/glasm/emit_glasm_warp.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_warp.cpp
@@ -2,8 +2,8 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.

-#include "shader_recompiler/backend/glasm/emit_context.h"
 #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"
 #include "shader_recompiler/profile.h"

--- a/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp
+++ b/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp
@@ -0,0 +1,156 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include <string_view>
+
+#include "shader_recompiler/backend/bindings.h"
+#include "shader_recompiler/backend/glasm/emit_glasm.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
+#include "shader_recompiler/frontend/ir/program.h"
+#include "shader_recompiler/profile.h"
+#include "shader_recompiler/runtime_info.h"
+
+namespace Shader::Backend::GLASM {
+namespace {
+std::string_view InterpDecorator(Interpolation interp) {
+    switch (interp) {
+    case Interpolation::Smooth:
+        return "";
+    case Interpolation::Flat:
+        return "FLAT ";
+    case Interpolation::NoPerspective:
+        return "NOPERSPECTIVE ";
+    }
+    throw InvalidArgument("Invalid interpolation {}", interp);
+}
+
+bool IsInputArray(Stage stage) {
+    return stage == Stage::Geometry || stage == Stage::TessellationControl ||
+           stage == Stage::TessellationEval;
+}
+} // Anonymous namespace
+
+EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
+                         const RuntimeInfo& runtime_info_)
+    : info{program.info}, profile{profile_}, runtime_info{runtime_info_} {
+    // FIXME: Temporary partial implementation
+    u32 cbuf_index{};
+    for (const auto& desc : info.constant_buffer_descriptors) {
+        if (desc.count != 1) {
+            throw NotImplementedException("Constant buffer descriptor array");
+        }
+        Add("CBUFFER c{}[]={{program.buffer[{}]}};", desc.index, cbuf_index);
+        ++cbuf_index;
+    }
+    u32 ssbo_index{};
+    for (const auto& desc : info.storage_buffers_descriptors) {
+        if (desc.count != 1) {
+            throw NotImplementedException("Storage buffer descriptor array");
+        }
+        if (runtime_info.glasm_use_storage_buffers) {
+            Add("STORAGE ssbo{}[]={{program.storage[{}]}};", ssbo_index, bindings.storage_buffer);
+            ++bindings.storage_buffer;
+            ++ssbo_index;
+        }
+    }
+    if (!runtime_info.glasm_use_storage_buffers) {
+        if (const size_t num = info.storage_buffers_descriptors.size(); num > 0) {
+            const size_t index{num + PROGRAM_LOCAL_PARAMETER_STORAGE_BUFFER_BASE};
+            Add("PARAM c[{}]={{program.local[0..{}]}};", index, index - 1);
+        }
+    }
+    stage = program.stage;
+    switch (program.stage) {
+    case Stage::VertexA:
+    case Stage::VertexB:
+        stage_name = "vertex";
+        attrib_name = "vertex";
+        break;
+    case Stage::TessellationControl:
+    case Stage::TessellationEval:
+        stage_name = "primitive";
+        attrib_name = "primitive";
+        break;
+    case Stage::Geometry:
+        stage_name = "primitive";
+        attrib_name = "vertex";
+        break;
+    case Stage::Fragment:
+        stage_name = "fragment";
+        attrib_name = "fragment";
+        break;
+    case Stage::Compute:
+        stage_name = "invocation";
+        break;
+    }
+    const std::string_view attr_stage{stage == Stage::Fragment ? "fragment" : "vertex"};
+    const VaryingState loads{info.loads.mask | info.passthrough.mask};
+    for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
+        if (loads.Generic(index)) {
+            Add("{}ATTRIB in_attr{}[]={{{}.attrib[{}..{}]}};",
+                InterpDecorator(info.interpolation[index]), index, attr_stage, index, index);
+        }
+    }
+    if (IsInputArray(stage) && loads.AnyComponent(IR::Attribute::PositionX)) {
+        Add("ATTRIB vertex_position=vertex.position;");
+    }
+    if (info.uses_invocation_id) {
+        Add("ATTRIB primitive_invocation=primitive.invocation;");
+    }
+    if (info.stores_tess_level_outer) {
+        Add("OUTPUT result_patch_tessouter[]={{result.patch.tessouter[0..3]}};");
+    }
+    if (info.stores_tess_level_inner) {
+        Add("OUTPUT result_patch_tessinner[]={{result.patch.tessinner[0..1]}};");
+    }
+    if (info.stores.ClipDistances()) {
+        Add("OUTPUT result_clip[]={{result.clip[0..7]}};");
+    }
+    for (size_t index = 0; index < info.uses_patches.size(); ++index) {
+        if (!info.uses_patches[index]) {
+            continue;
+        }
+        if (stage == Stage::TessellationControl) {
+            Add("OUTPUT result_patch_attrib{}[]={{result.patch.attrib[{}..{}]}};"
+                "ATTRIB primitive_out_patch_attrib{}[]={{primitive.out.patch.attrib[{}..{}]}};",
+                index, index, index, index, index, index);
+        } else {
+            Add("ATTRIB primitive_patch_attrib{}[]={{primitive.patch.attrib[{}..{}]}};", index,
+                index, index);
+        }
+    }
+    if (stage == Stage::Fragment) {
+        Add("OUTPUT frag_color0=result.color;");
+        for (size_t index = 1; index < info.stores_frag_color.size(); ++index) {
+            Add("OUTPUT frag_color{}=result.color[{}];", index, index);
+        }
+    }
+    for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
+        if (info.stores.Generic(index)) {
+            Add("OUTPUT out_attr{}[]={{result.attrib[{}..{}]}};", index, index, index);
+        }
+    }
+    image_buffer_bindings.reserve(info.image_buffer_descriptors.size());
+    for (const auto& desc : info.image_buffer_descriptors) {
+        image_buffer_bindings.push_back(bindings.image);
+        bindings.image += desc.count;
+    }
+    image_bindings.reserve(info.image_descriptors.size());
+    for (const auto& desc : info.image_descriptors) {
+        image_bindings.push_back(bindings.image);
+        bindings.image += desc.count;
+    }
+    texture_buffer_bindings.reserve(info.texture_buffer_descriptors.size());
+    for (const auto& desc : info.texture_buffer_descriptors) {
+        texture_buffer_bindings.push_back(bindings.texture);
+        bindings.texture += desc.count;
+    }
+    texture_bindings.reserve(info.texture_descriptors.size());
+    for (const auto& desc : info.texture_descriptors) {
+        texture_bindings.push_back(bindings.texture);
+        bindings.texture += desc.count;
+    }
+}
+
+} // namespace Shader::Backend::GLASM
--- a/src/shader_recompiler/backend/glasm/glasm_emit_context.h
+++ b/src/shader_recompiler/backend/glasm/glasm_emit_context.h
@@ -0,0 +1,80 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#pragma once
+
+#include <string>
+#include <utility>
+#include <vector>
+
+#include <fmt/format.h>
+
+#include "shader_recompiler/backend/glasm/reg_alloc.h"
+#include "shader_recompiler/stage.h"
+
+namespace Shader {
+struct Info;
+struct Profile;
+struct RuntimeInfo;
+} // namespace Shader
+
+namespace Shader::Backend {
+struct Bindings;
+}
+
+namespace Shader::IR {
+class Inst;
+struct Program;
+} // namespace Shader::IR
+
+namespace Shader::Backend::GLASM {
+
+class EmitContext {
+public:
+    explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
+                         const RuntimeInfo& runtime_info_);
+
+    template <typename... Args>
+    void Add(const char* format_str, IR::Inst& inst, Args&&... args) {
+        code += fmt::format(fmt::runtime(format_str), reg_alloc.Define(inst),
+                            std::forward<Args>(args)...);
+        // TODO: Remove this
+        code += '\n';
+    }
+
+    template <typename... Args>
+    void LongAdd(const char* format_str, IR::Inst& inst, Args&&... args) {
+        code += fmt::format(fmt::runtime(format_str), reg_alloc.LongDefine(inst),
+                            std::forward<Args>(args)...);
+        // TODO: Remove this
+        code += '\n';
+    }
+
+    template <typename... Args>
+    void Add(const char* format_str, Args&&... args) {
+        code += fmt::format(fmt::runtime(format_str), std::forward<Args>(args)...);
+        // TODO: Remove this
+        code += '\n';
+    }
+
+    std::string code;
+    RegAlloc reg_alloc{};
+    const Info& info;
+    const Profile& profile;
+    const RuntimeInfo& runtime_info;
+
+    std::vector<u32> texture_buffer_bindings;
+    std::vector<u32> image_buffer_bindings;
+    std::vector<u32> texture_bindings;
+    std::vector<u32> image_bindings;
+
+    Stage stage{};
+    std::string_view stage_name = "invalid";
+    std::string_view attrib_name = "invalid";
+
+    u32 num_safety_loop_vars{};
+    bool uses_y_direction{};
+};
+
+} // namespace Shader::Backend::GLASM
--- a/src/shader_recompiler/backend/glasm/reg_alloc.cpp
+++ b/src/shader_recompiler/backend/glasm/reg_alloc.cpp
@@ -6,7 +6,7 @@

 #include <fmt/format.h>

-#include "shader_recompiler/backend/glasm/emit_context.h"
+#include "shader_recompiler/backend/glasm/glasm_emit_context.h"
 #include "shader_recompiler/backend/glasm/reg_alloc.h"
 #include "shader_recompiler/exception.h"
 #include "shader_recompiler/frontend/ir/value.h"
--- a/src/shader_recompiler/backend/glsl/emit_glsl.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl.cpp
@@ -9,9 +9,9 @@

 #include "common/div_ceil.h"
 #include "common/settings.h"
-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/ir_emitter.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_barriers.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_barriers.cpp
@@ -2,8 +2,8 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_bitwise_conversion.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_bitwise_conversion.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_composite.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_composite.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"
 #include "shader_recompiler/profile.h"
 #include "shader_recompiler/runtime_info.h"
--- a/src/shader_recompiler/backend/glsl/emit_glsl_control_flow.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_control_flow.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/exception.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_convert.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_convert.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_floating_point.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_floating_point.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"
 #include "shader_recompiler/frontend/ir/value.h"

--- a/src/shader_recompiler/backend/glsl/emit_glsl_image.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_image.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"
 #include "shader_recompiler/frontend/ir/value.h"
 #include "shader_recompiler/profile.h"
--- a/src/shader_recompiler/backend/glsl/emit_glsl_integer.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_integer.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_logical.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_logical.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"
 #include "shader_recompiler/profile.h"

--- a/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 #ifdef _MSC_VER
--- a/src/shader_recompiler/backend/glsl/emit_glsl_select.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_select.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_shared_memory.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_shared_memory.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"

 namespace Shader::Backend::GLSL {
--- a/src/shader_recompiler/backend/glsl/emit_glsl_special.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_special.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/program.h"
 #include "shader_recompiler/frontend/ir/value.h"
 #include "shader_recompiler/profile.h"
--- a/src/shader_recompiler/backend/glsl/emit_glsl_undefined.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_undefined.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"

 namespace Shader::Backend::GLSL {

--- a/src/shader_recompiler/backend/glsl/emit_glsl_warp.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_warp.cpp
@@ -4,8 +4,8 @@

 #include <string_view>

-#include "shader_recompiler/backend/glsl/emit_context.h"
 #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
 #include "shader_recompiler/frontend/ir/value.h"
 #include "shader_recompiler/profile.h"

--- a/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp
+++ b/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp
@@ -0,0 +1,681 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include "shader_recompiler/backend/bindings.h"
+#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
+#include "shader_recompiler/frontend/ir/program.h"
+#include "shader_recompiler/profile.h"
+#include "shader_recompiler/runtime_info.h"
+
+namespace Shader::Backend::GLSL {
+namespace {
+u32 CbufIndex(size_t offset) {
+    return (offset / 4) % 4;
+}
+
+char Swizzle(size_t offset) {
+    return "xyzw"[CbufIndex(offset)];
+}
+
+std::string_view InterpDecorator(Interpolation interp) {
+    switch (interp) {
+    case Interpolation::Smooth:
+        return "";
+    case Interpolation::Flat:
+        return "flat ";
+    case Interpolation::NoPerspective:
+        return "noperspective ";
+    }
+    throw InvalidArgument("Invalid interpolation {}", interp);
+}
+
+std::string_view InputArrayDecorator(Stage stage) {
+    switch (stage) {
+    case Stage::Geometry:
+    case Stage::TessellationControl:
+    case Stage::TessellationEval:
+        return "[]";
+    default:
+        return "";
+    }
+}
+
+bool StoresPerVertexAttributes(Stage stage) {
+    switch (stage) {
+    case Stage::VertexA:
+    case Stage::VertexB:
+    case Stage::Geometry:
+    case Stage::TessellationEval:
+        return true;
+    default:
+        return false;
+    }
+}
+
+std::string OutputDecorator(Stage stage, u32 size) {
+    switch (stage) {
+    case Stage::TessellationControl:
+        return fmt::format("[{}]", size);
+    default:
+        return "";
+    }
+}
+
+std::string_view SamplerType(TextureType type, bool is_depth) {
+    if (is_depth) {
+        switch (type) {
+        case TextureType::Color1D:
+            return "sampler1DShadow";
+        case TextureType::ColorArray1D:
+            return "sampler1DArrayShadow";
+        case TextureType::Color2D:
+            return "sampler2DShadow";
+        case TextureType::ColorArray2D:
+            return "sampler2DArrayShadow";
+        case TextureType::ColorCube:
+            return "samplerCubeShadow";
+        case TextureType::ColorArrayCube:
+            return "samplerCubeArrayShadow";
+        default:
+            throw NotImplementedException("Texture type: {}", type);
+        }
+    }
+    switch (type) {
+    case TextureType::Color1D:
+        return "sampler1D";
+    case TextureType::ColorArray1D:
+        return "sampler1DArray";
+    case TextureType::Color2D:
+        return "sampler2D";
+    case TextureType::ColorArray2D:
+        return "sampler2DArray";
+    case TextureType::Color3D:
+        return "sampler3D";
+    case TextureType::ColorCube:
+        return "samplerCube";
+    case TextureType::ColorArrayCube:
+        return "samplerCubeArray";
+    case TextureType::Buffer:
+        return "samplerBuffer";
+    default:
+        throw NotImplementedException("Texture type: {}", type);
+    }
+}
+
+std::string_view ImageType(TextureType type) {
+    switch (type) {
+    case TextureType::Color1D:
+        return "uimage1D";
+    case TextureType::ColorArray1D:
+        return "uimage1DArray";
+    case TextureType::Color2D:
+        return "uimage2D";
+    case TextureType::ColorArray2D:
+        return "uimage2DArray";
+    case TextureType::Color3D:
+        return "uimage3D";
+    case TextureType::ColorCube:
+        return "uimageCube";
+    case TextureType::ColorArrayCube:
+        return "uimageCubeArray";
+    case TextureType::Buffer:
+        return "uimageBuffer";
+    default:
+        throw NotImplementedException("Image type: {}", type);
+    }
+}
+
+std::string_view ImageFormatString(ImageFormat format) {
+    switch (format) {
+    case ImageFormat::Typeless:
+        return "";
+    case ImageFormat::R8_UINT:
+        return ",r8ui";
+    case ImageFormat::R8_SINT:
+        return ",r8i";
+    case ImageFormat::R16_UINT:
+        return ",r16ui";
+    case ImageFormat::R16_SINT:
+        return ",r16i";
+    case ImageFormat::R32_UINT:
+        return ",r32ui";
+    case ImageFormat::R32G32_UINT:
+        return ",rg32ui";
+    case ImageFormat::R32G32B32A32_UINT:
+        return ",rgba32ui";
+    default:
+        throw NotImplementedException("Image format: {}", format);
+    }
+}
+
+std::string_view ImageAccessQualifier(bool is_written, bool is_read) {
+    if (is_written && !is_read) {
+        return "writeonly ";
+    }
+    if (is_read && !is_written) {
+        return "readonly ";
+    }
+    return "";
+}
+
+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_spacing";
+    case TessSpacing::FractionalOdd:
+        return "fractional_odd_spacing";
+    case TessSpacing::FractionalEven:
+        return "fractional_even_spacing";
+    }
+    throw InvalidArgument("Invalid tessellation spacing {}", spacing);
+}
+
+std::string_view InputPrimitive(InputTopology topology) {
+    switch (topology) {
+    case InputTopology::Points:
+        return "points";
+    case InputTopology::Lines:
+        return "lines";
+    case InputTopology::LinesAdjacency:
+        return "lines_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);
+}
+
+void SetupOutPerVertex(EmitContext& ctx, std::string& header) {
+    if (!StoresPerVertexAttributes(ctx.stage)) {
+        return;
+    }
+    if (ctx.uses_geometry_passthrough) {
+        return;
+    }
+    header += "out gl_PerVertex{vec4 gl_Position;";
+    if (ctx.info.stores[IR::Attribute::PointSize]) {
+        header += "float gl_PointSize;";
+    }
+    if (ctx.info.stores.ClipDistances()) {
+        header += "float gl_ClipDistance[];";
+    }
+    if (ctx.info.stores[IR::Attribute::ViewportIndex] &&
+        ctx.profile.support_viewport_index_layer_non_geometry && ctx.stage != Stage::Geometry) {
+        header += "int gl_ViewportIndex;";
+    }
+    header += "};";
+    if (ctx.info.stores[IR::Attribute::ViewportIndex] && ctx.stage == Stage::Geometry) {
+        header += "out int gl_ViewportIndex;";
+    }
+}
+
+void SetupInPerVertex(EmitContext& ctx, std::string& header) {
+    // Currently only required for TessellationControl to adhere to
+    // ARB_separate_shader_objects requirements
+    if (ctx.stage != Stage::TessellationControl) {
+        return;
+    }
+    const bool loads_position{ctx.info.loads.AnyComponent(IR::Attribute::PositionX)};
+    const bool loads_point_size{ctx.info.loads[IR::Attribute::PointSize]};
+    const bool loads_clip_distance{ctx.info.loads.ClipDistances()};
+    const bool loads_per_vertex{loads_position || loads_point_size || loads_clip_distance};
+    if (!loads_per_vertex) {
+        return;
+    }
+    header += "in gl_PerVertex{";
+    if (loads_position) {
+        header += "vec4 gl_Position;";
+    }
+    if (loads_point_size) {
+        header += "float gl_PointSize;";
+    }
+    if (loads_clip_distance) {
+        header += "float gl_ClipDistance[];";
+    }
+    header += "}gl_in[gl_MaxPatchVertices];";
+}
+
+} // Anonymous namespace
+
+EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
+                         const RuntimeInfo& runtime_info_)
+    : info{program.info}, profile{profile_}, runtime_info{runtime_info_}, stage{program.stage},
+      uses_geometry_passthrough{program.is_geometry_passthrough &&
+                                profile.support_geometry_shader_passthrough} {
+    if (profile.need_fastmath_off) {
+        header += "#pragma optionNV(fastmath off)\n";
+    }
+    SetupExtensions();
+    switch (program.stage) {
+    case Stage::VertexA:
+    case Stage::VertexB:
+        stage_name = "vs";
+        break;
+    case Stage::TessellationControl:
+        stage_name = "tcs";
+        header += fmt::format("layout(vertices={})out;", program.invocations);
+        break;
+    case Stage::TessellationEval:
+        stage_name = "tes";
+        header += fmt::format("layout({},{},{})in;", GetTessMode(runtime_info.tess_primitive),
+                              GetTessSpacing(runtime_info.tess_spacing),
+                              runtime_info.tess_clockwise ? "cw" : "ccw");
+        break;
+    case Stage::Geometry:
+        stage_name = "gs";
+        header += fmt::format("layout({})in;", InputPrimitive(runtime_info.input_topology));
+        if (uses_geometry_passthrough) {
+            header += "layout(passthrough)in gl_PerVertex{vec4 gl_Position;};";
+            break;
+        } else if (program.is_geometry_passthrough &&
+                   !profile.support_geometry_shader_passthrough) {
+            LOG_WARNING(Shader_GLSL, "Passthrough geometry program used but not supported");
+        }
+        header += fmt::format(
+            "layout({},max_vertices={})out;in gl_PerVertex{{vec4 gl_Position;}}gl_in[];",
+            OutputPrimitive(program.output_topology), program.output_vertices);
+        break;
+    case Stage::Fragment:
+        stage_name = "fs";
+        position_name = "gl_FragCoord";
+        if (runtime_info.force_early_z) {
+            header += "layout(early_fragment_tests)in;";
+        }
+        if (info.uses_sample_id) {
+            header += "in int gl_SampleID;";
+        }
+        if (info.stores_sample_mask) {
+            header += "out int gl_SampleMask[];";
+        }
+        break;
+    case Stage::Compute:
+        stage_name = "cs";
+        const u32 local_x{std::max(program.workgroup_size[0], 1u)};
+        const u32 local_y{std::max(program.workgroup_size[1], 1u)};
+        const u32 local_z{std::max(program.workgroup_size[2], 1u)};
+        header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;",
+                              local_x, local_y, local_z);
+        break;
+    }
+    SetupOutPerVertex(*this, header);
+    SetupInPerVertex(*this, header);
+
+    for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
+        if (!info.loads.Generic(index) || !runtime_info.previous_stage_stores.Generic(index)) {
+            continue;
+        }
+        const auto qualifier{uses_geometry_passthrough ? "passthrough"
+                                                       : fmt::format("location={}", index)};
+        header += fmt::format("layout({}){}in vec4 in_attr{}{};", qualifier,
+                              InterpDecorator(info.interpolation[index]), index,
+                              InputArrayDecorator(stage));
+    }
+    for (size_t index = 0; index < info.uses_patches.size(); ++index) {
+        if (!info.uses_patches[index]) {
+            continue;
+        }
+        const auto qualifier{stage == Stage::TessellationControl ? "out" : "in"};
+        header += fmt::format("layout(location={})patch {} vec4 patch{};", index, qualifier, index);
+    }
+    if (stage == Stage::Fragment) {
+        for (size_t index = 0; index < info.stores_frag_color.size(); ++index) {
+            if (!info.stores_frag_color[index] && !profile.need_declared_frag_colors) {
+                continue;
+            }
+            header += fmt::format("layout(location={})out vec4 frag_color{};", index, index);
+        }
+    }
+    for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
+        if (info.stores.Generic(index)) {
+            DefineGenericOutput(index, program.invocations);
+        }
+    }
+    if (info.uses_rescaling_uniform) {
+        header += "layout(location=0) uniform vec4 scaling;";
+    }
+    DefineConstantBuffers(bindings);
+    DefineStorageBuffers(bindings);
+    SetupImages(bindings);
+    SetupTextures(bindings);
+    DefineHelperFunctions();
+    DefineConstants();
+}
+
+void EmitContext::SetupExtensions() {
+    header += "#extension GL_ARB_separate_shader_objects : enable\n";
+    if (info.uses_shadow_lod && profile.support_gl_texture_shadow_lod) {
+        header += "#extension GL_EXT_texture_shadow_lod : enable\n";
+    }
+    if (info.uses_int64 && profile.support_int64) {
+        header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
+    }
+    if (info.uses_int64_bit_atomics) {
+        header += "#extension GL_NV_shader_atomic_int64 : enable\n";
+    }
+    if (info.uses_atomic_f32_add) {
+        header += "#extension GL_NV_shader_atomic_float : enable\n";
+    }
+    if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) {
+        header += "#extension GL_NV_shader_atomic_fp16_vector : enable\n";
+    }
+    if (info.uses_fp16) {
+        if (profile.support_gl_nv_gpu_shader_5) {
+            header += "#extension GL_NV_gpu_shader5 : enable\n";
+        }
+        if (profile.support_gl_amd_gpu_shader_half_float) {
+            header += "#extension GL_AMD_gpu_shader_half_float : enable\n";
+        }
+    }
+    if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote ||
+        info.uses_subgroup_shuffles || info.uses_fswzadd) {
+        header += "#extension GL_ARB_shader_ballot : enable\n"
+                  "#extension GL_ARB_shader_group_vote : enable\n";
+        if (!info.uses_int64 && profile.support_int64) {
+            header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
+        }
+        if (profile.support_gl_warp_intrinsics) {
+            header += "#extension GL_NV_shader_thread_shuffle : enable\n";
+        }
+    }
+    if ((info.stores[IR::Attribute::ViewportIndex] || info.stores[IR::Attribute::Layer]) &&
+        profile.support_viewport_index_layer_non_geometry && stage != Stage::Geometry) {
+        header += "#extension GL_ARB_shader_viewport_layer_array : enable\n";
+    }
+    if (info.uses_sparse_residency && profile.support_gl_sparse_textures) {
+        header += "#extension GL_ARB_sparse_texture2 : enable\n";
+    }
+    if (info.stores[IR::Attribute::ViewportMask] && profile.support_viewport_mask) {
+        header += "#extension GL_NV_viewport_array2 : enable\n";
+    }
+    if (info.uses_typeless_image_reads) {
+        header += "#extension GL_EXT_shader_image_load_formatted : enable\n";
+    }
+    if (info.uses_derivatives && profile.support_gl_derivative_control) {
+        header += "#extension GL_ARB_derivative_control : enable\n";
+    }
+    if (uses_geometry_passthrough) {
+        header += "#extension GL_NV_geometry_shader_passthrough : enable\n";
+    }
+}
+
+void EmitContext::DefineConstantBuffers(Bindings& bindings) {
+    if (info.constant_buffer_descriptors.empty()) {
+        return;
+    }
+    for (const auto& desc : info.constant_buffer_descriptors) {
+        header += fmt::format(
+            "layout(std140,binding={}) uniform {}_cbuf_{}{{vec4 {}_cbuf{}[{}];}};",
+            bindings.uniform_buffer, stage_name, desc.index, stage_name, desc.index, 4 * 1024);
+        bindings.uniform_buffer += desc.count;
+    }
+}
+
+void EmitContext::DefineStorageBuffers(Bindings& bindings) {
+    if (info.storage_buffers_descriptors.empty()) {
+        return;
+    }
+    u32 index{};
+    for (const auto& desc : info.storage_buffers_descriptors) {
+        header += fmt::format("layout(std430,binding={}) buffer {}_ssbo_{}{{uint {}_ssbo{}[];}};",
+                              bindings.storage_buffer, stage_name, bindings.storage_buffer,
+                              stage_name, index);
+        bindings.storage_buffer += desc.count;
+        index += desc.count;
+    }
+}
+
+void EmitContext::DefineGenericOutput(size_t index, u32 invocations) {
+    static constexpr std::string_view swizzle{"xyzw"};
+    const size_t base_index{static_cast<size_t>(IR::Attribute::Generic0X) + index * 4};
+    u32 element{0};
+    while (element < 4) {
+        std::string definition{fmt::format("layout(location={}", index)};
+        const u32 remainder{4 - element};
+        const TransformFeedbackVarying* xfb_varying{};
+        if (!runtime_info.xfb_varyings.empty()) {
+            xfb_varying = &runtime_info.xfb_varyings[base_index + element];
+            xfb_varying = xfb_varying && xfb_varying->components > 0 ? xfb_varying : nullptr;
+        }
+        const u32 num_components{xfb_varying ? xfb_varying->components : remainder};
+        if (element > 0) {
+            definition += fmt::format(",component={}", element);
+        }
+        if (xfb_varying) {
+            definition +=
+                fmt::format(",xfb_buffer={},xfb_stride={},xfb_offset={}", xfb_varying->buffer,
+                            xfb_varying->stride, xfb_varying->offset);
+        }
+        std::string name{fmt::format("out_attr{}", index)};
+        if (num_components < 4 || element > 0) {
+            name += fmt::format("_{}", swizzle.substr(element, num_components));
+        }
+        const auto type{num_components == 1 ? "float" : fmt::format("vec{}", num_components)};
+        definition += fmt::format(")out {} {}{};", type, name, OutputDecorator(stage, invocations));
+        header += definition;
+
+        const GenericElementInfo element_info{
+            .name = name,
+            .first_element = element,
+            .num_components = num_components,
+        };
+        std::fill_n(output_generics[index].begin() + element, num_components, element_info);
+        element += num_components;
+    }
+}
+
+void EmitContext::DefineHelperFunctions() {
+    header += "\n#define ftoi floatBitsToInt\n#define ftou floatBitsToUint\n"
+              "#define itof intBitsToFloat\n#define utof uintBitsToFloat\n";
+    if (info.uses_global_increment || info.uses_shared_increment) {
+        header += "uint CasIncrement(uint op_a,uint op_b){return op_a>=op_b?0u:(op_a+1u);}";
+    }
+    if (info.uses_global_decrement || info.uses_shared_decrement) {
+        header += "uint CasDecrement(uint op_a,uint op_b){"
+                  "return op_a==0||op_a>op_b?op_b:(op_a-1u);}";
+    }
+    if (info.uses_atomic_f32_add) {
+        header += "uint CasFloatAdd(uint op_a,float op_b){"
+                  "return ftou(utof(op_a)+op_b);}";
+    }
+    if (info.uses_atomic_f32x2_add) {
+        header += "uint CasFloatAdd32x2(uint op_a,vec2 op_b){"
+                  "return packHalf2x16(unpackHalf2x16(op_a)+op_b);}";
+    }
+    if (info.uses_atomic_f32x2_min) {
+        header += "uint CasFloatMin32x2(uint op_a,vec2 op_b){return "
+                  "packHalf2x16(min(unpackHalf2x16(op_a),op_b));}";
+    }
+    if (info.uses_atomic_f32x2_max) {
+        header += "uint CasFloatMax32x2(uint op_a,vec2 op_b){return "
+                  "packHalf2x16(max(unpackHalf2x16(op_a),op_b));}";
+    }
+    if (info.uses_atomic_f16x2_add) {
+        header += "uint CasFloatAdd16x2(uint op_a,f16vec2 op_b){return "
+                  "packFloat2x16(unpackFloat2x16(op_a)+op_b);}";
+    }
+    if (info.uses_atomic_f16x2_min) {
+        header += "uint CasFloatMin16x2(uint op_a,f16vec2 op_b){return "
+                  "packFloat2x16(min(unpackFloat2x16(op_a),op_b));}";
+    }
+    if (info.uses_atomic_f16x2_max) {
+        header += "uint CasFloatMax16x2(uint op_a,f16vec2 op_b){return "
+                  "packFloat2x16(max(unpackFloat2x16(op_a),op_b));}";
+    }
+    if (info.uses_atomic_s32_min) {
+        header += "uint CasMinS32(uint op_a,uint op_b){return uint(min(int(op_a),int(op_b)));}";
+    }
+    if (info.uses_atomic_s32_max) {
+        header += "uint CasMaxS32(uint op_a,uint op_b){return uint(max(int(op_a),int(op_b)));}";
+    }
+    if (info.uses_global_memory && profile.support_int64) {
+        header += DefineGlobalMemoryFunctions();
+    }
+    if (info.loads_indexed_attributes) {
+        const bool is_array{stage == Stage::Geometry};
+        const auto vertex_arg{is_array ? ",uint vertex" : ""};
+        std::string func{
+            fmt::format("float IndexedAttrLoad(int offset{}){{int base_index=offset>>2;uint "
+                        "masked_index=uint(base_index)&3u;switch(base_index>>2){{",
+                        vertex_arg)};
+        if (info.loads.AnyComponent(IR::Attribute::PositionX)) {
+            const auto position_idx{is_array ? "gl_in[vertex]." : ""};
+            func += fmt::format("case {}:return {}{}[masked_index];",
+                                static_cast<u32>(IR::Attribute::PositionX) >> 2, position_idx,
+                                position_name);
+        }
+        const u32 base_attribute_value = static_cast<u32>(IR::Attribute::Generic0X) >> 2;
+        for (u32 index = 0; index < IR::NUM_GENERICS; ++index) {
+            if (!info.loads.Generic(index)) {
+                continue;
+            }
+            const auto vertex_idx{is_array ? "[vertex]" : ""};
+            func += fmt::format("case {}:return in_attr{}{}[masked_index];",
+                                base_attribute_value + index, index, vertex_idx);
+        }
+        func += "default: return 0.0;}}";
+        header += func;
+    }
+    if (info.stores_indexed_attributes) {
+        // TODO
+    }
+}
+
+std::string EmitContext::DefineGlobalMemoryFunctions() {
+    const auto define_body{[&](std::string& func, size_t index, std::string_view return_statement) {
+        const auto& ssbo{info.storage_buffers_descriptors[index]};
+        const u32 size_cbuf_offset{ssbo.cbuf_offset + 8};
+        const auto ssbo_addr{fmt::format("ssbo_addr{}", index)};
+        const auto cbuf{fmt::format("{}_cbuf{}", stage_name, ssbo.cbuf_index)};
+        std::array<std::string, 2> addr_xy;
+        std::array<std::string, 2> size_xy;
+        for (size_t i = 0; i < addr_xy.size(); ++i) {
+            const auto addr_loc{ssbo.cbuf_offset + 4 * i};
+            const auto size_loc{size_cbuf_offset + 4 * i};
+            addr_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, addr_loc / 16, Swizzle(addr_loc));
+            size_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, size_loc / 16, Swizzle(size_loc));
+        }
+        const auto addr_pack{fmt::format("packUint2x32(uvec2({},{}))", addr_xy[0], addr_xy[1])};
+        const auto addr_statment{fmt::format("uint64_t {}={};", ssbo_addr, addr_pack)};
+        func += addr_statment;
+
+        const auto size_vec{fmt::format("uvec2({},{})", size_xy[0], size_xy[1])};
+        const auto comp_lhs{fmt::format("(addr>={})", ssbo_addr)};
+        const auto comp_rhs{fmt::format("(addr<({}+uint64_t({})))", ssbo_addr, size_vec)};
+        const auto comparison{fmt::format("if({}&&{}){{", comp_lhs, comp_rhs)};
+        func += comparison;
+
+        const auto ssbo_name{fmt::format("{}_ssbo{}", stage_name, index)};
+        func += fmt::format(fmt::runtime(return_statement), ssbo_name, ssbo_addr);
+    }};
+    std::string write_func{"void WriteGlobal32(uint64_t addr,uint data){"};
+    std::string write_func_64{"void WriteGlobal64(uint64_t addr,uvec2 data){"};
+    std::string write_func_128{"void WriteGlobal128(uint64_t addr,uvec4 data){"};
+    std::string load_func{"uint LoadGlobal32(uint64_t addr){"};
+    std::string load_func_64{"uvec2 LoadGlobal64(uint64_t addr){"};
+    std::string load_func_128{"uvec4 LoadGlobal128(uint64_t addr){"};
+    const size_t num_buffers{info.storage_buffers_descriptors.size()};
+    for (size_t index = 0; index < num_buffers; ++index) {
+        if (!info.nvn_buffer_used[index]) {
+            continue;
+        }
+        define_body(write_func, index, "{0}[uint(addr-{1})>>2]=data;return;}}");
+        define_body(write_func_64, index,
+                    "{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;return;}}");
+        define_body(write_func_128, index,
+                    "{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;{0}[uint("
+                    "addr-{1}+8)>>2]=data.z;{0}[uint(addr-{1}+12)>>2]=data.w;return;}}");
+        define_body(load_func, index, "return {0}[uint(addr-{1})>>2];}}");
+        define_body(load_func_64, index,
+                    "return uvec2({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2]);}}");
+        define_body(load_func_128, index,
+                    "return uvec4({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2],{0}["
+                    "uint(addr-{1}+8)>>2],{0}[uint(addr-{1}+12)>>2]);}}");
+    }
+    write_func += '}';
+    write_func_64 += '}';
+    write_func_128 += '}';
+    load_func += "return 0u;}";
+    load_func_64 += "return uvec2(0);}";
+    load_func_128 += "return uvec4(0);}";
+    return write_func + write_func_64 + write_func_128 + load_func + load_func_64 + load_func_128;
+}
+
+void EmitContext::SetupImages(Bindings& bindings) {
+    image_buffers.reserve(info.image_buffer_descriptors.size());
+    for (const auto& desc : info.image_buffer_descriptors) {
+        image_buffers.push_back({bindings.image, desc.count});
+        const auto format{ImageFormatString(desc.format)};
+        const auto qualifier{ImageAccessQualifier(desc.is_written, desc.is_read)};
+        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
+        header += fmt::format("layout(binding={}{}) uniform {}uimageBuffer img{}{};",
+                              bindings.image, format, qualifier, bindings.image, array_decorator);
+        bindings.image += desc.count;
+    }
+    images.reserve(info.image_descriptors.size());
+    for (const auto& desc : info.image_descriptors) {
+        images.push_back({bindings.image, desc.count});
+        const auto format{ImageFormatString(desc.format)};
+        const auto image_type{ImageType(desc.type)};
+        const auto qualifier{ImageAccessQualifier(desc.is_written, desc.is_read)};
+        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
+        header += fmt::format("layout(binding={}{})uniform {}{} img{}{};", bindings.image, format,
+                              qualifier, image_type, bindings.image, array_decorator);
+        bindings.image += desc.count;
+    }
+}
+
+void EmitContext::SetupTextures(Bindings& bindings) {
+    texture_buffers.reserve(info.texture_buffer_descriptors.size());
+    for (const auto& desc : info.texture_buffer_descriptors) {
+        texture_buffers.push_back({bindings.texture, desc.count});
+        const auto sampler_type{SamplerType(TextureType::Buffer, false)};
+        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
+        header += fmt::format("layout(binding={}) uniform {} tex{}{};", bindings.texture,
+                              sampler_type, bindings.texture, array_decorator);
+        bindings.texture += desc.count;
+    }
+    textures.reserve(info.texture_descriptors.size());
+    for (const auto& desc : info.texture_descriptors) {
+        textures.push_back({bindings.texture, desc.count});
+        const auto sampler_type{SamplerType(desc.type, desc.is_depth)};
+        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
+        header += fmt::format("layout(binding={}) uniform {} tex{}{};", bindings.texture,
+                              sampler_type, bindings.texture, array_decorator);
+        bindings.texture += desc.count;
+    }
+}
+
+void EmitContext::DefineConstants() {
+    if (info.uses_fswzadd) {
+        header += "const float FSWZ_A[]=float[4](-1.f,1.f,-1.f,0.f);"
+                  "const float FSWZ_B[]=float[4](-1.f,-1.f,1.f,-1.f);";
+    }
+}
+
+} // namespace Shader::Backend::GLSL
--- a/src/shader_recompiler/backend/glsl/glsl_emit_context.h
+++ b/src/shader_recompiler/backend/glsl/glsl_emit_context.h
@@ -0,0 +1,174 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#pragma once
+
+#include <string>
+#include <utility>
+#include <vector>
+
+#include <fmt/format.h>
+
+#include "shader_recompiler/backend/glsl/var_alloc.h"
+#include "shader_recompiler/stage.h"
+
+namespace Shader {
+struct Info;
+struct Profile;
+struct RuntimeInfo;
+} // namespace Shader
+
+namespace Shader::Backend {
+struct Bindings;
+}
+
+namespace Shader::IR {
+class Inst;
+struct Program;
+} // namespace Shader::IR
+
+namespace Shader::Backend::GLSL {
+
+struct GenericElementInfo {
+    std::string name;
+    u32 first_element{};
+    u32 num_components{};
+};
+
+struct TextureImageDefinition {
+    u32 binding;
+    u32 count;
+};
+
+class EmitContext {
+public:
+    explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
+                         const RuntimeInfo& runtime_info_);
+
+    template <GlslVarType type, typename... Args>
+    void Add(const char* format_str, IR::Inst& inst, Args&&... args) {
+        const auto var_def{var_alloc.AddDefine(inst, type)};
+        if (var_def.empty()) {
+            // skip assigment.
+            code += fmt::format(fmt::runtime(format_str + 3), std::forward<Args>(args)...);
+        } else {
+            code += fmt::format(fmt::runtime(format_str), var_def, std::forward<Args>(args)...);
+        }
+        // TODO: Remove this
+        code += '\n';
+    }
+
+    template <typename... Args>
+    void AddU1(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::U1>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddF16x2(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::F16x2>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddU32(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::U32>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddF32(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::F32>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddU64(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::U64>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddF64(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::F64>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddU32x2(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::U32x2>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddF32x2(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::F32x2>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddU32x3(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::U32x3>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddF32x3(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::F32x3>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddU32x4(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::U32x4>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddF32x4(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::F32x4>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddPrecF32(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::PrecF32>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void AddPrecF64(const char* format_str, IR::Inst& inst, Args&&... args) {
+        Add<GlslVarType::PrecF64>(format_str, inst, args...);
+    }
+
+    template <typename... Args>
+    void Add(const char* format_str, Args&&... args) {
+        code += fmt::format(fmt::runtime(format_str), std::forward<Args>(args)...);
+        // TODO: Remove this
+        code += '\n';
+    }
+
+    std::string header;
+    std::string code;
+    VarAlloc var_alloc;
+    const Info& info;
+    const Profile& profile;
+    const RuntimeInfo& runtime_info;
+
+    Stage stage{};
+    std::string_view stage_name = "invalid";
+    std::string_view position_name = "gl_Position";
+
+    std::vector<TextureImageDefinition> texture_buffers;
+    std::vector<TextureImageDefinition> image_buffers;
+    std::vector<TextureImageDefinition> textures;
+    std::vector<TextureImageDefinition> images;
+    std::array<std::array<GenericElementInfo, 4>, 32> output_generics{};
+
+    u32 num_safety_loop_vars{};
+
+    bool uses_y_direction{};
+    bool uses_cc_carry{};
+    bool uses_geometry_passthrough{};
+
+private:
+    void SetupExtensions();
+    void DefineConstantBuffers(Bindings& bindings);
+    void DefineStorageBuffers(Bindings& bindings);
+    void DefineGenericOutput(size_t index, u32 invocations);
+    void DefineHelperFunctions();
+    void DefineConstants();
+    std::string DefineGlobalMemoryFunctions();
+    void SetupImages(Bindings& bindings);
+    void SetupTextures(Bindings& bindings);
+};
+
+} // namespace Shader::Backend::GLSL
--- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
@@ -11,6 +11,7 @@
 #include "common/settings.h"
 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
 #include "shader_recompiler/frontend/ir/basic_block.h"
 #include "shader_recompiler/frontend/ir/program.h"

--- a/src/shader_recompiler/backend/spirv/emit_spirv.h
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.h
@@ -6,13 +6,11 @@

 #include <vector>

-#include <sirit/sirit.h>
-
 #include "common/common_types.h"
 #include "shader_recompiler/backend/bindings.h"
-#include "shader_recompiler/backend/spirv/emit_context.h"
 #include "shader_recompiler/frontend/ir/program.h"
 #include "shader_recompiler/profile.h"
+#include "shader_recompiler/runtime_info.h"

 namespace Shader::Backend::SPIRV {

--- a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {
 namespace {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_barriers.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_barriers.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"

 namespace Shader::Backend::SPIRV {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_bitwise_conversion.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_bitwise_conversion.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {

--- a/src/shader_recompiler/backend/spirv/emit_spirv_composite.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_composite.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"

 namespace Shader::Backend::SPIRV {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
@@ -7,6 +7,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {
 namespace {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {

--- a/src/shader_recompiler/backend/spirv/emit_spirv_convert.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_convert.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {
 namespace {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"

 namespace Shader::Backend::SPIRV {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_image.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_image.cpp
@@ -6,6 +6,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"

 namespace Shader::Backend::SPIRV {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_image_atomic.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_image_atomic.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"

 namespace Shader::Backend::SPIRV {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_integer.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_integer.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {
 namespace {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_logical.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_logical.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {

--- a/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp
@@ -6,6 +6,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {
 namespace {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_select.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_select.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {

--- a/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {
 namespace {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {
 namespace {
--- a/src/shader_recompiler/backend/spirv/emit_spirv_undefined.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_undefined.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {

--- a/src/shader_recompiler/backend/spirv/emit_spirv_warp.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_warp.cpp
@@ -4,6 +4,7 @@

 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
+#include "shader_recompiler/backend/spirv/spirv_emit_context.h"

 namespace Shader::Backend::SPIRV {
 namespace {
--- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
+++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
--- a/src/shader_recompiler/backend/spirv/spirv_emit_context.h
+++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h
@@ -0,0 +1,320 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#pragma once
+
+#include <array>
+#include <string_view>
+
+#include <sirit/sirit.h>
+
+#include "shader_recompiler/backend/bindings.h"
+#include "shader_recompiler/frontend/ir/program.h"
+#include "shader_recompiler/profile.h"
+#include "shader_recompiler/runtime_info.h"
+#include "shader_recompiler/shader_info.h"
+
+namespace Shader::Backend::SPIRV {
+
+using Sirit::Id;
+
+class VectorTypes {
+public:
+    void Define(Sirit::Module& sirit_ctx, Id base_type, std::string_view name);
+
+    [[nodiscard]] Id operator[](size_t size) const noexcept {
+        return defs[size - 1];
+    }
+
+private:
+    std::array<Id, 4> defs{};
+};
+
+struct TextureDefinition {
+    Id id;
+    Id sampled_type;
+    Id pointer_type;
+    Id image_type;
+    u32 count;
+};
+
+struct TextureBufferDefinition {
+    Id id;
+    u32 count;
+};
+
+struct ImageBufferDefinition {
+    Id id;
+    Id image_type;
+    u32 count;
+};
+
+struct ImageDefinition {
+    Id id;
+    Id image_type;
+    u32 count;
+};
+
+struct UniformDefinitions {
+    Id U8{};
+    Id S8{};
+    Id U16{};
+    Id S16{};
+    Id U32{};
+    Id F32{};
+    Id U32x2{};
+    Id U32x4{};
+};
+
+struct StorageTypeDefinition {
+    Id array{};
+    Id element{};
+};
+
+struct StorageTypeDefinitions {
+    StorageTypeDefinition U8{};
+    StorageTypeDefinition S8{};
+    StorageTypeDefinition U16{};
+    StorageTypeDefinition S16{};
+    StorageTypeDefinition U32{};
+    StorageTypeDefinition U64{};
+    StorageTypeDefinition F32{};
+    StorageTypeDefinition U32x2{};
+    StorageTypeDefinition U32x4{};
+};
+
+struct StorageDefinitions {
+    Id U8{};
+    Id S8{};
+    Id U16{};
+    Id S16{};
+    Id U32{};
+    Id F32{};
+    Id U64{};
+    Id U32x2{};
+    Id U32x4{};
+};
+
+struct GenericElementInfo {
+    Id id{};
+    u32 first_element{};
+    u32 num_components{};
+};
+
+class EmitContext final : public Sirit::Module {
+public:
+    explicit EmitContext(const Profile& profile, const RuntimeInfo& runtime_info,
+                         IR::Program& program, Bindings& binding);
+    ~EmitContext();
+
+    [[nodiscard]] Id Def(const IR::Value& value);
+
+    [[nodiscard]] Id BitOffset8(const IR::Value& offset);
+    [[nodiscard]] Id BitOffset16(const IR::Value& offset);
+
+    Id Const(u32 value) {
+        return Constant(U32[1], value);
+    }
+
+    Id Const(u32 element_1, u32 element_2) {
+        return ConstantComposite(U32[2], Const(element_1), Const(element_2));
+    }
+
+    Id Const(u32 element_1, u32 element_2, u32 element_3) {
+        return ConstantComposite(U32[3], Const(element_1), Const(element_2), Const(element_3));
+    }
+
+    Id Const(u32 element_1, u32 element_2, u32 element_3, u32 element_4) {
+        return ConstantComposite(U32[4], Const(element_1), Const(element_2), Const(element_3),
+                                 Const(element_4));
+    }
+
+    Id SConst(s32 value) {
+        return Constant(S32[1], value);
+    }
+
+    Id SConst(s32 element_1, s32 element_2) {
+        return ConstantComposite(S32[2], SConst(element_1), SConst(element_2));
+    }
+
+    Id SConst(s32 element_1, s32 element_2, s32 element_3) {
+        return ConstantComposite(S32[3], SConst(element_1), SConst(element_2), SConst(element_3));
+    }
+
+    Id SConst(s32 element_1, s32 element_2, s32 element_3, s32 element_4) {
+        return ConstantComposite(S32[4], SConst(element_1), SConst(element_2), SConst(element_3),
+                                 SConst(element_4));
+    }
+
+    Id Const(f32 value) {
+        return Constant(F32[1], value);
+    }
+
+    const Profile& profile;
+    const RuntimeInfo& runtime_info;
+    Stage stage{};
+
+    Id void_id{};
+    Id U1{};
+    Id U8{};
+    Id S8{};
+    Id U16{};
+    Id S16{};
+    Id U64{};
+    VectorTypes F32;
+    VectorTypes U32;
+    VectorTypes S32;
+    VectorTypes F16;
+    VectorTypes F64;
+
+    Id true_value{};
+    Id false_value{};
+    Id u32_zero_value{};
+    Id f32_zero_value{};
+
+    UniformDefinitions uniform_types;
+    StorageTypeDefinitions storage_types;
+
+    Id private_u32{};
+
+    Id shared_u8{};
+    Id shared_u16{};
+    Id shared_u32{};
+    Id shared_u64{};
+    Id shared_u32x2{};
+    Id shared_u32x4{};
+
+    Id input_f32{};
+    Id input_u32{};
+    Id input_s32{};
+
+    Id output_f32{};
+    Id output_u32{};
+
+    Id image_buffer_type{};
+    Id sampled_texture_buffer_type{};
+    Id image_u32{};
+
+    std::array<UniformDefinitions, Info::MAX_CBUFS> cbufs{};
+    std::array<StorageDefinitions, Info::MAX_SSBOS> ssbos{};
+    std::vector<TextureBufferDefinition> texture_buffers;
+    std::vector<ImageBufferDefinition> image_buffers;
+    std::vector<TextureDefinition> textures;
+    std::vector<ImageDefinition> images;
+
+    Id workgroup_id{};
+    Id local_invocation_id{};
+    Id invocation_id{};
+    Id sample_id{};
+    Id is_helper_invocation{};
+    Id subgroup_local_invocation_id{};
+    Id subgroup_mask_eq{};
+    Id subgroup_mask_lt{};
+    Id subgroup_mask_le{};
+    Id subgroup_mask_gt{};
+    Id subgroup_mask_ge{};
+    Id instance_id{};
+    Id instance_index{};
+    Id base_instance{};
+    Id vertex_id{};
+    Id vertex_index{};
+    Id base_vertex{};
+    Id front_face{};
+    Id point_coord{};
+    Id tess_coord{};
+    Id clip_distances{};
+    Id layer{};
+    Id viewport_index{};
+    Id viewport_mask{};
+    Id primitive_id{};
+
+    Id fswzadd_lut_a{};
+    Id fswzadd_lut_b{};
+
+    Id indexed_load_func{};
+    Id indexed_store_func{};
+
+    Id rescaling_uniform_constant{};
+    Id rescaling_push_constants{};
+    Id rescaling_textures_type{};
+    Id rescaling_images_type{};
+    u32 rescaling_textures_member_index{};
+    u32 rescaling_images_member_index{};
+    u32 rescaling_downfactor_member_index{};
+    u32 texture_rescaling_index{};
+    u32 image_rescaling_index{};
+
+    Id local_memory{};
+
+    Id shared_memory_u8{};
+    Id shared_memory_u16{};
+    Id shared_memory_u32{};
+    Id shared_memory_u64{};
+    Id shared_memory_u32x2{};
+    Id shared_memory_u32x4{};
+
+    Id shared_memory_u32_type{};
+
+    Id shared_store_u8_func{};
+    Id shared_store_u16_func{};
+    Id increment_cas_shared{};
+    Id increment_cas_ssbo{};
+    Id decrement_cas_shared{};
+    Id decrement_cas_ssbo{};
+    Id f32_add_cas{};
+    Id f16x2_add_cas{};
+    Id f16x2_min_cas{};
+    Id f16x2_max_cas{};
+    Id f32x2_add_cas{};
+    Id f32x2_min_cas{};
+    Id f32x2_max_cas{};
+
+    Id load_global_func_u32{};
+    Id load_global_func_u32x2{};
+    Id load_global_func_u32x4{};
+    Id write_global_func_u32{};
+    Id write_global_func_u32x2{};
+    Id write_global_func_u32x4{};
+
+    Id input_position{};
+    std::array<Id, 32> input_generics{};
+
+    Id output_point_size{};
+    Id output_position{};
+    std::array<std::array<GenericElementInfo, 4>, 32> output_generics{};
+
+    Id output_tess_level_outer{};
+    Id output_tess_level_inner{};
+    std::array<Id, 30> patches{};
+
+    std::array<Id, 8> frag_color{};
+    Id sample_mask{};
+    Id frag_depth{};
+
+    std::vector<Id> interfaces;
+
+private:
+    void DefineCommonTypes(const Info& info);
+    void DefineCommonConstants();
+    void DefineInterfaces(const IR::Program& program);
+    void DefineLocalMemory(const IR::Program& program);
+    void DefineSharedMemory(const IR::Program& program);
+    void DefineSharedMemoryFunctions(const IR::Program& program);
+    void DefineConstantBuffers(const Info& info, u32& binding);
+    void DefineStorageBuffers(const Info& info, u32& binding);
+    void DefineTextureBuffers(const Info& info, u32& binding);
+    void DefineImageBuffers(const Info& info, u32& binding);
+    void DefineTextures(const Info& info, u32& binding, u32& scaling_index);
+    void DefineImages(const Info& info, u32& binding, u32& scaling_index);
+    void DefineAttributeMemAccess(const Info& info);
+    void DefineGlobalMemoryFunctions(const Info& info);
+    void DefineRescalingInput(const Info& info);
+    void DefineRescalingInputPushConstant();
+    void DefineRescalingInputUniformConstant();
+
+    void DefineInputs(const IR::Program& program);
+    void DefineOutputs(const IR::Program& program);
+};
+
+} // namespace Shader::Backend::SPIRV