diff --git a/dm/gfx/nvn/dmsh_abi.h b/dm/gfx/nvn/dmsh_abi.h
index 68825c0365..27941882bf 100644
--- a/dm/gfx/nvn/dmsh_abi.h
+++ b/dm/gfx/nvn/dmsh_abi.h
@@ -31,6 +31,8 @@
 #define DMSH_C0_SSBO_MS 0x770            // NVN2 mesh stage storage buffer table: its mesh shaders read c0 0x770, 0x780, 0x790 and 0x7b0, and it binds slots from 0
 #define DMSH_MESH_DRAW_SLOT 29           // a task or object function's buffer: the draw's {count, first} workgroups, NVN's DrawMeshTasksIndirect record (dmsh_meshstage.py, nvn_geometry.inc mtk_encode)
 #define DMSH_MESH_GRID 1024              // Metal's limit of threadgroups in a mesh grid: a mesh shader with no task shader is drawn by an object function launching blocks of this many (dmsh_meshstage.py)
+#define DMSH_MESH_PAYLOAD_MAX 16384      // Metal's limit of an object function's payload: a task shader's larger payload goes through a device buffer (dmsh_meshstage.py spill)
+#define DMSH_MESH_SPILL_SLOT 11          // that buffer, at this index of the object and vertex (mesh) tables: word 0 its slots, then one slice of the payload's words rounded to 4 a task workgroup (nvn_geometry.inc mtk_encode)
 #define DMSH_C0_DISPATCH 0x200           // compute: the dispatch's workgroup counts x, y, z (32 bits each; Link's Awakening's light culling)
 #define DMSH_C0_DISPATCH_NVN2 4          // NVN2 compute keeps them here, below its per-shader layout at 0x10 (x and y read by Fire Emblem Fortune's Weave's light culling; z by their order)
 #define DMSH_C0_CS_UBO 0x210             // compute: uniform slot r's address and size at 0x210 + 0x10 r, read through global memory (the build's ubo_ptr bindings); it ends where the storage buffer table starts
diff --git a/dm/gfx/nvn/dmsh_chain.h b/dm/gfx/nvn/dmsh_chain.h
index 8e31bc7099..53897756c6 100644
--- a/dm/gfx/nvn/dmsh_chain.h
+++ b/dm/gfx/nvn/dmsh_chain.h
@@ -1,3 +1,3 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header): the B shader library translation chain version this source tree expects (default options), one for every game;
 // the host loads only <library dir>/b-<V>/. Run dmsh_chain.py --header after changing the chain's tools and commit this file with the change
