shader_recompiler: translate Switch 2 task and mesh programs

Mario Kart World draws some of its geometry with NVN task and mesh programs
(the task program culls and writes the mesh workgroup count and a payload
to its stage buffer; each mesh workgroup reads it back and writes vertices
and primitives). Add Task and Mesh stages and translate them for
VK_EXT_mesh_shader:

- the ISBE pseudo instruction from the lowering becomes reads and writes
  of an emulated stage buffer: the task's output region is the task
  payload, the mesh output regions live in workgroup memory
- every EXIT of these stages jumps to one final return, so all
  invocations reach the epilogue: a barrier, then OpEmitMeshTasksEXT, or
  OpSetMeshOutputsEXT and a copy of the 32-element vertex and primitive
  groups (lines named by the mesh header's and the pass-through geometry
  header's output maps) to arrayed and per-primitive outputs
- Sirit lacks the two mesh instructions: they are emitted as OpStore
  markers and replaced after assembly
- fragment inputs a mesh program writes per primitive get PerPrimitiveEXT

The captured task and mesh programs of the title translate to SPIR-V that
spirv-val accepts.

30 files changed, 703 insertions, 8 deletions

diff --git a/src/shader_recompiler/CMakeLists.txt b/src/shader_recompiler/CMakeLists.txt
index e60f26d47f..e4df9c045d 100644
--- a/src/shader_recompiler/CMakeLists.txt
+++ b/src/shader_recompiler/CMakeLists.txt
@@ -24,6 +24,7 @@ add_library(shader_recompiler STATIC
     backend/spirv/emit_spirv_memory.cpp
     backend/spirv/emit_spirv_select.cpp
     backend/spirv/emit_spirv_shared_memory.cpp
+    backend/spirv/emit_spirv_mesh.cpp
     backend/spirv/emit_spirv_special.cpp
     backend/spirv/emit_spirv_undefined.cpp
     backend/spirv/emit_spirv_warp.cpp
diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
index b795c01797..c18eec63cd 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp
@@ -341,6 +341,9 @@ std::string_view StageHeader(Stage stage) {
         return "!!NVfp5.0\n";
     case Stage::Compute:
         return "!!NVcp5.0\n";
+    case Stage::Task:
+    case Stage::Mesh:
+        break;
     }
     throw InvalidArgument("Invalid stage {}", stage);
 }
diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_instructions.h b/src/shader_recompiler/backend/glasm/emit_glasm_instructions.h
index 6f522ab5c2..d66fe2f3d6 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm_instructions.h
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_instructions.h
@@ -58,6 +58,10 @@ void EmitGetAttributeIndexed(EmitContext& ctx, IR::Inst& inst, ScalarS32 offset,
 void EmitSetAttributeIndexed(EmitContext& ctx, ScalarU32 offset, ScalarF32 value, ScalarU32 vertex);
 void EmitGetPatch(EmitContext& ctx, IR::Inst& inst, IR::Patch patch);
 void EmitSetPatch(EmitContext& ctx, IR::Patch patch, ScalarF32 value);
+void EmitIsbeRead(EmitContext& ctx);
+void EmitIsbeWriteAttribute(EmitContext& ctx);
+void EmitIsbeWriteMapU8(EmitContext& ctx);
+void EmitIsbeWriteMapU32(EmitContext& ctx);
 void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, ScalarF32 value);
 void EmitSetSampleMask(EmitContext& ctx, ScalarS32 value);
 void EmitSetFragDepth(EmitContext& ctx, ScalarF32 value);
diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp
index 1f4ffdd624..41c829ca5d 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp
@@ -97,4 +97,20 @@ void EmitGetInBoundsFromOp(EmitContext& ctx) {
     NotImplemented();
 }
 
+void EmitIsbeRead(EmitContext& ctx) {
+    NotImplemented();
+}
+
+void EmitIsbeWriteAttribute(EmitContext& ctx) {
+    NotImplemented();
+}
+
+void EmitIsbeWriteMapU8(EmitContext& ctx) {
+    NotImplemented();
+}
+
+void EmitIsbeWriteMapU32(EmitContext& ctx) {
+    NotImplemented();
+}
+
 } // namespace Shader::Backend::GLASM
diff --git a/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp b/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp
index 333a91cc58..02b6992b52 100644
--- a/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp
+++ b/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp
@@ -80,6 +80,9 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
     case Stage::Compute:
         stage_name = "invocation";
         break;
+    case Stage::Task:
+    case Stage::Mesh:
+        throw NotImplementedException("GLASM task and mesh programs");
     }
     const std::string_view attr_stage{stage == Stage::Fragment ? "fragment" : "vertex"};
     const VaryingState loads{info.loads.mask | info.passthrough.mask};
diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_instructions.h b/src/shader_recompiler/backend/glsl/emit_glsl_instructions.h
index 96aac502d5..06069d4c56 100644
--- a/src/shader_recompiler/backend/glsl/emit_glsl_instructions.h
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_instructions.h
@@ -72,6 +72,10 @@ void EmitSetAttributeIndexed(EmitContext& ctx, std::string_view offset, std::str
                              std::string_view vertex);
 void EmitGetPatch(EmitContext& ctx, IR::Inst& inst, IR::Patch patch);
 void EmitSetPatch(EmitContext& ctx, IR::Patch patch, std::string_view value);
+void EmitIsbeRead(EmitContext& ctx);
+void EmitIsbeWriteAttribute(EmitContext& ctx);
+void EmitIsbeWriteMapU8(EmitContext& ctx);
+void EmitIsbeWriteMapU32(EmitContext& ctx);
 void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, std::string_view value);
 void EmitSetSampleMask(EmitContext& ctx, std::string_view value);
 void EmitSetFragDepth(EmitContext& ctx, std::string_view value);
diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp
index 9f1ed95a47..18779e51b8 100644
--- a/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp
@@ -97,4 +97,20 @@ void EmitGetInBoundsFromOp(EmitContext& ctx) {
     NotImplemented();
 }
 
+void EmitIsbeRead(EmitContext& ctx) {
+    NotImplemented();
+}
+
+void EmitIsbeWriteAttribute(EmitContext& ctx) {
+    NotImplemented();
+}
+
+void EmitIsbeWriteMapU8(EmitContext& ctx) {
+    NotImplemented();
+}
+
+void EmitIsbeWriteMapU32(EmitContext& ctx) {
+    NotImplemented();
+}
+
 } // namespace Shader::Backend::GLSL
diff --git a/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp b/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp
index a2df6159fb..e4b11d6844 100644
--- a/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp
+++ b/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp
@@ -317,6 +317,9 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
             header += "layout(early_fragment_tests)in;";
         }
         break;
+    case Stage::Task:
+    case Stage::Mesh:
+        throw NotImplementedException("GLSL task and mesh programs");
     case Stage::Compute:
         stage_name = "cs";
         const u32 local_x{(std::max)(program.workgroup_size[0], 1u)};
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
index a903480c78..633ab25935 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
@@ -248,6 +248,20 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
     case Stage::VertexB:
         execution_model = spv::ExecutionModel::Vertex;
         break;
