diff --git a/src/shader_recompiler/backend/msl/emit_msl.cpp b/src/shader_recompiler/backend/msl/emit_msl.cpp
index 45c2141528..ce6b902c53 100644
--- a/src/shader_recompiler/backend/msl/emit_msl.cpp
+++ b/src/shader_recompiler/backend/msl/emit_msl.cpp
@@ -205,213 +205,12 @@ void DefineVariables(const EmitContext& ctx, std::string& header) {
         header += fmt::format("int loop{}=0x2000;", i);
     }
 }
-
-// TODO
-std::string_view DepthSamplerType(TextureType type) {
-    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);
-    }
-}
-
-// TODO: emit sampler as well
-// TODO: handle multisample
-// TODO: handle texture buffer
-std::string_view ColorSamplerType(TextureType type, bool is_multisample = false) {
-    if (is_multisample) {
-        ASSERT(type == TextureType::Color2D || type == TextureType::ColorArray2D);
-    }
-    switch (type) {
-    case TextureType::Color1D:
-        return "texture1d";
-    case TextureType::ColorArray1D:
-        return "texture1d_array";
-    case TextureType::Color2D:
-    case TextureType::Color2DRect:
-        return "texture2d";
-    case TextureType::ColorArray2D:
-        return "texture2d_array";
-    case TextureType::Color3D:
-        return "texture3d";
-    case TextureType::ColorCube:
-        return "texturecube";
-    case TextureType::ColorArrayCube:
-        return "texturecube_array";
-    default:
-        throw NotImplementedException("Texture type: {}", type);
-    }
-}
-
-// TODO: handle texture buffer
-std::string_view ImageType(TextureType type) {
-    switch (type) {
-    case TextureType::Color1D:
-        return "texture1d";
-    case TextureType::ColorArray1D:
-        return "texture1d_array";
-    case TextureType::Color2D:
-        return "texture2d";
-    case TextureType::ColorArray2D:
-        return "texture2d_array";
-    case TextureType::Color3D:
-        return "texture3d";
-    case TextureType::ColorCube:
-        return "texturecube";
-    case TextureType::ColorArrayCube:
-        return "texturecube_array";
-    default:
-        throw NotImplementedException("Image type: {}", type);
-    }
-}
-
-// TODO: is this needed?
-/*
-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 "access::read, access::write";
-    }
-    if (is_written) {
-        return "access::write";
-    }
-    if (is_read) {
-        return "access::read";
-    }
-    return "";
-}
-
-void DefineInputs(EmitContext& ctx, std::string& header, Bindings& bindings) {
-    bool added{false};
-
-    // Constant buffers
-    for (const auto& desc : ctx.info.constant_buffer_descriptors) {
-        const u32 cbuf_used_size{
-            Common::DivCeil(ctx.info.constant_buffer_used_sizes[desc.index], 16U)};
-        const u32 cbuf_binding_size{ctx.info.uses_global_memory ? 0x1000U : cbuf_used_size};
-        if (added)
-            header += ",";
-        header += fmt::format("constant float4& cbuf{}[{}] [[buffer({})]]", desc.index,
-                              cbuf_binding_size, bindings.uniform_buffer);
-        bindings.uniform_buffer += desc.count;
-        added = true;
-    }
-
-    // Constant buffer indirect
-    // TODO
-
-    // Storage space buffers
-    u32 index{};
-    for (const auto& desc : ctx.info.storage_buffers_descriptors) {
-        if (added)
-            header += ",";
-        header +=
-            fmt::format("device uint& ssbo{}[] [[buffer({})]]", index, bindings.storage_buffer);
-        bindings.storage_buffer += desc.count;
-        index += desc.count;
-        added = true;
-    }
-
-    // Images
-    // TODO
-    /*
-    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;
-    }
-    */
-    ctx.images.reserve(ctx.info.image_descriptors.size());
-    for (const auto& desc : ctx.info.image_descriptors) {
-        ctx.images.push_back({bindings.image, desc.count});
-        // TODO: do we need format?
-        // 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) : ""};
-        if (added)
-            header += ",";
-        header += fmt::format("{}<{}> img{}{} [[texture({})]]", qualifier, image_type,
-                              bindings.image, array_decorator, bindings.image);
-        bindings.image += desc.count;
-        added = true;
-    }
-
-    // Textures
-    // TODO
-    /*
-    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{ColorSamplerType(TextureType::Buffer)};
-        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;
-    }
-    */
-    ctx.textures.reserve(ctx.info.texture_descriptors.size());
-    for (const auto& desc : ctx.info.texture_descriptors) {
-        ctx.textures.push_back({bindings.texture, desc.count});
-        const auto texture_type{desc.is_depth ? DepthSamplerType(desc.type)
-                                              : ColorSamplerType(desc.type, desc.is_multisample)};
-        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
-        if (added)
-            header += ",";
-        header += fmt::format("{} tex{}{} [[texture({})]]", texture_type, bindings.texture,
-                              array_decorator, bindings.texture);
-        header += fmt::format(",sampler samp{}{} [[sampler({})]]", bindings.texture,
-                              array_decorator, bindings.texture);
-        bindings.texture += desc.count;
-        added = true;
-    }
-}
 } // Anonymous namespace
 
 std::string EmitMSL(const Profile& profile, const RuntimeInfo& runtime_info, IR::Program& program,
                     Bindings& bindings) {
     EmitContext ctx{program, bindings, profile, runtime_info};
     std::string inputs;
-    DefineInputs(ctx, inputs, bindings);
     Precolor(program);
     EmitCode(ctx, program);
     ctx.header.insert(0, "#include <metal_stdlib>\nusing namespace metal;\n");
@@ -426,8 +225,9 @@ std::string EmitMSL(const Profile& profile, const RuntimeInfo& runtime_info, IR:
         const auto smem_size{needs_clamp ? max_size : requested_size};
         ctx.header += fmt::format("shared uint smem[{}];", Common::DivCeil(smem_size, 4U));
     }
-    ctx.header += "void main_(";
-    ctx.header += inputs;
+    ctx.header += ctx.stage_name;
+    ctx.header += " __Output main_(";
+    ctx.header += ctx.input_str;
     ctx.header += "){\n";
     if (program.local_memory_size > 0) {
         ctx.header += fmt::format("uint lmem[{}];", Common::DivCeil(program.local_memory_size, 4U));
diff --git a/src/shader_recompiler/backend/msl/msl_emit_context.cpp b/src/shader_recompiler/backend/msl/msl_emit_context.cpp
index 9dc25df7a9..203b87fdfd 100644
--- a/src/shader_recompiler/backend/msl/msl_emit_context.cpp
+++ b/src/shader_recompiler/backend/msl/msl_emit_context.cpp
@@ -107,6 +107,115 @@ std::string_view OutputPrimitive(OutputTopology topology) {
     }
     throw InvalidArgument("Invalid output topology {}", topology);
 }
+
+// TODO
+std::string_view DepthSamplerType(TextureType type) {
+    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);
+    }
+}
+
+// TODO: emit sampler as well
+// TODO: handle multisample
+// TODO: handle texture buffer
+std::string_view ColorSamplerType(TextureType type, bool is_multisample = false) {
+    if (is_multisample) {
+        ASSERT(type == TextureType::Color2D || type == TextureType::ColorArray2D);
+    }
+    switch (type) {
+    case TextureType::Color1D:
+        return "texture1d";
+    case TextureType::ColorArray1D:
+        return "texture1d_array";
+    case TextureType::Color2D:
+    case TextureType::Color2DRect:
+        return "texture2d";
+    case TextureType::ColorArray2D:
+        return "texture2d_array";
+    case TextureType::Color3D:
+        return "texture3d";
+    case TextureType::ColorCube:
+        return "texturecube";
+    case TextureType::ColorArrayCube:
+        return "texturecube_array";
+    default:
+        throw NotImplementedException("Texture type: {}", type);
+    }
+}
+
+// TODO: handle texture buffer
+std::string_view ImageType(TextureType type) {
+    switch (type) {
+    case TextureType::Color1D:
+        return "texture1d";
+    case TextureType::ColorArray1D:
+        return "texture1d_array";
+    case TextureType::Color2D:
+        return "texture2d";
+    case TextureType::ColorArray2D:
+        return "texture2d_array";
+    case TextureType::Color3D:
+        return "texture3d";
+    case TextureType::ColorCube:
+        return "texturecube";
+    case TextureType::ColorArrayCube:
+        return "texturecube_array";
+    default:
+        throw NotImplementedException("Image type: {}", type);
+    }
+}
+
+// TODO: is this needed?
+/*
+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 "access::read, access::write";
+    }
+    if (is_written) {
+        return "access::write";
+    }
+    if (is_read) {
+        return "access::read";
+    }
+    return "";
+}
 } // Anonymous namespace
 
 EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
@@ -120,20 +229,20 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
     switch (program.stage) {
     case Stage::VertexA:
     case Stage::VertexB:
-        stage_name = "vs";
+        stage_name = "vertex";
         break;
     case Stage::TessellationControl:
-        stage_name = "tcs";
+        stage_name = "kernel";
         header += fmt::format("layout(vertices={})out;", program.invocations);
         break;
     case Stage::TessellationEval:
-        stage_name = "tes";
+        stage_name = "vertex";
         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";
+        stage_name = "vertex";
         header += fmt::format("layout({})in;", InputPrimitive(runtime_info.input_topology));
         if (uses_geometry_passthrough) {
             header += "layout(passthrough)in gl_PerVertex{vec4 gl_Position;};";
@@ -147,13 +256,13 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
             OutputPrimitive(program.output_topology), program.output_vertices);
         break;
     case Stage::Fragment:
-        stage_name = "fs";
+        stage_name = "fragment";
         if (runtime_info.force_early_z) {
             header += "layout(early_fragment_tests)in;";
         }
         break;
     case Stage::Compute:
-        stage_name = "cs";
+        stage_name = "kernel";
         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)};
@@ -190,21 +299,122 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
             header += fmt::format("layout(location={})out vec4 frag_color{};", index, index);
         }
     }
+    header += "struct __Output {\n";
     for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
         if (info.stores.Generic(index)) {
             DefineGenericOutput(index, program.invocations);
         }
     }
+    header += "};\n";
+    bool added = DefineInputs(bindings);
     if (info.uses_rescaling_uniform) {
-        header += "layout(location=0) uniform vec4 scaling;";
+        if (added)
+            input_str += ",";
+        input_str += "constant float4& scaling";
+        added = true;
     }
     if (info.uses_render_area) {
-        header += "layout(location=1) uniform vec4 render_area;";
+        if (added)
+            input_str += ",";
+        input_str += "constant float4& render_area";
+        added = true;
     }
     DefineHelperFunctions();
     DefineConstants();
 }
 
+bool EmitContext::DefineInputs(Bindings& bindings) {
+    bool added{false};
+
+    // Constant buffers
+    for (const auto& desc : info.constant_buffer_descriptors) {
+        const u32 cbuf_used_size{Common::DivCeil(info.constant_buffer_used_sizes[desc.index], 16U)};
+        const u32 cbuf_binding_size{info.uses_global_memory ? 0x1000U : cbuf_used_size};
+        if (added)
+            input_str += ",";
+        input_str += fmt::format("constant float4& cbuf{}[{}] [[buffer({})]]", desc.index,
+                                 cbuf_binding_size, bindings.uniform_buffer);
+        bindings.uniform_buffer += desc.count;
+        added = true;
+    }
+
+    // Constant buffer indirect
+    // TODO
+
+    // Storage space buffers
+    u32 index{};
+    for (const auto& desc : info.storage_buffers_descriptors) {
+        if (added)
+            input_str += ",";
+        input_str +=
+            fmt::format("device uint& ssbo{}[] [[buffer({})]]", index, bindings.storage_buffer);
+        bindings.storage_buffer += desc.count;
+        index += desc.count;
+        added = true;
+    }
+
+    // Images
+    // TODO
+    /*
+    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) : ""};
+        input_str += 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});
+        // TODO: do we need format?
+        // 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) : ""};
+        if (added)
+            input_str += ",";
+        input_str += fmt::format("{}<{}> img{}{} [[texture({})]]", qualifier, image_type,
+                                 bindings.image, array_decorator, bindings.image);
+        bindings.image += desc.count;
+        added = true;
+    }
+
+    // Textures
+    // TODO
+    /*
+    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{ColorSamplerType(TextureType::Buffer)};
+        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
+        input_str += 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 texture_type{desc.is_depth ? DepthSamplerType(desc.type)
+                                              : ColorSamplerType(desc.type, desc.is_multisample)};
+        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
+        if (added)
+            input_str += ",";
+        input_str += fmt::format("{} tex{}{} [[texture({})]]", texture_type, bindings.texture,
+                                 array_decorator, bindings.texture);
+        input_str += fmt::format(",sampler samp{}{} [[sampler({})]]", bindings.texture,
+                                 array_decorator, bindings.texture);
+        bindings.texture += desc.count;
+        added = true;
+    }
+
+    return added;
+}
+
 // TODO
 void EmitContext::DefineGenericOutput(size_t index, u32 invocations) {
     static constexpr std::string_view swizzle{"xyzw"};
diff --git a/src/shader_recompiler/backend/msl/msl_emit_context.h b/src/shader_recompiler/backend/msl/msl_emit_context.h
index 233b4bd606..830e97ce9a 100644
--- a/src/shader_recompiler/backend/msl/msl_emit_context.h
+++ b/src/shader_recompiler/backend/msl/msl_emit_context.h
@@ -136,6 +136,7 @@ public:
     }
 
     std::string header;
+    std::string input_str;
     std::string code;
     VarAlloc var_alloc;
     const Info& info;
@@ -158,6 +159,8 @@ public:
     bool uses_geometry_passthrough{};
 
 private:
+    // TODO: break down into smaller functions
+    bool DefineInputs(Bindings& bindings);
     void DefineGenericOutput(size_t index, u32 invocations);
     void DefineHelperFunctions();
     void DefineConstants();