-#define DMSH_CHAIN "bc8830d4ba5a"
+#define DMSH_CHAIN "eb9a3f4c3394"
diff --git a/dm/gfx/nvn/nvn_geometry.inc b/dm/gfx/nvn/nvn_geometry.inc
index 5a487ad6fc..cc71877f71 100644
--- a/dm/gfx/nvn/nvn_geometry.inc
+++ b/dm/gfx/nvn/nvn_geometry.inc
@@ -556,6 +556,11 @@ static void geo_mesh_encode(id<MTLRenderCommandEncoder> en_) {
 // equivalent fallback: the demanded pipeline must finish compiling before its first draw, so each one built at a draw is recorded as an M row (pso_user.tsv) and the prewarm builds
 // the rows before the game reaches them (docs/reference/shaders.md, Native mesh pipelines compiled at their first draw)
 static struct { int on, miss; NvShader *ts, *ms; uint32_t first, count; uint64_t ind, ind_rec; } g_mtk;   // ind: an indirect draw's threadgroups (mdx_mask), ind_rec: its {count, first} record
+// A task shader's payload past Metal's object payload (DMSH_MESH_PAYLOAD_MAX) goes through a buffer at DMSH_MESH_SPILL_SLOT of both stages (dmsh_meshstage.py spill): the Metal payload is
+// the slot's number only. The buffer has a slice for each task workgroup of the draw, up to MTK_SPILL_MAX; an indirect draw's count is written on the GPU, so it gets MTK_SPILL_MAX
+enum { MTK_SPILL_MAX = 1024 };
+static int mtk_spill(const NvShader *ts) { return ts && ts->mesh_payload > DMSH_MESH_PAYLOAD_MAX; }
+static NSUInteger mtk_payload_len(const NvShader *ts) { return !ts ? 16 : mtk_spill(ts) ? 16 : (ts->mesh_payload + 15u) & ~15u; }
 static void draw_(const p_draw *dw, const p_draw_idx *di);
 static uint64_t g_mtk_n, g_mtk_wait;
 static nvn_hmap g_pso_mtk = {.klen = sizeof(pso_key)};   // only the submit thread uses it: prewarm workers build pipelines (mtk_build) and the submit thread claims them (warm_claim)
@@ -574,7 +579,7 @@ static MTL4MeshRenderPipelineDescriptor *mtk_desc4(MTLRenderPipelineDescriptor *
     if (!mf || !of || !f4) return nil;
     MTL4MeshRenderPipelineDescriptor *m = [MTL4MeshRenderPipelineDescriptor new];
     m.meshFunctionDescriptor = mf; m.fragmentFunctionDescriptor = f4; m.maxTotalThreadsPerMeshThreadgroup = ms->mesh_threads;
-    m.objectFunctionDescriptor = of; m.payloadMemoryLength = ts ? (ts->mesh_payload + 15u) & ~15u : 16; m.maxTotalThreadsPerObjectThreadgroup = ts ? ts->mesh_threads : 1;   // with no task: a block's first workgroup, one thread a block
+    m.objectFunctionDescriptor = of; m.payloadMemoryLength = mtk_payload_len(ts); m.maxTotalThreadsPerObjectThreadgroup = ts ? ts->mesh_threads : 1;   // with no task: a block's first workgroup, one thread a block
     m.rasterSampleCount = d.rasterSampleCount; m.alphaToCoverageState = d.alphaToCoverageEnabled ? MTL4AlphaToCoverageStateEnabled : MTL4AlphaToCoverageStateDisabled;
     for (int i = 0; i < 8; i++) pso_col4(m.colorAttachments[i], d.colorAttachments[i]);
     return m;
@@ -604,7 +609,7 @@ static id<MTLRenderPipelineState> mtk_build(const pso_key *k, NvShader *ts, NvSh
     }
     MTLMeshRenderPipelineDescriptor *m = [MTLMeshRenderPipelineDescriptor new];
     m.meshFunction = fn_dev(ms->fn); m.fragmentFunction = fn_dev(ff); m.maxTotalThreadsPerMeshThreadgroup = ms->mesh_threads;
-    m.objectFunction = fn_dev(ts ? ts->fn : ms->mesh_obj); m.payloadMemoryLength = ts ? (ts->mesh_payload + 15u) & ~15u : 16; m.maxTotalThreadsPerObjectThreadgroup = ts ? ts->mesh_threads : 1;
+    m.objectFunction = fn_dev(ts ? ts->fn : ms->mesh_obj); m.payloadMemoryLength = mtk_payload_len(ts); m.maxTotalThreadsPerObjectThreadgroup = ts ? ts->mesh_threads : 1;
     m.rasterSampleCount = d.rasterSampleCount; m.alphaToCoverageEnabled = d.alphaToCoverageEnabled;
     for (int i = 0; i < 8; i++) m.colorAttachments[i] = d.colorAttachments[i];
     m.depthAttachmentPixelFormat = d.depthAttachmentPixelFormat; m.stencilAttachmentPixelFormat = d.stencilAttachmentPixelFormat;
@@ -659,6 +664,15 @@ static void mtk_encode(id<MTLRenderCommandEncoder> en_) {
     if (rec_on()) rec_table_op(RO_T_ADDR, g_ot, draw_gpu, DMSH_MESH_DRAW_SLOT); else [g_ot setAddress:draw_gpu atIndex:DMSH_MESH_DRAW_SLOT];
     g_args_dirty[4] = 1; encode_args(en, ms, 4, 1);
     // A task shader's threadgroups are the draw's workgroups; with no task, the object function's are blocks of DMSH_MESH_GRID of them (mdx_mask writes their count at +16)
+    if (mtk_spill(ts)) {   // the task's payload slices: word 0 their count, the slices from word 4
+        uint32_t words = ((ts->mesh_payload + 15u) & ~15u) >> 2, cap = g_mtk.ind ? MTK_SPILL_MAX : MIN(g_mtk.count, (uint32_t)MTK_SPILL_MAX);
+        if (!g_mtk.ind && g_mtk.count > MTK_SPILL_MAX) { static int n; if (n++ < 4) MLOG("[RED] task shader %s: %u workgroups, payload slices for %u: the rest launch no mesh workgroups", ts->key, g_mtk.count, cap); }
+        id<MTLBuffer> sb; size_t so = arena_alloc(16 + (size_t)MAX(cap, 1u) * words * 4, &sb); *(uint32_t *)((uint8_t *)sb.contents + so) = cap;
+        uint64_t sg = sb.gpuAddress + so;
+        if (rec_on()) { rec_table_op(RO_T_ADDR, g_ot, sg, DMSH_MESH_SPILL_SLOT); rec_table_op(RO_T_ADDR, g_vt, sg, DMSH_MESH_SPILL_SLOT); }
+        else { [g_ot setAddress:sg atIndex:DMSH_MESH_SPILL_SLOT]; [g_vt setAddress:sg atIndex:DMSH_MESH_SPILL_SLOT]; }
+        static int said; if (!said++) MLOG("task shader %s: its %u-byte payload goes through a buffer (%u workgroups, %u-byte slices)", ts->key, ts->mesh_payload, g_mtk.ind ? 0 : g_mtk.count, words * 4);
+    }
     MTLSize to = MTLSizeMake(ts ? ts->mesh_threads : 1, 1, 1), tm = MTLSizeMake(ms->mesh_threads, 1, 1);
     if (g_mtk.ind) [en drawMeshThreadgroupsWithIndirectBuffer:g_mtk.ind + (ts ? 0 : 16) threadsPerObjectThreadgroup:to threadsPerMeshThreadgroup:tm];
     else [en drawMeshThreadgroups:MTLSizeMake(ts ? g_mtk.count : (g_mtk.count + DMSH_MESH_GRID - 1) / DMSH_MESH_GRID, 1, 1) threadsPerObjectThreadgroup:to threadsPerMeshThreadgroup:tm];
diff --git a/dm/gfx/tests/test_dmsh_meshstage.py b/dm/gfx/tests/test_dmsh_meshstage.py
index ee383b6b6a..45007ff8c4 100644
--- a/dm/gfx/tests/test_dmsh_meshstage.py
+++ b/dm/gfx/tests/test_dmsh_meshstage.py
@@ -33,6 +33,16 @@
     _23[%d] = (*spvDescriptorSet0.cp_c3)._m0[0];
 }
 """
+# A spill mesh shader's kernel: the translator's dm_stage_in follows the stage memory, and the body reads it (ISBERD 0x38)
+KERNEL_IN = HEAD + """
+kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
+{
+    threadgroup spvUnsafeArray<uint, %d> _23;
+    threadgroup spvUnsafeArray<uint, %d> _24;
+    if (gl_LocalInvocationID.x != 0u) return;
+    _23[%d] = _24[int(gl_LocalInvocationID.x) + 1] + _24[0];
+}
+"""
 
 
 def sph(pos=0xf0, generics=b''):
@@ -61,9 +71,18 @@ def test_per_primitive_sph(self):   # after the instructions, at control +0x64 w
         self.assertEqual((i['kind'], i['threads'], i['payload'], i['task'], i['max_v'], i['max_p'], i['primitive']), ('ms', 8, 1408, True, 32, 16, [(0, 0)]))
         self.assertEqual(ms.info(control(7, max_v=128, max_p=64), area)['primitive'], [])
 
-    def test_task_memory(self):   # the payload's bytes hold the mesh count at +4, as the game's stores do
+    def test_task_memory(self):   # the payload's bytes hold the mesh count at +4, as the game's stores do. A task reads no input: its outputs follow word 0's block
         lay, size = ms.memory(ms.info(control(6), bytes(0x80)))
-        self.assertEqual(lay, [0, 0x10, 0x110, 0x210]); self.assertEqual(size, 0x210)
+        self.assertEqual(lay, [0, 0x10, 0x10, 0x110]); self.assertEqual(size, 0x110)
+
+    def test_spill_memory(self):   # Nintendo Switch Sports Resort's 07abb052a670b8ba / d60092d7424cb5b4: 27,904 payload bytes, past Metal's 16 KB
+        t = ms.info(control(6, payload=27904), bytes(0x80))
+        self.assertTrue(ms.spill(t)); self.assertFalse(ms.spill(ms.info(control(6, payload=abi.MESH_PAYLOAD_MAX), bytes(0x80))))
+        self.assertEqual(ms.memory(t), ([0, 0x10, 0x10, 0x10 + 27904], 0x10 + 27904))
+        m = ms.info(control(7, payload=27904, task=1, max_v=64, max_p=126), sph(0xf0, b'\x30'))
+        lay, size = ms.memory(m)
+        self.assertEqual(lay[2], 0x10); self.assertEqual(lay[4], 27904 // 4); self.assertLessEqual(size, ms.METAL_MAX_SHARED)   # the input is not in the stage memory
+        self.assertEqual(ms.spill_words(m), 6976)
 
     def test_mesh_memory(self):   # 128 vertices: four blocks of 32. Six components a block, 0x80 bytes each; the primitive bytes follow, 3 a triangle after the count's word
         i = ms.info(control(7, payload=128, max_v=128, max_p=64), sph(0xf0, b'\x30'))
@@ -93,10 +112,16 @@ def test_own_shared_memory(self):
 
 
 class Stage(unittest.TestCase):
-    def task(self):
-        i = ms.info(control(6), bytes(0x80))
+    def task(self, payload=256):
+        i = ms.info(control(6, payload=payload), bytes(0x80))
         return ms.to_stage(KERNEL % (ms.memory(i)[1] // 4, ms.memory(i)[0][2] // 4 + 1), i), i
 
+    def spill_mesh(self, with_input=True):
+        i = ms.info(control(7, payload=27904, task=1, max_v=64, max_p=126), sph(0xf0, b'\x30'))
+        lay, size = ms.memory(i)
+        src = KERNEL_IN % (size // 4, lay[4], lay[3] // 4) if with_input else KERNEL % (size // 4, lay[3] // 4)
+        return ms.to_stage(src, i), i
+
     def mesh(self, task=False, primitive=False, vertices=128, primitives=64, pos=0xf0, verts=0, reads=None):
         area = sph(pos, b'\x30') + (sph(0, b'\x01') if primitive else b'')
         i = ms.info(control(7, payload=128, task=int(task), max_v=vertices, max_p=primitives, prim_at=0x80 if primitive else 0, verts=verts), area)
@@ -135,6 +160,27 @@ def test_mesh_count_and_indices(self):
         self.assertIn(f'uint np = min(_23[{prim >> 2}] >> 24, 64u);', src)
         self.assertIn(f'uint at = {prim + 4}u + n;', src)
 
+    def test_spill_task_stores_its_payload_in_its_slice(self):
+        src, i = self.task(payload=27904); out = ms.memory(i)[0][2] // 4
+        self.assertIn(f'device uint* dm_spill [[buffer({abi.MESH_SPILL_SLOT})]]', src)
+        self.assertIn('struct dm_payload_t { uint slot; };', src)
+        self.assertIn('if (dm_tg >= dm_spill[0])', src)   # past the host's slices: no launch
+        self.assertIn('device uint* dm_slice = dm_spill + 4u + dm_tg * 6976u;', src)
+        self.assertIn(f'dm_slice[i] = _23[{out}u + i];', src)
+        self.assertIn('dm_payload.slot = dm_tg;', src)
+        self.assertNotIn('dm_payload.w', src)
+
+    def test_spill_mesh_reads_its_input_from_the_slice(self):
+        src, i = self.spill_mesh()
+        self.assertNotIn('threadgroup spvUnsafeArray<uint, 6976> _24;', src)   # the array left the body
+        self.assertIn('dm_in_t _24)', src)
+        self.assertIn('dm_in_t dm_in = { dm_spill + 4u + dm_payload.slot * 6976u, dm_tg };', src)   # word 0: the workgroup's number among its task's launches
+        self.assertIn(', _23, dm_in);', src)
+        self.assertIn(f'const device uint* dm_spill [[buffer({abi.MESH_SPILL_SLOT})]]', src)
+
+    def test_spill_mesh_without_its_input_refused(self):
+        with self.assertRaises(ValueError): self.spill_mesh(with_input=False)
+
     def test_mesh_with_a_task_reads_its_number_from_the_payload(self):
         src = self.mesh(task=True)[0]
         self.assertIn('_23[4] = dm_tg;', src)
@@ -217,7 +263,7 @@ def test_position_is_invariant(self):   # two mesh shaders of one geometry must
 
     @unittest.skipUnless(shutil.which('xcrun') and subprocess.run(['xcrun', '-sdk', 'macosx', '-f', 'metal'], capture_output=True).returncode == 0, 'no Metal compiler')
     def test_compiles(self):
-        cases = [self.task()[0], self.mesh()[0], self.mesh(task=True)[0], self.mesh(primitive=True)[0], self.mesh(vertices=32, primitives=32, pos=0xf8, verts=1, reads={5: 0x7})[0], self.mesh(primitive=True, reads={0: 1, 1: 0xf, 5: 0xf})[0], self.local_word()]
+        cases = [self.task()[0], self.task(payload=27904)[0], self.spill_mesh()[0], self.mesh()[0], self.mesh(task=True)[0], self.mesh(primitive=True)[0], self.mesh(vertices=32, primitives=32, pos=0xf8, verts=1, reads={5: 0x7})[0], self.mesh(primitive=True, reads={0: 1, 1: 0xf, 5: 0xf})[0], self.local_word()]
         with tempfile.TemporaryDirectory() as d:
             for n, src in enumerate(cases):
                 f = Path(d, f's{n}.metal'); f.write_text(src)
diff --git a/dm/gfx/tools/dmsh_chain.lock b/dm/gfx/tools/dmsh_chain.lock
index 1e4a8ddaa8..ef481fc5e4 100644
--- a/dm/gfx/tools/dmsh_chain.lock
+++ b/dm/gfx/tools/dmsh_chain.lock
@@ -1,5 +1,5 @@
 # Translation chain tool fingerprints (written by tool_lock.py --update-lock; regenerate after changing tools and commit it)
-dmsh_xlate	0406b50b2856b1b4
+dmsh_xlate	cebc4fc4564b4a4f
 xlate_source	657232fd201f16a1
 spirv-cross	e37945d8e441d34d
 metal_macosx	Apple metal version 32023.922 (metalfe-32023.922.4)
diff --git a/dm/gfx/tools/dmsh_meshstage.py b/dm/gfx/tools/dmsh_meshstage.py
index ad2d4f2223..d617eda610 100644
--- a/dm/gfx/tools/dmsh_meshstage.py
+++ b/dm/gfx/tools/dmsh_meshstage.py
@@ -16,6 +16,12 @@
                 or a mesh shader's attributes (0xf8): all blocks of 32 vertices (the components the SPH's OMAP enables, 0x80 bytes each, a word per vertex), followed by all blocks of 32
                 primitives (the components the per-primitive SPH's OMAP enables, 0x80 bytes each, a word per primitive)
   primitives    a mesh shader's primitive bytes: the count at +3 (0x88), index bytes from +4 (0x80, control +0x5f per primitive: 1 point, 2 line, 3 triangle)
+
+A payload past Metal's (abi.MESH_PAYLOAD_MAX, 16 KB) goes through a device buffer the host binds at abi.MESH_SPILL_SLOT in both stages (spill): word 0 its slots, then a slice of the
+payload's words (rounded to 4) for each of the draw's task workgroups. The task stores its payload in its threadgroup's slice and hands its mesh threadgroups the slice's number in the
+Metal payload; the mesh shader's input is not in its stage memory then but a shared memory of its own (job stage_memory's fifth entry, the translator's dm_stage_in), read from the
+slice. Nintendo Switch Sports Resort's task shaders hand 21-28 KB to their mesh shaders: with the payload in Metal's payload and the mesh shader's input in its threadgroup memory,
+both stages were past Metal's limits and the build refused their six shaders. A task's stage memory has no input area either (a task reads no ISBERD 0x38): its outputs follow word 0
 """
 import re
 import struct
@@ -29,6 +35,16 @@
 METAL_MAX_SHARED = 0x8000   # threadgroup memory of a Metal function
 
 
+def spill(i):
+    """Whether a task or mesh shader's payload goes through the device buffer (the module's docstring)."""
+    return i['payload'] > abi.MESH_PAYLOAD_MAX
+
+
+def spill_words(i):
+    """The words of one task workgroup's slice of that buffer."""
+    return (i['payload'] + 15 & ~15) >> 2
+
+
 def is_stage(ctl):
     """'ts', 'ms' or None, by the control section's stage byte."""
     return {6: 'ts', 7: 'ms'}.get(ctl[0x56]) if len(ctl) > 0x56 else None
@@ -60,9 +76,10 @@ def info(ctl, area):
 
 def memory(i):
     """(the stage memory's offsets for the translator: the workgroup's number, the input, the outputs, the primitive bytes; the shared memory's bytes)."""
-    out = INPUT + (i['payload'] + 15 & ~15)
     if i['kind'] == 'ts':
+        out = INPUT
         return [WORKGROUP, INPUT, out, out + i['payload']], out + (i['payload'] + 3 & ~3)
+    out = INPUT if spill(i) else INPUT + (i['payload'] + 15 & ~15)
     # The guest's attribute stores: all vertex blocks of 32 vertices first, then all primitive blocks of 32 primitives, each block its own SPH's components (Ocarina of Time's
     # ee4ef51a4fddc28e, 128 vertices of 6 components and 64 primitives of 4, writes vertex v at (v / 32) * 768 + (v % 32) * 4 and primitive p at 3072 + (p / 32) * 512 + (p % 32) * 4)
     v_blocks = (max(i['max_v'], 1) + BLOCK - 1) // BLOCK
@@ -70,7 +87,7 @@ def memory(i):
     prim = out + max(v_blocks * len(i['vertex']) + p_blocks * len(i['primitive']), 1) * BLOCK * 4
     # The guest writes the index bytes of primitive min(count - 1, lane + 32k) in each of its primitive blocks, and a meshlet count of 0 wraps count - 1, so every lane of every primitive
     # block writes: the bytes are reserved for whole blocks of 32 primitives, a primitive's vertices each (a mesh with max_p 126 wrote past 4 + 3 * 126, into the next workgroup's memory)
-    return [WORKGROUP, INPUT, out, prim], prim + (4 + i.get('verts', 3) * p_blocks * BLOCK + 3 & ~3)
+    return [WORKGROUP, INPUT, out, prim] + ([spill_words(i)] if spill(i) else []), prim + (4 + i.get('verts', 3) * p_blocks * BLOCK + 3 & ~3)
 
 
 def own_shared(code):
@@ -82,6 +99,8 @@ def own_shared(code):
 # The kernel, its parameters, the thread variables spirv-cross declares ahead of the threadgroup memory (a private variable's initializer: GlobalToStorage.TryGenerateRawStorageOp's last entry,
 # uint _N = 0u;), and its stage memory, the first threadgroup array
 _KERNEL = re.compile(r'kernel\s+void\s+main0\s*\(((?:[^()]|\([^()]*\))*)\)\s*\{\s*((?:(?!threadgroup\b)[^\n{};]+;\n\s*)*)threadgroup\s+spvUnsafeArray<uint,\s*(\d+)>\s+(\w+);\n')
+# A spill mesh shader's input (the translator's dm_stage_in), the threadgroup array that follows the stage memory
+_STAGE_IN = re.compile(r'\s*threadgroup\s+spvUnsafeArray<uint,\s*\d+>\s+(\w+);\n')
 _NAME = re.compile(r'(\w+)\s*(?:\[\[.*\]\])?$')
 _THREAD = re.compile(r'(uint3?)\s+(\w+)\s*\[\[(thread_position_in_threadgroup|threadgroup_position_in_grid)\]\]$')
 
@@ -141,11 +160,20 @@ def to_stage(msl, i, reads=None):
     params, local, words, mem = _params(m[1]), m[2].strip(), int(m[3]), m[4]
     names = [_name(p) for p in params]
     plain = [re.sub(r'\s*\[\[.*\]\]$', '', p) for p in params]
+    sp, rest = spill(i), msl[m.end():]
+    # A spill mesh shader's input: its array leaves the body, which reads the task's slice through dm_in_t instead (word 0 its workgroup's number, as the hardware leaves it)
+    sin = None
+    if sp and kind == 'ms' and i['task']:
+        n = _STAGE_IN.match(rest)
+        if n: sin, rest = n[1], rest[n.end():]
+    body_params = plain + [f'threadgroup spvUnsafeArray<uint, {words}>& {mem}'] + ([f'dm_in_t {sin}'] if sin else [])
     # The kernel's body, called with the stage memory the wrapper owns: its early returns (a thread with nothing to write) come back to the wrapper's copy of the outputs. Its thread variables
     # stay in it: one call a thread, as the kernel ran
-    src = (msl[:m.start()] + f'static inline __attribute__((always_inline)) void dm_stage_body({", ".join(plain + [f"threadgroup spvUnsafeArray<uint, {words}>& {mem}"])})\n{{\n' +
-           (f'    {local}\n' if local else '') + msl[m.end():])
-    call = f'dm_stage_body({", ".join(names + [mem])});'
+    in_t = ('struct dm_in_t { const device uint* p; uint w0; uint operator[](int k) const { return k ? p[k] : w0; } uint operator[](uint k) const { return k ? p[k] : w0; } };\n'
+            if sin else '')
+    src = (msl[:m.start()] + in_t + f'static inline __attribute__((always_inline)) void dm_stage_body({", ".join(body_params)})\n{{\n' +
+           (f'    {local}\n' if local else '') + rest)
+    call = f'dm_stage_body({", ".join(names + [mem] + (["dm_in"] if sin else []))});'
     # The thread's and the threadgroup's numbers: the kernel's parameters where it has them (an attribute may appear once in a function), else parameters of the wrapper's own
     own = {t[3]: t[2] + ('.x' if t[1] == 'uint3' else '') for t in map(_THREAD.match, params) if t}
     lid, tg = own.get('thread_position_in_threadgroup', 'dm_lid'), own.get('threadgroup_position_in_grid', 'dm_tg')
@@ -156,16 +184,23 @@ def to_stage(msl, i, reads=None):
     start = (f'    threadgroup spvUnsafeArray<uint, {words}> {mem};\n'
              f'    for (uint i = {lid}; i < {words}u; i += {thr}u) {mem}[i] = 0u;\n'
              '    threadgroup_barrier(mem_flags::mem_threadgroup);\n')
-    head = f'struct dm_payload_t {{ uint w[{max(pay, 1)}]; }};\n'
-    sig = ', '.join(params + extra)
+    slice_w = spill_words(i)
+    head = 'struct dm_payload_t { uint slot; };\n' if sp else f'struct dm_payload_t {{ uint w[{max(pay, 1)}]; }};\n'
+    sig = ', '.join(params + extra + ([f'{"device" if kind == "ts" else "const device"} uint* dm_spill [[buffer({abi.MESH_SPILL_SLOT})]]'] if sp else []))
     draw = f'constant uint* dm_draw [[buffer({abi.MESH_DRAW_SLOT})]]'   # the draw's {count, first}, which the host binds (nvn_geometry.inc mtk_encode)
     if kind == 'ts':
+        # A spill task past the buffer's slots launches nothing (the host sizes it to the draw; an indirect draw's count is not known when it is encoded)
+        store = (f'    if ({tg} >= dm_spill[0]) {{ if ({lid} == 0u) dm_grid.set_threadgroups_per_grid(uint3(0u, 1u, 1u)); return; }}\n'
+                 f'    device uint* dm_slice = dm_spill + 4u + {tg} * {slice_w}u;\n'
+                 f'    for (uint i = {lid}; i < {pay}u; i += {thr}u) dm_slice[i] = {mem}[{out}u + i];\n'
+                 '    threadgroup_barrier(mem_flags::mem_device);\n'
+                 f'    if ({lid} == 0u) dm_payload.slot = {tg};\n') if sp else f'    for (uint i = {lid}; i < {pay}u; i += {thr}u) dm_payload.w[i] = {mem}[{out}u + i];\n'
         return src + '\n' + head + (
             f'[[object]] void main0({sig}, object_data dm_payload_t& dm_payload [[payload]], mesh_grid_properties dm_grid, {draw})\n{{\n{start}'
             f'    if ({lid} == 0u) {mem}[{wg}] = dm_draw[1] + {tg};\n'
-            f'    threadgroup_barrier(mem_flags::mem_threadgroup);\n    {call}\n    threadgroup_barrier(mem_flags::mem_threadgroup);\n'
-            f'    for (uint i = {lid}; i < {pay}u; i += {thr}u) dm_payload.w[i] = {mem}[{out}u + i];\n'
+            f'    threadgroup_barrier(mem_flags::mem_threadgroup);\n    {call}\n    threadgroup_barrier(mem_flags::mem_threadgroup);\n' + store +
             f'    if ({lid} == 0u) dm_grid.set_threadgroups_per_grid(uint3(min({mem}[{out + 1}], {METAL_MAX_GRID}u), 1u, 1u));\n}}\n')
+    if sp and not sin: raise ValueError('ms: a payload past Metal\'s with no task, or no dm_stage_in array after the stage memory')
     vertex, primitive = i['vertex'], i['primitive']
     max_v, max_p = i['max_v'], i['max_p']
     verts = i.get('verts', 3)
@@ -202,7 +237,8 @@ def to_stage(msl, i, reads=None):
 
     def entry(name, gl):
         # With a task the payload is the input, its word 0 the workgroup's number among the task's launches, as the hardware leaves it (+0 is not a word the task writes: it starts at +4)
-        fill = (f'    for (uint i = {lid}; i < {pay}u; i += {thr}u) {mem}[{inp}u + i] = dm_payload.w[i];\n'
+        fill = (f'    dm_in_t dm_in = {{ dm_spill + 4u + dm_payload.slot * {slice_w}u, {tg} }};\n    if ({lid} == 0u) {mem}[{wg}] = {tg};\n' if sin else
+                f'    for (uint i = {lid}; i < {pay}u; i += {thr}u) {mem}[{inp}u + i] = dm_payload.w[i];\n'
                 '    threadgroup_barrier(mem_flags::mem_threadgroup);\n'
                 f'    if ({lid} == 0u) {{ {mem}[{wg}] = {tg}; {mem}[{inp}] = {tg}; }}\n') if i['task'] else f'    if ({lid} == 0u) {mem}[{wg}] = dm_block.first + {tg};\n'
         return (f'[[mesh]] void {name}({sig}, dm_mesh_t dm_out' + (', const object_data dm_payload_t& dm_payload [[payload]]' if i['task'] else ', const object_data dm_block_t& dm_block [[payload]]') +
diff --git a/dm/gfx/tools/dmsh_xlate_bin/dmsh_xlate b/dm/gfx/tools/dmsh_xlate_bin/dmsh_xlate
index 9f265792d3..0bb1c34677 100755
Binary files a/dm/gfx/tools/dmsh_xlate_bin/dmsh_xlate and b/dm/gfx/tools/dmsh_xlate_bin/dmsh_xlate differ
diff --git a/dm/third_party/ryujinx/README.md b/dm/third_party/ryujinx/README.md
index 9e0f4c1609..4ba99032ba 100644
--- a/dm/third_party/ryujinx/README.md
+++ b/dm/third_party/ryujinx/README.md
@@ -17,5 +17,7 @@ Three projects from Ryujinx 1.3.2: `Ryujinx.Graphics.Shader`, `Ryujinx.Common` a
 
 - `Ryujinx.Graphics.Shader/CodeGen/Spirv/Instructions.cs` (`TryGenerateBoundedUniformLoad`), `Translation/ResourceManager.cs` (`NoteUniformRead`), `StructuredIr/StructuredProgram.cs`, `StructuredIr/StructuredProgramContext.cs`, `StructuredIr/ShaderProperties.cs`, `BufferDescriptor.cs`, `Constants.cs` and `Translation/TranslationFlags.cs`: the Metal path (`TranslationFlags.MetalUniformBounds`) gives 0 for a read of a game uniform slot (c[3] to c[16]) past the size the game bound, as Maxwell does. A read at a computed offset compares the word's index with the slot's size from a table in bank 25 that the host fills; the reads at fixed offsets are reported as an extent of the slot (`BufferDescriptor.FixedSize`), up to which the host gives a shorter binding zeros. Upstream relied on Vulkan's robust buffer access, which Metal does not have. See [reads past a uniform binding's size](../../../docs/reference/shaders.md#reads-past-a-uniform-bindings-size), found on Paper Mario: The Thousand-Year Door (GitHub issue #1860) and Donkey Kong Bananza (GitHub issue #2066).
 
+- `Ryujinx.Graphics.Shader/Instructions/InstEmitAmpereMem.cs` (`AIsberd`) and `Translation/ResourceManager.cs` (`GetStageInputMemoryId`): when the job's stage memory has a fifth entry, a mesh shader's stage input (ISBERD mode 0x38) is a shared memory of its own, `dm_stage_in`, which `dm/gfx/tools/dmsh_meshstage.py` replaces with the device buffer its task stored the payload in. Metal hands a mesh threadgroup at most 16 KB of payload; Nintendo Switch Sports Resort's task shaders hand theirs 21-28 KB
+
 The shader translation tool in `dm/gfx/tools/dmsh_xlate` references them to translate the original Switch shaders to SPIR-V.
 License: MIT, see `LICENSE.txt`. Upstream: <https://github.com/Ryubing>.
diff --git a/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Instructions/InstEmitAmpereMem.cs b/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Instructions/InstEmitAmpereMem.cs
index 6795611cc6..42ee3725f6 100644
--- a/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Instructions/InstEmitAmpereMem.cs
+++ b/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Instructions/InstEmitAmpereMem.cs
@@ -142,6 +142,19 @@ public static void AIsberd(EmitterContext context)
                     throw AFail(context, $"ISBERD mode {mode:x2} in a task or mesh shader");
                 }
 
+                if (mode == 0x38 && stage.Length > 4 && stage[4] > 0)
+                {
+                    // The input kept apart (ResourceManager.GetStageInputMemoryId): a word of its own shared memory, which the wrapper reads from the task's device buffer
+                    if (x.Rd < RegisterConsts.RegisterZeroIndex)
+                    {
+                        int id = context.ResourceManager.GetStageInputMemoryId(stage[4]);
+                        Operand byteOffset = context.IAdd(AGpr(x.Ra), Const((int)x.SignedField(40, 64)));
+                        context.Copy(Register(x.Rd, RegisterType.Gpr), context.Load(StorageKind.SharedMemory, id, context.ShiftRightU32(byteOffset, Const(2))));
+                    }
+
+                    return;
+                }
+
                 Operand addr = mode == 0x78 ? Const(stage[0]) : context.IAdd(AGpr(x.Ra), Const(stage[mode == 0xb8 ? 2 : 1] + (int)x.SignedField(40, 64)));
                 EmitLoad(context, StorageKind.SharedMemory, LsSize2.B32, addr, x.Rd, 0);
                 return;
diff --git a/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/ResourceManager.cs b/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/ResourceManager.cs
index 5bc2348827..1c7700d8bf 100644
--- a/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/ResourceManager.cs
+++ b/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/ResourceManager.cs
@@ -142,6 +142,21 @@ public int GetHmmaScratchMemoryId(int subgroups)
             return HmmaScratchMemoryId;
         }
 
+        // nxnative: a mesh shader's stage input (ISBERD mode 0x38, the payload its task wrote) as a shared memory of its own, when the job keeps it apart from the stage memory
+        // (QueryAmpereStageMemory's fifth entry: its words). dmsh_meshstage.py replaces it with the device buffer the task stored its payload in: Metal hands an object
+        // function's payload to its mesh threadgroups only up to 16 KB, and with the input in the stage memory a mesh shader's threadgroup memory went past 32 KB
+        public int StageInputMemoryId { get; private set; } = -1;
+
+        public int GetStageInputMemoryId(int words)
+        {
+            if (StageInputMemoryId < 0)
+            {
+                StageInputMemoryId = Properties.AddSharedMemory(new MemoryDefinition("dm_stage_in", AggregateType.Array | AggregateType.U32, words));
+            }
+
+            return StageInputMemoryId;
+        }
+
         public void SetVertexAsComputeLocalMemories(ShaderStage stage, InputTopology inputTopology)
         {
             LocalVertexDataMemoryId = AddMemoryDefinition("local_vertex_data", AggregateType.Array | AggregateType.FP32, Reservations.OutputSizePerInvocation);
diff --git a/docs/agents/symptoms.md b/docs/agents/symptoms.md
index c314a75ead..78cebac3f5 100644
--- a/docs/agents/symptoms.md
+++ b/docs/agents/symptoms.md
@@ -956,6 +956,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - after a pair is added to `geo_mesh.tsv`, `[RED] geometry mesh pipeline: ... Total mesh size (37248) exceeds the maximum mesh size allowed (32768)`, and the frames that draw the pair get much slower than on the compute path (Donkey Kong Bananza: translation from about 10 to 22–31 ms a frame). → [Mesh geometry groups sized to the vertex output](../reference/shaders.md#mesh-geometry-groups-sized-to-the-vertex-output)
 - a depth prepass of many small geometry-shader draws, drawn as mesh draws, leaves the GPU nearly idle: the per-pass profile (`DM_NVN_GPUPROF`) shows an idle gap before the pass that grows with its draw count (about 2.3 µs a draw), and the GPU counters show vertex-only time at about 4 % occupancy. → [Mesh geometry draws gathered into one dispatch](../reference/shaders.md#mesh-geometry-draws-gathered-into-one-dispatch)
 - white blotches over Kokiri Forest's foliage, walkway and ground in Ocarina of Time, which are the house light tiles. The light-tile pass of task `a5288d0a6c4971fd` with meshes `e8c667282ba7bd9b` and `d4171cf9adf0504b` writes constant white through fragment `ea1bf96b72e72476`, and its triangles were corrupt. → [Mesh stage memory reserves the primitive index bytes for every primitive block](../reference/shaders.md#mesh-stage-memory-reserves-the-primitive-index-bytes-for-every-primitive-block)
+- the shader build refuses a task shader and the mesh shaders it launches: `mesh stage failed: shared memory of its own, or stage memory past Metal's limit` in `dmsh_build.log`, and the game's draws with them are skipped. → [A task shader's payload past Metal's 16 KB](../reference/shaders.md#a-task-shaders-payload-past-metals-16-kb)
 - a mesh shader's per-primitive outputs read as other values once it has more than 32 vertices: Ocarina of Time's Poe Sisters and Phantom Ganon paintings (mesh shaders `ee4ef51a4fddc28e` and `83c0dbc51c5481e3`, 128 vertices, 64 primitives, a per-primitive colour) came out solid green or magenta (pull request #1590). A mesh shader that draws points with a point size takes the size as its position's x. → [Mesh attributes: every vertex block before the primitive blocks](../reference/shaders.md#mesh-attributes-every-vertex-block-before-the-primitive-blocks)
 - `primitive type 14 not implemented`; patch draws skipped. → [Tessellation stages are not drawn](../reference/shaders.md#tessellation-stages-are-not-drawn)
 - a Switch 2 game's shaders do not translate (the translator reads Maxwell: 64-bit instructions behind scheduling words, a 0x50-byte SPH), or, rewritten as Maxwell code, UI text draws as solid boxes or not at all. → [Ampere shader code (Switch 2)](../reference/shaders.md#ampere-shader-code-switch-2)
diff --git a/docs/reference/shaders.md b/docs/reference/shaders.md
index e46607025d..70cb524dc4 100644
--- a/docs/reference/shaders.md
+++ b/docs/reference/shaders.md
@@ -1108,6 +1108,14 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b
 - **Fix:** `dmsh_meshstage.memory` reserves `4 + 3 * ceil(max_p / 32) * 32` bytes (rounded to a word) for the index bytes: whole primitive blocks, not the vertex blocks a mesh with more vertices than primitives has as well. Meshes of 128 vertices and 64 primitives did not write past `4 + 3 * 64` before the fix, so the guest's loop goes by primitive blocks. All 500 of Ocarina's mesh shaders were run standalone on zeroed inputs, each with a guard band past its stage array (Metal shader validation does not check static threadgroup arrays): 387 meshes wrote past the array before the fix, every `max_p` 126 mesh, and none after (pull request #1575, measured with every vertex and primitive block reserved; the primitive-block size was not run through the guard band). In the game on an M2 Pro the blotches are gone and the house doorways and windows glow as intended. This does not fix the intermittent GPU hangs seen in buffers with these meshes; those are a separate open issue.
 - **Found on:** [Ocarina of Time](../../dm/games/0400486024DEC000/README.md). The other mesh-shader games in the games table (Yoshi and the Mysterious Book, Mario Kart World, Splatoon Raiders) rebuild their shaders with the chain change; none was checked in the game.
 
+### A task shader's payload past Metal's 16 KB
+
+- **Symptom:** the shader build refuses a task shader and the mesh shaders it launches: `mesh stage failed: shared memory of its own, or stage memory past Metal's limit` in `dmsh_build.log`, and the game's draws with them are skipped.
+- **Cause:** Metal holds an object function's payload to 16 KB and a function's threadgroup memory to 32 KB. Nintendo Switch Sports Resort's task shaders hand their mesh shaders 21-28 KB. The payload went through Metal's payload, and the mesh shader copied it into its stage memory (the input area, ISBERD mode 0x38) on top of its attributes and primitives, about 33 KB. The task shader's stage memory also reserved an input area as large as its payload, although a task shader reads no input: 42-56 KB.
+- **Fix:** a payload past `DMSH_MESH_PAYLOAD_MAX` goes through a device buffer the host binds at `DMSH_MESH_SPILL_SLOT` of both stages (`nvn_geometry.inc` `mtk_encode`): word 0 the number of slices, then a slice of the payload's words for each task workgroup. The task stores its payload in its workgroup's slice and passes the slice's number in Metal's payload (`dmsh_meshstage.to_stage`). The mesh shader's input is then a shared memory of its own (the job's fifth stage memory entry; `InstEmitAmpereMem.cs` `AIsberd`, `ResourceManager.GetStageInputMemoryId`), which the wrapper reads from the slice, so it is not in the mesh shader's stage memory. A task shader's stage memory has no input area. A direct draw gets a slice for each workgroup up to 1,024 (`MTK_SPILL_MAX`; past that the rest launch no mesh workgroups, with a `[RED]` line), an indirect draw 1,024, since its count is written on the GPU.
+- **Found on:** Nintendo Switch Sports Resort (`04008E001F098000`, v65536; no profile on main): task shaders `4d57f62a663c4cc9`, `3c8eaa19f720a02f`, `07abb052a670b8ba`, `8c49b3da7fd802df` and mesh shaders `d60092d7424cb5b4`, `a68f2301bbaceb8c` (issue #1753) now build: 11,059 of 11,059 shaders. The host logs the first draw through the buffer; no play session has logged one yet, so the path has not run in the game.
+- **Check:** `python3 dm/gfx/tests/test_dmsh_meshstage.py` (the layouts, the task's slice store, the mesh shader's read of its input from the slice, and the Metal compile of both).
+
 ### Mesh attributes: every vertex block before the primitive blocks
 
 - **Symptom:** a mesh shader's per-primitive outputs read as other values once it has more than 32 vertices: Ocarina of Time's Poe Sisters and Phantom Ganon paintings (mesh shaders `ee4ef51a4fddc28e` and `83c0dbc51c5481e3`, 128 vertices, 64 primitives, a per-primitive colour) came out solid green or magenta (pull request #1590). A mesh shader that draws points with a point size takes the size as its position's x.