+    case Stage::Task:
+        execution_model = spv::ExecutionModel::TaskEXT;
+        ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, program.workgroup_size[0], 1U,
+                             1U);
+        break;
+    case Stage::Mesh:
+        execution_model = spv::ExecutionModel::MeshEXT;
+        ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, program.workgroup_size[0], 1U,
+                             1U);
+        ctx.AddExecutionMode(main, spv::ExecutionMode::OutputVertices, program.mesh.max_vertices);
+        ctx.AddExecutionMode(main, spv::ExecutionMode::OutputPrimitivesEXT,
+                             program.mesh.max_primitives);
+        ctx.AddExecutionMode(main, spv::ExecutionMode::OutputTrianglesEXT);
+        break;
     case Stage::TessellationControl:
         execution_model = spv::ExecutionModel::TessellationControl;
         ctx.AddCapability(spv::Capability::Tessellation);
@@ -548,6 +562,51 @@ void RewriteOpcodes(std::vector<u32>& code, std::span<const std::pair<u32, spv::
         offset += word_count;
     }
 }
+
+// Replaces the OpStore markers of the mesh instructions Sirit cannot emit: OpStore(vertices,
+// primitives) becomes OpSetMeshOutputsEXT, and OpStore(x, y), OpStore(z, payload), OpReturn
+// becomes the OpEmitMeshTasksEXT terminator.
+void PatchMeshInstructions(std::vector<u32>& code, const EmitContext& ctx) {
+    const bool set_outputs{Sirit::ValidId(ctx.mesh_outputs_marker[0])};
+    const bool emit_tasks{Sirit::ValidId(ctx.mesh_tasks_marker[0])};
+    if (!set_outputs && !emit_tasks) {
+        return;
+    }
+    const auto store_at = [&](size_t offset, Id pointer, Id object) {
+        return offset + 2 < code.size() &&
+               code[offset] == ((3U << 16) | static_cast<u32>(spv::Op::OpStore)) &&
+               code[offset + 1] == pointer.value && code[offset + 2] == object.value;
+    };
+    const auto& tasks{ctx.mesh_tasks_marker};
+    size_t offset = 5;
+    bool patched{};
+    while (offset < code.size()) {
+        const u32 word_count{code[offset] >> 16};
+        if (word_count == 0) {
+            break;
+        }
+        if (set_outputs &&
+            store_at(offset, ctx.mesh_outputs_marker[0], ctx.mesh_outputs_marker[1])) {
+            code[offset] = (3U << 16) | static_cast<u32>(spv::Op::OpSetMeshOutputsEXT);
+            patched = true;
+        } else if (emit_tasks && store_at(offset, tasks[0], tasks[1]) &&
+                   store_at(offset + 3, tasks[2], tasks[3]) && offset + 6 < code.size() &&
+                   code[offset + 6] == ((1U << 16) | static_cast<u32>(spv::Op::OpReturn))) {
+            const std::array<u32, 5> emit{(5U << 16) | static_cast<u32>(spv::Op::OpEmitMeshTasksEXT),
+                                          tasks[0].value, tasks[1].value, tasks[2].value,
+                                          tasks[3].value};
+            code.erase(code.begin() + offset, code.begin() + offset + 7);
+            code.insert(code.begin() + offset, emit.begin(), emit.end());
+            patched = true;
+            offset += emit.size();
+            continue;
+        }
+        offset += word_count;
+    }
+    if (!patched) {
+        throw LogicError("Mesh instruction markers not found");
+    }
+}
 } // Anonymous namespace
 
 std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info, IR::Program& program, Bindings& bindings) {
@@ -564,6 +623,7 @@ std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_in
     PatchPhiNodes(program, ctx);
     std::vector<u32> code{ctx.Assemble()};
     RewriteOpcodes(code, ctx.opcode_rewrites);
+    PatchMeshInstructions(code, ctx);
     return code;
 }
 
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h b/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h
index 165ec3fe84..ec526ff354 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h
@@ -61,6 +61,12 @@ Id EmitGetAttributeIndexed(EmitContext& ctx, Id offset, Id vertex);
 void EmitSetAttributeIndexed(EmitContext& ctx, Id offset, Id value, Id vertex);
 Id EmitGetPatch(EmitContext& ctx, IR::Patch patch);
 void EmitSetPatch(EmitContext& ctx, IR::Patch patch, Id value);
+Id EmitIsbeRead(EmitContext& ctx, Id address);
+void EmitIsbeWriteAttribute(EmitContext& ctx, Id address, Id value);
+void EmitIsbeWriteMapU8(EmitContext& ctx, Id address, Id value);
+void EmitIsbeWriteMapU32(EmitContext& ctx, const IR::Value& address, Id value);
+void EmitMeshPrologue(EmitContext& ctx);
+void EmitMeshEpilogue(EmitContext& ctx);
 void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value);
 void EmitSetSampleMask(EmitContext& ctx, Id value);
 void EmitSetFragDepth(EmitContext& ctx, Id value);
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_mesh.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_mesh.cpp
new file mode 100644
index 0000000000..6c599528cc
--- /dev/null
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_mesh.cpp
@@ -0,0 +1,227 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+// Switch 2 (NVN) task and mesh programs. They talk to the hardware through the stage buffer
+// (ISBE): a task program writes the number of mesh workgroups at byte 4 of its output region and
+// its payload after it; each mesh workgroup reads that region (its own index at byte 0) and
+// writes its vertices and primitives to an attribute region in groups of 32, one 32-word line per
+// attribute, plus the primitive count and 8-bit vertex indices to a map region. Here the task's
+// output region is the task payload, the mesh output regions live in workgroup memory, and the
+// epilogue (reached by every invocation, see structured_control_flow.cpp) hands them to Vulkan.
+
+#include <algorithm>
+#include <span>
+
+#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 {
+void Barrier(EmitContext& ctx) {
+    ctx.OpControlBarrier(ctx.Const(static_cast<u32>(spv::Scope::Workgroup)),
+                         ctx.Const(static_cast<u32>(spv::Scope::Workgroup)),
+                         ctx.Const(static_cast<u32>(spv::MemorySemanticsMask::AcquireRelease |
+                                                    spv::MemorySemanticsMask::WorkgroupMemory)));
+}
+
+Id LocalInvocation(EmitContext& ctx) {
+    return ctx.OpCompositeExtract(ctx.U32[1], ctx.OpLoad(ctx.U32[3], ctx.local_invocation_id), 0U);
+}
+
+Id Clamped(EmitContext& ctx, Id index, u32 count) {
+    return ctx.OpUMin(ctx.U32[1], index, ctx.Const(count - 1));
+}
+
+Id PayloadWord(EmitContext& ctx, Id index) {
+    return ctx.OpAccessChain(ctx.mesh_payload_u32, ctx.mesh_payload,
+                             Clamped(ctx, index, ctx.mesh_info.payload_words));
+}
+
+Id SharedWord(EmitContext& ctx, Id array, Id index) {
+    return ctx.OpAccessChain(ctx.mesh_shared_u32, array, index);
+}
+
+// Runs body under `if (condition)`.
+template <typename Func>
+void If(EmitContext& ctx, Id condition, Func&& body) {
+    const Id then_label{ctx.OpLabel()};
+    const Id merge_label{ctx.OpLabel()};
+    ctx.OpSelectionMerge(merge_label, spv::SelectionControlMask::MaskNone);
+    ctx.OpBranchConditional(condition, then_label, merge_label);
+    ctx.AddLabel(then_label);
+    body();
+    ctx.OpBranch(merge_label);
+    ctx.AddLabel(merge_label);
+}
+
+// Copies the lines of element `index` (a vertex or a primitive) from the attribute region.
+void CopyLines(EmitContext& ctx, Id index, u32 base, u32 num_lines,
+               const std::vector<EmitContext::MeshLineTarget>& targets) {
+    // Word of line 0: base + (index / 32) * num_lines * 32 + index % 32.
+    const Id group{ctx.OpShiftRightLogical(ctx.U32[1], index, ctx.Const(5U))};
+    const Id slot{ctx.OpBitwiseAnd(ctx.U32[1], index, ctx.Const(31U))};
+    const Id first{ctx.OpIAdd(ctx.U32[1],
+                              ctx.OpIMul(ctx.U32[1], group, ctx.Const(num_lines * 32)),
+                              ctx.OpIAdd(ctx.U32[1], slot, ctx.Const(base)))};
+    for (u32 line = 0; line < num_lines; ++line) {
+        const EmitContext::MeshLineTarget& target{targets[line]};
+        if (!Sirit::ValidId(target.variable)) {
+            continue;
+        }
+        const Id word{ctx.OpLoad(
+            ctx.U32[1],
+            SharedWord(ctx, ctx.mesh_attributes, ctx.OpIAdd(ctx.U32[1], first, ctx.Const(line * 32))))};
+        const Id value{target.is_float ? ctx.OpBitcast(ctx.F32[1], word) : word};
+        const Id pointer_type{target.is_float ? ctx.output_f32 : ctx.output_u32};
+        const Id pointer{target.element == ~0U
+                             ? ctx.OpAccessChain(pointer_type, target.variable, index)
+                             : ctx.OpAccessChain(pointer_type, target.variable, index,
+                                                 ctx.Const(target.element))};
+        ctx.OpStore(pointer, value);
+    }
+}
+
+// OpenGL-style clip space: z = (z + w) / 2, from the stage buffer's position lines.
+void ConvertDepth(EmitContext& ctx, Id vertex) {
+    const auto& lines{ctx.mesh_info.vertex_lines};
+    const auto z_line{std::ranges::find(lines, static_cast<u32>(IR::Attribute::PositionZ))};
+    const auto w_line{std::ranges::find(lines, static_cast<u32>(IR::Attribute::PositionW))};
+    if (z_line == lines.end() || w_line == lines.end()) {
+        return;
+    }
+    const u32 num_lines{static_cast<u32>(lines.size())};
+    const Id group{ctx.OpShiftRightLogical(ctx.U32[1], vertex, ctx.Const(5U))};
+    const Id slot{ctx.OpBitwiseAnd(ctx.U32[1], vertex, ctx.Const(31U))};
+    const Id first{ctx.OpIAdd(ctx.U32[1], ctx.OpIMul(ctx.U32[1], group, ctx.Const(num_lines * 32)),
+                              slot)};
+    const auto read = [&](auto line) {
+        const u32 offset{static_cast<u32>(line - lines.begin()) * 32};
+        return ctx.OpBitcast(ctx.F32[1],
+                             ctx.OpLoad(ctx.U32[1], SharedWord(ctx, ctx.mesh_attributes,
+                                                               ctx.OpIAdd(ctx.U32[1], first,
+                                                                          ctx.Const(offset)))));
+    };
+    const Id z{ctx.OpFMul(ctx.F32[1], ctx.OpFAdd(ctx.F32[1], read(z_line), read(w_line)),
+                          ctx.Const(0.5f))};
+    ctx.OpStore(ctx.OpAccessChain(ctx.output_f32, ctx.output_position, vertex, ctx.Const(2U)), z);
+}
+
+void TaskEpilogue(EmitContext& ctx) {
+    Barrier(ctx);
+    const Id count{ctx.OpLoad(ctx.U32[1], PayloadWord(ctx, ctx.Const(1U)))};
+    const Id one{ctx.Const(1U)};
+    // Marker for OpEmitMeshTasksEXT count, 1, 1, payload (PatchMeshInstructions).
+    ctx.OpStore(count, one);
+    ctx.OpStore(one, ctx.mesh_payload);
+    ctx.mesh_tasks_marker = {count, one, one, ctx.mesh_payload};
+}
+
+void MeshEpilogue(EmitContext& ctx) {
+    const MeshInfo& mesh{ctx.mesh_info};
+    Barrier(ctx);
+    const Id count{ctx.OpUMin(ctx.U32[1], ctx.OpLoad(ctx.U32[1], ctx.mesh_primitive_count),
+                              ctx.Const(mesh.max_primitives))};
+    const Id max_vertices{ctx.Const(mesh.max_vertices)};
+    // Marker for OpSetMeshOutputsEXT max_vertices, count (PatchMeshInstructions).
+    ctx.OpStore(max_vertices, count);
+    ctx.mesh_outputs_marker = {max_vertices, count};
+
+    const Id invocation{LocalInvocation(ctx)};
+    const u32 threads{ctx.mesh_threads};
+    const bool convert_depth{ctx.runtime_info.convert_depth_mode && !ctx.profile.support_native_ndc &&
+                             Sirit::ValidId(ctx.output_position)};
+    const u32 vertex_lines{static_cast<u32>(mesh.vertex_lines.size())};
+    for (u32 first = 0; first < mesh.max_vertices; first += threads) {
+        const Id vertex{ctx.OpIAdd(ctx.U32[1], invocation, ctx.Const(first))};
+        const auto copy = [&] {
+            CopyLines(ctx, vertex, 0, vertex_lines, ctx.mesh_vertex_targets);
+            if (convert_depth) {
+                ConvertDepth(ctx, vertex);
+            }
+        };
+        if (first + threads > mesh.max_vertices) {
+            If(ctx, ctx.OpULessThan(ctx.U1, vertex, max_vertices), copy);
+        } else {
+            copy();
+        }
+    }
+    const u32 primitive_lines{static_cast<u32>(mesh.primitive_lines.size())};
+    for (u32 first = 0; first < mesh.max_primitives; first += threads) {
+        const Id primitive{ctx.OpIAdd(ctx.U32[1], invocation, ctx.Const(first))};
+        If(ctx, ctx.OpULessThan(ctx.U1, primitive, count), [&] {
+            // Vertex indices: map bytes 4 + 3 * primitive.
+            const Id byte{ctx.OpIAdd(ctx.U32[1], ctx.OpIMul(ctx.U32[1], primitive, ctx.Const(3U)),
+                                     ctx.Const(4U))};
+            std::array<Id, 3> indices;
+            for (u32 element = 0; element < 3; ++element) {
+                const Id address{ctx.OpIAdd(ctx.U32[1], byte, ctx.Const(element))};
+                indices[element] = ctx.OpLoad(ctx.U32[1], SharedWord(ctx, ctx.mesh_map, address));
+            }
+            const Id output_indices{ctx.TypePointer(spv::StorageClass::Output, ctx.U32[3])};
+            ctx.OpStore(ctx.OpAccessChain(output_indices, ctx.mesh_primitive_indices, primitive),
+                        ctx.OpCompositeConstruct(ctx.U32[3], std::span<const Id>(indices)));
+            CopyLines(ctx, primitive, ctx.mesh_primitive_base, primitive_lines,
+                      ctx.mesh_primitive_targets);
+        });
+    }
+}
+} // Anonymous namespace
+
+Id EmitIsbeRead(EmitContext& ctx, Id address) {
+    const Id workgroup{
+        ctx.OpCompositeExtract(ctx.U32[1], ctx.OpLoad(ctx.U32[3], ctx.workgroup_id), 0U)};
+    if (ctx.stage == Stage::Task) {
+        return workgroup; // the input region holds the workgroup index (VERTEX_ID)
+    }
+    const Id index{ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.Const(2U))};
+    const Id word{ctx.OpLoad(ctx.U32[1], PayloadWord(ctx, index))};
+    return ctx.OpSelect(ctx.U32[1], ctx.OpIEqual(ctx.U1, address, ctx.u32_zero_value), workgroup,
+                        word);
+}
+
+void EmitIsbeWriteAttribute(EmitContext& ctx, Id address, Id value) {
+    const Id index{ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.Const(2U))};
+    if (ctx.stage == Stage::Task) {
+        ctx.OpStore(PayloadWord(ctx, index), value);
+        return;
+    }
+    ctx.OpStore(SharedWord(ctx, ctx.mesh_attributes, Clamped(ctx, index, ctx.mesh_attribute_words)),
+                value);
+}
+
+void EmitIsbeWriteMapU8(EmitContext& ctx, Id address, Id value) {
+    if (ctx.stage != Stage::Mesh) {
+        throw NotImplementedException("Stage buffer map write outside mesh programs");
+    }
+    ctx.OpStore(SharedWord(ctx, ctx.mesh_map, Clamped(ctx, address, ctx.mesh_map_bytes)),
+                ctx.OpBitwiseAnd(ctx.U32[1], value, ctx.Const(0xffU)));
+}
+
+void EmitIsbeWriteMapU32(EmitContext& ctx, const IR::Value& address, Id value) {
+    // The primitive count: the map is written from its end, so its first word is at byte 3.
+    if (ctx.stage != Stage::Mesh || !address.IsImmediate() || address.U32() != 3) {
+        throw NotImplementedException("32-bit stage buffer map write");
+    }
+    ctx.OpStore(ctx.mesh_primitive_count, value);
+}
+
+void EmitMeshPrologue(EmitContext& ctx) {
+    if (ctx.stage != Stage::Mesh) {
+        return;
+    }
+    // Every invocation stores the zero: no branch here, the entry block's label is a phi
+    // predecessor for later blocks.
+    ctx.OpStore(ctx.mesh_primitive_count, ctx.u32_zero_value);
+    Barrier(ctx);
+}
+
+void EmitMeshEpilogue(EmitContext& ctx) {
+    if (ctx.stage == Stage::Task) {
+        TaskEpilogue(ctx);
+    } else if (ctx.stage == Stage::Mesh) {
+        MeshEpilogue(ctx);
+    }
+}
+
+} // namespace Shader::Backend::SPIRV
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp
index 8eca1fac29..2b36c1fd8d 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp
@@ -137,9 +137,16 @@ void EmitPrologue(EmitContext& ctx) {
     if (ctx.stage == Stage::VertexB || ctx.stage == Stage::Geometry) {
         SetFixedPipelinePointSize(ctx);
     }
+    if (ctx.stage == Stage::Task || ctx.stage == Stage::Mesh) {
+        EmitMeshPrologue(ctx);
+    }
 }
 
 void EmitEpilogue(EmitContext& ctx) {
+    if (ctx.stage == Stage::Task || ctx.stage == Stage::Mesh) {
+        EmitMeshEpilogue(ctx);
+        return;
+    }
     if (ctx.stage == Stage::VertexB && ctx.runtime_info.convert_depth_mode &&
         !ctx.profile.support_native_ndc) {
         ConvertDepthMode(ctx);
diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
index b562e82983..f80f0440df 100644
--- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
+++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
@@ -257,6 +257,10 @@ std::string_view StageName(Stage stage) {
         return "fs";
     case Stage::Compute:
         return "cs";
+    case Stage::Task:
+        return "ts";
+    case Stage::Mesh:
+        return "ms";
     }
     throw InvalidArgument("Invalid stage {}", stage);
 }
@@ -475,6 +479,7 @@ EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_inf
     DefineCommonTypes(program.info);
     DefineCommonConstants();
     DefineInterfaces(program);
+    DefineMesh(program);
     DefineLocalMemory(program);
     DefineSharedMemory(program);
     DefineSharedMemoryFunctions(program);
@@ -1586,6 +1591,16 @@ void EmitContext::DefineInputs(const IR::Program& program) {
         }
         const bool is_integer = input_type == AttributeType::SignedInt ||
                                 input_type == AttributeType::UnsignedInt;
+        if (runtime_info.per_primitive_generics.test(index)) {
+            // Written per primitive by a mesh program: constant over the primitive.
+            AddExtension("SPV_EXT_mesh_shader");
+            AddCapability(spv::Capability::MeshShadingEXT);
+            Decorate(id, spv::Decoration::PerPrimitiveEXT);
+            if (is_integer) {
+                Decorate(id, spv::Decoration::Flat);
+            }
+            continue;
+        }
         if (is_integer) {
             Decorate(id, spv::Decoration::Flat);
         } else {
@@ -1614,8 +1629,120 @@ void EmitContext::DefineInputs(const IR::Program& program) {
     }
 }
 
+void EmitContext::DefineMesh(const IR::Program& program) {
+    if (stage != Stage::Task && stage != Stage::Mesh) {
+        return;
+    }
+    AddExtension("SPV_EXT_mesh_shader");
+    AddCapability(spv::Capability::MeshShadingEXT);
+    mesh_info = program.mesh;
+    mesh_threads = program.workgroup_size[0];
+    // The task's output region, which its mesh workgroups read as their input region.
+    const Id payload_type{TypeArray(U32[1], Const(mesh_info.payload_words))};
+    mesh_payload = DefineVariable(*this, payload_type, std::nullopt,
+                                  spv::StorageClass::TaskPayloadWorkgroupEXT);
+    Name(mesh_payload, "task_payload");
+    mesh_payload_u32 = TypePointer(spv::StorageClass::TaskPayloadWorkgroupEXT, U32[1]);
+    if (stage == Stage::Task) {
+        return;
+    }
+    // Output attribute region: groups of 32 vertices, then groups of 32 primitives, one 32-word
+    // line per attribute. The map region keeps one map byte per word.
+    const u32 vertex_lines{static_cast<u32>(mesh_info.vertex_lines.size())};
+    const u32 primitive_lines{static_cast<u32>(mesh_info.primitive_lines.size())};
+    mesh_primitive_base = Common::DivCeil(mesh_info.max_vertices, 32U) * vertex_lines * 32;
+    mesh_attribute_words = (std::max)(
+        mesh_primitive_base + Common::DivCeil(mesh_info.max_primitives, 32U) * primitive_lines * 32,
+        1U);
+    mesh_map_bytes = 4 + 3 * mesh_info.max_primitives;
+    mesh_shared_u32 = TypePointer(spv::StorageClass::Workgroup, U32[1]);
+    mesh_attributes = DefineVariable(*this, TypeArray(U32[1], Const(mesh_attribute_words)),
+                                     std::nullopt, spv::StorageClass::Workgroup);
+    Name(mesh_attributes, "isbe_attributes");
+    mesh_map = DefineVariable(*this, TypeArray(U32[1], Const(mesh_map_bytes)), std::nullopt,
+                              spv::StorageClass::Workgroup);
+    Name(mesh_map, "isbe_map");
+    mesh_primitive_count =
+        DefineVariable(*this, U32[1], std::nullopt, spv::StorageClass::Workgroup);
+    Name(mesh_primitive_count, "isbe_primitive_count");
+
+    const Id vertex_count{Const(mesh_info.max_vertices)};
+    const Id primitive_count{Const(mesh_info.max_primitives)};
+    mesh_primitive_indices =
+        DefineVariable(*this, TypeArray(U32[3], primitive_count),
+                       spv::BuiltIn::PrimitiveTriangleIndicesEXT, spv::StorageClass::Output);
+    std::array<Id, IR::NUM_GENERICS> generics[2]{};
+    Id point_size{}, primitive_layer{}, primitive_viewport{}, primitive_id_output{};
+    const auto per_primitive_builtin = [&](Id& variable, spv::BuiltIn builtin) {
+        if (!Sirit::ValidId(variable)) {
+            variable = DefineVariable(*this, TypeArray(U32[1], primitive_count), builtin,
+                                      spv::StorageClass::Output);
+            Decorate(variable, spv::Decoration::PerPrimitiveEXT);
+        }
+        return MeshLineTarget{variable, ~0U, false};
+    };
+    const auto target = [&](u32 line, bool per_primitive) -> MeshLineTarget {
+        const IR::Attribute attribute{line};
+        if (IR::IsGeneric(attribute)) {
+            const u32 index{IR::GenericAttributeIndex(attribute)};
+            Id& variable{generics[per_primitive][index]};
+            if (!Sirit::ValidId(variable)) {
+                variable = DefineVariable(
+                    *this, TypeArray(F32[4], per_primitive ? primitive_count : vertex_count),
+                    std::nullopt, spv::StorageClass::Output);
+                Decorate(variable, spv::Decoration::Location, index);
+                if (per_primitive) {
+                    Decorate(variable, spv::Decoration::PerPrimitiveEXT);
+                }
+            }
+            return {variable, IR::GenericAttributeElement(attribute), true};
+        }
+        if (per_primitive) {
+            switch (attribute) {
+            case IR::Attribute::PrimitiveId:
+                return per_primitive_builtin(primitive_id_output, spv::BuiltIn::PrimitiveId);
+            case IR::Attribute::Layer:
+                return per_primitive_builtin(primitive_layer, spv::BuiltIn::Layer);
+            case IR::Attribute::ViewportIndex:
+                return per_primitive_builtin(primitive_viewport, spv::BuiltIn::ViewportIndex);
+            default:
+                return {};
+            }
+        }
+        switch (attribute) {
+        case IR::Attribute::PositionX:
+        case IR::Attribute::PositionY:
+        case IR::Attribute::PositionZ:
+        case IR::Attribute::PositionW:
+            if (!Sirit::ValidId(output_position)) {
+                output_position = DefineVariable(*this, TypeArray(F32[4], vertex_count),
+                                                 spv::BuiltIn::Position, spv::StorageClass::Output);
+            }
+            return {output_position,
+                    static_cast<u32>(attribute) - static_cast<u32>(IR::Attribute::PositionX), true};
+        case IR::Attribute::PointSize:
+            if (!Sirit::ValidId(point_size)) {
+                point_size = DefineVariable(*this, TypeArray(F32[1], vertex_count),
+                                            spv::BuiltIn::PointSize, spv::StorageClass::Output);
+            }
+            return {point_size, ~0U, true};
+        default:
+            return {}; // per-primitive values in vertex lines have no vertex output
+        }
+    };
+    for (const u32 line : mesh_info.vertex_lines) {
+        mesh_vertex_targets.push_back(target(line, false));
+    }
+    for (const u32 line : mesh_info.primitive_lines) {
+        mesh_primitive_targets.push_back(target(line, true));
+    }
+}
+
 void EmitContext::DefineOutputs(const IR::Program& program) {
     const Info& info{program.info};
+    if (stage == Stage::Task || stage == Stage::Mesh) {
+        return; // DefineMesh
+    }
     const std::optional<u32> invocations{program.invocations};
     if (runtime_info.convert_depth_mode || info.stores.AnyComponent(IR::Attribute::PositionX) ||
         stage == Stage::VertexB) {
diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.h b/src/shader_recompiler/backend/spirv/spirv_emit_context.h
index a5821e9e93..1f9b534aaf 100644
--- a/src/shader_recompiler/backend/spirv/spirv_emit_context.h
+++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h
@@ -363,6 +363,32 @@ public:
     std::vector<Id> interfaces;
     std::vector<std::pair<u32, spv::Op>> opcode_rewrites;
 
+    // Task and mesh programs. Their stage buffer: the task's output region is the task payload
+    // (the mesh stage's input region); the mesh output regions live in workgroup memory until
+    // the epilogue copies them to the outputs. The two mesh instructions Sirit lacks are emitted
+    // as OpStore markers with these operands and replaced after assembly (PatchMeshInstructions).
+    struct MeshLineTarget {
+        Id variable;    ///< arrayed output
+        u32 element;    ///< vector element, or ~0 for a scalar output
+        bool is_float;  ///< bitcast the stage buffer word to F32
+    };
+    Id mesh_payload{};
+    Id mesh_payload_u32{};  ///< pointer type to one payload word
+    Id mesh_attributes{};
+    Id mesh_map{};
+    Id mesh_primitive_count{};
+    Id mesh_shared_u32{};   ///< pointer type to one workgroup word
+    Id mesh_primitive_indices{};
+    MeshInfo mesh_info{};
+    u32 mesh_threads{};
+    u32 mesh_attribute_words{};
+    u32 mesh_map_bytes{};
+    u32 mesh_primitive_base{}; ///< first word of the per-primitive groups
+    std::vector<MeshLineTarget> mesh_vertex_targets;
+    std::vector<MeshLineTarget> mesh_primitive_targets;
+    std::array<Id, 2> mesh_outputs_marker{};
+    std::array<Id, 4> mesh_tasks_marker{};
+
     Id load_const_func_u8{};
     Id load_const_func_u16{};
     Id load_const_func_u32{};
@@ -402,6 +428,7 @@ private:
 
     void DefineInputs(const IR::Program& program);
     void DefineOutputs(const IR::Program& program);
+    void DefineMesh(const IR::Program& program);
 };
 
 } // namespace Shader::Backend::SPIRV
diff --git a/src/shader_recompiler/environment.h b/src/shader_recompiler/environment.h
index 67486430c8..d186b0c1d3 100644
--- a/src/shader_recompiler/environment.h
+++ b/src/shader_recompiler/environment.h
@@ -121,9 +121,22 @@ public:
         return is_proprietary_driver;
     }
 
+    /// Task and mesh programs: SET_MESH_SHADER_A, SET_MESH_SHADER_B and SET_MESH_INIT_SHADER.
+    [[nodiscard]] const std::array<u32, 3>& MeshRegisters() const noexcept {
+        return mesh_registers;
+    }
+
+    /// Mesh programs: the header of the pass-through geometry program bound with them, whose
+    /// output map lists the per-primitive outputs.
+    [[nodiscard]] const ProgramHeader& MeshPrimitiveHeader() const noexcept {
+        return mesh_primitive_header;
+    }
+
 protected:
     ProgramHeader sph{};
     std::array<u32, 8> gp_passthrough_mask{};
+    std::array<u32, 3> mesh_registers{};
+    ProgramHeader mesh_primitive_header{};
     Stage stage{};
     u32 start_address{};
     bool is_proprietary_driver{};
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.cpp b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
index 95815d7d51..23d1db5bf4 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.cpp
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
@@ -329,6 +329,22 @@ void IREmitter::SetPatch(Patch patch, const F32& value) {
     Inst(Opcode::SetPatch, patch, value);
 }
 
+U32 IREmitter::IsbeRead(const U32& address) {
+    return Inst<U32>(Opcode::IsbeRead, address);
+}
+
+void IREmitter::IsbeWriteAttribute(const U32& address, const U32& value) {
+    Inst(Opcode::IsbeWriteAttribute, address, value);
+}
+
+void IREmitter::IsbeWriteMapU8(const U32& address, const U32& value) {
+    Inst(Opcode::IsbeWriteMapU8, address, value);
+}
+
+void IREmitter::IsbeWriteMapU32(const U32& address, const U32& value) {
+    Inst(Opcode::IsbeWriteMapU32, address, value);
+}
+
 void IREmitter::SetFragColor(u32 index, u32 component, const F32& value) {
     Inst(Opcode::SetFragColor, Imm32(index), Imm32(component), value);
 }
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.h b/src/shader_recompiler/frontend/ir/ir_emitter.h
index f22f0f75e2..d0cd88ca3c 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.h
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.h
@@ -88,6 +88,11 @@ public:
     [[nodiscard]] F32 GetPatch(Patch patch);
     void SetPatch(Patch patch, const F32& value);
 
+    [[nodiscard]] U32 IsbeRead(const U32& address);
+    void IsbeWriteAttribute(const U32& address, const U32& value);
+    void IsbeWriteMapU8(const U32& address, const U32& value);
+    void IsbeWriteMapU32(const U32& address, const U32& value);
+
     void SetFragColor(u32 index, u32 component, const F32& value);
     void SetSampleMask(const U32& value);
     void SetFragDepth(const F32& value);
diff --git a/src/shader_recompiler/frontend/ir/microinstruction.cpp b/src/shader_recompiler/frontend/ir/microinstruction.cpp
index 84417980bc..a579ed1057 100644
--- a/src/shader_recompiler/frontend/ir/microinstruction.cpp
+++ b/src/shader_recompiler/frontend/ir/microinstruction.cpp
@@ -82,6 +82,9 @@ bool Inst::MayHaveSideEffects() const noexcept {
     case Opcode::SetAttribute:
     case Opcode::SetAttributeIndexed:
     case Opcode::SetPatch:
+    case Opcode::IsbeWriteAttribute:
+    case Opcode::IsbeWriteMapU8:
+    case Opcode::IsbeWriteMapU32:
     case Opcode::SetFragColor:
     case Opcode::SetSampleMask:
     case Opcode::SetFragDepth:
diff --git a/src/shader_recompiler/frontend/ir/opcodes.cpp b/src/shader_recompiler/frontend/ir/opcodes.cpp
index 94ba0f6ae0..1999f877fe 100644
--- a/src/shader_recompiler/frontend/ir/opcodes.cpp
+++ b/src/shader_recompiler/frontend/ir/opcodes.cpp
@@ -10,7 +10,7 @@ namespace Shader::IR {
 
 namespace Detail {
 
-OpcodeMeta META_TABLE[534] = {
+OpcodeMeta META_TABLE[538] = {
 #define OPCODE(name_token, type_token, ...)                                                        \
     {                                                                                              \
         .name{#name_token},                                                                        \
@@ -21,7 +21,7 @@ OpcodeMeta META_TABLE[534] = {
 #undef OPCODE
 };
 
-u8 NUM_ARGS[534] = {
+u8 NUM_ARGS[538] = {
 #define OPCODE(name_token, type_token, ...) u8(CalculateNumArgsOf(Opcode::name_token)),
 #include "opcodes.inc"
 #undef OPCODE
diff --git a/src/shader_recompiler/frontend/ir/opcodes.h b/src/shader_recompiler/frontend/ir/opcodes.h
index 50f9d7b155..c59f50adfe 100644
--- a/src/shader_recompiler/frontend/ir/opcodes.h
+++ b/src/shader_recompiler/frontend/ir/opcodes.h
@@ -57,12 +57,12 @@ static constexpr Type F64x2{Type::F64x2};
 static constexpr Type F64x3{Type::F64x3};
 static constexpr Type F64x4{Type::F64x4};
 
-extern OpcodeMeta META_TABLE[534];
+extern OpcodeMeta META_TABLE[538];
 constexpr size_t CalculateNumArgsOf(Opcode op) noexcept {
     const auto& arg_types = META_TABLE[size_t(op)].arg_types;
     return size_t(std::distance(arg_types.begin(), std::ranges::find(arg_types, Type::Void)));
 }
-extern u8 NUM_ARGS[534];
+extern u8 NUM_ARGS[538];
 } // namespace Detail
 
 /// Get return type of an opcode
diff --git a/src/shader_recompiler/frontend/ir/opcodes.inc b/src/shader_recompiler/frontend/ir/opcodes.inc
index fee0afaf33..552ec4fbb6 100644
--- a/src/shader_recompiler/frontend/ir/opcodes.inc
+++ b/src/shader_recompiler/frontend/ir/opcodes.inc
@@ -48,6 +48,10 @@ OPCODE(GetAttributeIndexed,                                 F32,            U32,
 OPCODE(SetAttributeIndexed,                                 Void,           U32,            F32,            U32,                                            )
 OPCODE(GetPatch,                                            F32,            Patch,                                                                          )
 OPCODE(SetPatch,                                            Void,           Patch,          F32,                                                            )
+OPCODE(IsbeRead,                                            U32,            U32,                                                                            )
+OPCODE(IsbeWriteAttribute,                                  Void,           U32,            U32,                                                            )
+OPCODE(IsbeWriteMapU8,                                      Void,           U32,            U32,                                                            )
+OPCODE(IsbeWriteMapU32,                                     Void,           U32,            U32,                                                            )
 OPCODE(SetFragColor,                                        Void,           U32,            U32,            F32,                                            )
 OPCODE(SetSampleMask,                                       Void,           U32,                                                                            )
 OPCODE(SetFragDepth,                                        Void,           F32,                                                                            )
diff --git a/src/shader_recompiler/frontend/ir/program.h b/src/shader_recompiler/frontend/ir/program.h
index 6b4a05c598..fee5cee721 100644
--- a/src/shader_recompiler/frontend/ir/program.h
+++ b/src/shader_recompiler/frontend/ir/program.h
@@ -27,6 +27,7 @@ struct Program {
     u32 local_memory_size{};
     u32 shared_memory_size{};
     bool is_geometry_passthrough{};
+    MeshInfo mesh{};
 };
 
 [[nodiscard]] std::string DumpProgram(const Program& program);
diff --git a/src/shader_recompiler/frontend/maxwell/maxwell.inc b/src/shader_recompiler/frontend/maxwell/maxwell.inc
index e5dc9c8005..5facfb4803 100644
--- a/src/shader_recompiler/frontend/maxwell/maxwell.inc
+++ b/src/shader_recompiler/frontend/maxwell/maxwell.inc
@@ -55,6 +55,7 @@ INST(IMNMX_reg,    "IMNMX (reg)",    "0101 1100 0010 0---")
 INST(IMUL_cbuf,    "IMUL (cbuf)",    "0100 1100 0011 1---")
 INST(IMUL_reg,     "IMUL (reg)",     "0101 1100 0011 1---")
 INST(ISBERD,       "ISBERD",         "1110 1111 1101 0---")
+INST(ISBE_nx2,     "ISBE (nx2)",     "1110 1111 1100 0---")
 INST(ISCADD_cbuf,  "ISCADD (cbuf)",  "0100 1100 0001 1---")
 INST(ISCADD_reg,   "ISCADD (reg)",   "0101 1100 0001 1---")
 INST(LDC,          "LDC",            "1110 1111 1001 0---")
diff --git a/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp b/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
index 329d44443e..50cff52f88 100644
--- a/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
+++ b/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
@@ -314,7 +314,8 @@ bool NeedsLift(Node goto_stmt, Node label_stmt) noexcept {
 
 class GotoPass {
 public:
-    explicit GotoPass(Flow::CFG& cfg, ObjectPool<Statement>& stmt_pool) : pool{stmt_pool} {
+    explicit GotoPass(Flow::CFG& cfg, ObjectPool<Statement>& stmt_pool, bool converge_exits_)
+        : pool{stmt_pool}, converge_exits{converge_exits_} {
         std::vector gotos{BuildTree(cfg)};
         const auto end{gotos.rend()};
         for (auto goto_stmt = gotos.rbegin(); goto_stmt != end; ++goto_stmt) {
@@ -381,7 +382,17 @@ private:
         u32 label_id{0};
         std::vector<Node> gotos;
         Flow::Function& first_function{cfg.Functions().front()};
-        BuildTree(cfg, first_function, label_id, gotos, root_stmt.children.end(), std::nullopt);
+        Node insert_point{root_stmt.children.end()};
+        if (converge_exits) {
+            // Every EXIT jumps to one return at the end, so all invocations reach the same
+            // epilogue (task and mesh programs emit their outputs there, after a barrier).
+            Tree& root{root_stmt.children};
+            exit_label = root.insert(root.end(), *pool.Create(Label{}, label_id, &root_stmt));
+            ++label_id;
+            root.insert(root.end(), *pool.Create(Return{}, &root_stmt));
+            insert_point = *exit_label;
+        }
+        BuildTree(cfg, first_function, label_id, gotos, insert_point, std::nullopt);
         return gotos;
     }
 
@@ -454,7 +465,14 @@ private:
                 break;
             }
             case Flow::EndClass::Exit:
-                root.insert(ip, *pool.Create(Return{}, &root_stmt));
+                if (exit_label) {
+                    Statement* const always_cond{
+                        pool.Create(Identity{}, IR::Condition{true}, &root_stmt)};
+                    gotos.push_back(root.insert(
+                        ip, *pool.Create(Goto{}, always_cond, *exit_label, &root_stmt)));
+                } else {
+                    root.insert(ip, *pool.Create(Return{}, &root_stmt));
+                }
                 break;
             case Flow::EndClass::Return: {
                 Statement* const always_cond{pool.Create(Identity{}, block.cond, &root_stmt)};
@@ -618,6 +636,8 @@ private:
 
     ObjectPool<Statement>& pool;
     Statement root_stmt{FunctionTag{}};
+    bool converge_exits{};
+    std::optional<Node> exit_label;
 };
 
 [[nodiscard]] Statement* TryFindForwardBlock(Statement& stmt) {
@@ -987,7 +1007,8 @@ IR::AbstractSyntaxList BuildASL(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::
                                 Environment& env, Flow::CFG& cfg,
                                 const HostTranslateInfo& host_info) {
     ObjectPool<Statement> stmt_pool{64};
-    GotoPass goto_pass{cfg, stmt_pool};
+    const Stage stage{env.ShaderStage()};
+    GotoPass goto_pass{cfg, stmt_pool, stage == Stage::Task || stage == Stage::Mesh};
     Statement& root{goto_pass.RootStatement()};
     IR::AbstractSyntaxList syntax_list;
     TranslatePass{inst_pool, block_pool, stmt_pool, env, root, syntax_list, host_info};
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/internal_stage_buffer_entry_read.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/internal_stage_buffer_entry_read.cpp
index 46660563a1..27de2a0171 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/internal_stage_buffer_entry_read.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/internal_stage_buffer_entry_read.cpp
@@ -125,4 +125,40 @@ void TranslatorVisitor::ISBERD(u64 insn) {
     X(isberd.dest_reg.Value(), X(isberd.src_reg.Value()));
 }
 
+// nx2 pseudo instruction for the ISBERD/ISBEWR of Switch 2 task and mesh programs (see
+// frontend/ampere/lower_to_maxwell.h). Addresses are linear bytes in the stage buffer: reads come
+// from the input attribute region, writes go to the output attribute or map region.
+void TranslatorVisitor::ISBE_nx2(u64 insn) {
+    union {
+        u64 raw;
+        BitField<0, 8, IR::Reg> reg;
+        BitField<8, 8, IR::Reg> address_reg;
+        BitField<20, 24, u64> offset;
+        BitField<44, 1, u64> is_32;
+        BitField<45, 1, u64> attribute;
+        BitField<48, 1, u64> write;
+    } const isbe{insn};
+
+    const Stage stage{env.ShaderStage()};
+    if (stage != Stage::Task && stage != Stage::Mesh) {
+        throw NotImplementedException("Stage buffer access outside task and mesh programs");
+    }
+    const IR::U32 address{ir.IAdd(X(isbe.address_reg), ir.Imm32(static_cast<u32>(isbe.offset)))};
+    if (isbe.write == 0) {
+        X(isbe.reg, ir.IsbeRead(address));
+        return;
+    }
+    const IR::U32 value{X(isbe.reg)};
+    if (isbe.attribute != 0) {
+        if (isbe.is_32 == 0) {
+            throw NotImplementedException("U8 stage buffer attribute write");
+        }
+        ir.IsbeWriteAttribute(address, value);
+    } else if (isbe.is_32 != 0) {
+        ir.IsbeWriteMapU32(address, value);
+    } else {
+        ir.IsbeWriteMapU8(address, value);
+    }
+}
+
 } // namespace Shader::Maxwell
diff --git a/src/shader_recompiler/frontend/maxwell/translate_program.cpp b/src/shader_recompiler/frontend/maxwell/translate_program.cpp
index cd4678cef6..1fcb533ae7 100644
--- a/src/shader_recompiler/frontend/maxwell/translate_program.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate_program.cpp
@@ -101,10 +101,12 @@ void AddNVNStorageBuffers(IR::Program& program) {
         switch (program.stage) {
         case Stage::VertexA:
         case Stage::VertexB:
+        case Stage::Task: // bound in the vertex slot
             return 0x110u;
         case Stage::TessellationControl:
             return 0x210u;
         case Stage::TessellationEval:
+        case Stage::Mesh: // bound in the tessellation slot
             return 0x310u;
         case Stage::Geometry:
             return 0x410u;
@@ -304,6 +306,30 @@ void LowerGeometryPassthrough(const IR::Program& program, const HostTranslateInf
     }
 }
 
+// The lines of a 32-element group in a mesh program's stage buffer: one per attribute an output
+// map enables, in attribute order (system values, then generics).
+void AddMeshLines(const ProgramHeader& sph, boost::container::static_vector<u32, 136>& lines) {
+    const auto& c{sph.vtg.omap_systemc};
+    if (c.clip_distances != 0 || c.point_sprite_s != 0 || c.point_sprite_t != 0 ||
+        c.fog_coordinate != 0 || c.tessellation_eval_point_u != 0 ||
+        c.tessellation_eval_point_v != 0 || c.instance_id != 0 || c.vertex_id != 0) {
+        throw NotImplementedException("Mesh outputs past the generic attributes");
+    }
+    for (u32 bit = 0; bit < 8; ++bit) {
+        if (((sph.vtg.omap_systemb.raw >> bit) & 1) != 0) {
+            lines.push_back(static_cast<u32>(IR::Attribute::PrimitiveId) + bit);
+        }
+    }
+    for (u32 generic = 0; generic < 32; ++generic) {
+        const std::array<bool, 4> mask{sph.vtg.OutputGeneric(generic)};
+        for (u32 element = 0; element < 4; ++element) {
+            if (mask[element]) {
+                lines.push_back(static_cast<u32>(IR::Attribute::Generic0X) + generic * 4 + element);
+            }
+        }
+    }
+}
+
 void TightenOutputVertices(IR::Program& program) {
     if (program.stage != Stage::Geometry || program.is_geometry_passthrough) {
         return;
@@ -368,6 +394,29 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
         program.workgroup_size = env.WorkgroupSize();
         program.shared_memory_size = env.SharedMemorySize();
         break;
+    case Stage::Task:
+    case Stage::Mesh: {
+        // SET_MESH_SHADER_A: topology 3:0, max vertex 12:4, max primitive 22:13;
+        // SET_MESH_SHADER_B: mesh threads 25:20; SET_MESH_INIT_SHADER: task threads 5:0, local
+        // buffer lines 20:12, lines passed to the mesh stage 30:22 (lines are 128 bytes).
+        const auto& registers{env.MeshRegisters()};
+        const u32 shader_a{registers[0]}, shader_b{registers[1]}, init{registers[2]};
+        const u32 lines{(std::max)((init >> 12) & 0x1ff, (init >> 22) & 0x1ff)};
+        program.mesh.payload_words = (std::max)(lines, 1U) * 32;
+        program.mesh.topology = shader_a & 0xf;
+        program.mesh.max_vertices = (shader_a >> 4) & 0x1ff;
+        program.mesh.max_primitives = (shader_a >> 13) & 0x3ff;
+        const u32 threads{program.stage == Stage::Task ? init & 0x3f : (shader_b >> 20) & 0x3f};
+        program.workgroup_size = {(std::max)(threads, 1U), 1, 1};
+        if (program.stage == Stage::Mesh) {
+            if (program.mesh.topology != 4) {
+                throw NotImplementedException("Mesh output topology {}", program.mesh.topology);
+            }
+            AddMeshLines(env.SPH(), program.mesh.vertex_lines);
+            AddMeshLines(env.MeshPrimitiveHeader(), program.mesh.primitive_lines);
+        }
+        break;
+    }
     default:
         break;
     }
@@ -406,6 +455,18 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
         Optimization::VerificationPass(program);
     }
     Optimization::CollectShaderInfoPass(env, program);
+    program.info.is_task = program.stage == Stage::Task;
+    program.info.is_mesh = program.stage == Stage::Mesh;
+    if (program.stage == Stage::Mesh) {
+        // Its outputs come from the stage buffer lines, not attribute stores; the next stage
+        // reads what they name.
+        for (const u32 line : program.mesh.vertex_lines) {
+            program.info.stores.Set(static_cast<IR::Attribute>(line), true);
+        }
+        for (const u32 line : program.mesh.primitive_lines) {
+            program.info.stores.Set(static_cast<IR::Attribute>(line), true);
+        }
+    }
     Optimization::LayerPass(program, normalized_host_info);
     Optimization::VendorWorkaroundPass(program);
 
diff --git a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
index 9010c2bf68..c212a25d37 100644
--- a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
+++ b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
@@ -963,10 +963,12 @@ void CollectShaderInfoPass(Environment& env, IR::Program& program) {
         switch (program.stage) {
         case Stage::VertexA:
         case Stage::VertexB:
+        case Stage::Task: // bound in the vertex slot
             return 0x110u;
         case Stage::TessellationControl:
             return 0x210u;
         case Stage::TessellationEval:
+        case Stage::Mesh: // bound in the tessellation slot
             return 0x310u;
         case Stage::Geometry:
             return 0x410u;
@@ -978,6 +980,11 @@ void CollectShaderInfoPass(Environment& env, IR::Program& program) {
         throw InvalidArgument("Invalid stage {}", program.stage);
     }()};
     info.nvn_buffer_base = base;
+    if (program.stage == Stage::Task || program.stage == Stage::Mesh) {
+        // The stage buffer reads the workgroup index; the epilogue splits work by invocation.
+        info.uses_workgroup_id = true;
+        info.uses_local_invocation_id = true;
+    }
 
     for (IR::Block* const block : program.post_order_blocks) {
         for (IR::Inst& inst : block->Instructions()) {
diff --git a/src/shader_recompiler/runtime_info.h b/src/shader_recompiler/runtime_info.h
index e6e1284762..560afb3648 100644
--- a/src/shader_recompiler/runtime_info.h
+++ b/src/shader_recompiler/runtime_info.h
@@ -7,6 +7,7 @@
 #pragma once
 
 #include <array>
+#include <bitset>
 #include <map>
 #include <optional>
 #include <vector>
@@ -85,6 +86,8 @@ struct TransformFeedbackVarying {
 struct RuntimeInfo {
     std::array<AttributeType, 32> generic_input_types{};
     VaryingState previous_stage_stores;
+    /// Fragment stage after a mesh program: generic inputs it writes per primitive.
+    std::bitset<32> per_primitive_generics{};
     std::map<IR::Attribute, IR::Attribute> previous_stage_legacy_stores_mapping;
 
     bool convert_depth_mode{};
diff --git a/src/shader_recompiler/shader_info.h b/src/shader_recompiler/shader_info.h
index 71309e040e..ea70b2913e 100644
--- a/src/shader_recompiler/shader_info.h
+++ b/src/shader_recompiler/shader_info.h
@@ -239,11 +239,26 @@ struct ImageDescriptor {
 };
 using ImageDescriptors = boost::container::small_vector<ImageDescriptor, 4>;
 
+/// Switch 2 (NVN) task and mesh programs: the sizes their mesh registers set and, for mesh
+/// programs, the attribute (IR::Attribute index) held by each 128-byte line of a 32-element group
+/// in the emulated stage buffer: per-vertex lines from the program's output map, per-primitive
+/// lines from the pass-through geometry header bound with it, in attribute order.
+struct MeshInfo {
+    u32 payload_words{}; ///< Task output region / mesh input region, in 32-bit words
+    u32 max_vertices{};
+    u32 max_primitives{};
+    u32 topology{}; ///< SET_MESH_SHADER_A output topology (4 = triangles)
+    boost::container::static_vector<u32, 136> vertex_lines;
+    boost::container::static_vector<u32, 136> primitive_lines;
+};
+
 struct Info {
     static constexpr size_t MAX_INDIRECT_CBUFS{14};
     static constexpr size_t MAX_CBUFS{18};
     static constexpr size_t MAX_SSBOS{32};
 
+    bool is_task{}; ///< Switch 2 task program (Stage::Task)
+    bool is_mesh{}; ///< Switch 2 mesh program (Stage::Mesh)
     bool uses_workgroup_id{};
     bool uses_local_invocation_id{};
     bool uses_invocation_id{};
diff --git a/src/shader_recompiler/stage.h b/src/shader_recompiler/stage.h
index 9af0a1b720..34172cfee2 100644
--- a/src/shader_recompiler/stage.h
+++ b/src/shader_recompiler/stage.h
@@ -17,6 +17,11 @@ enum class Stage : u32 {
     Compute,
 
     VertexA,
+
+    // Switch 2 (NVN) task and mesh programs, bound in the vertex and tessellation slots while
+    // SET_MESH_CONTROL is enabled.
+    Task,
+    Mesh,
 };
 constexpr u32 MaxStageTypes = 6;