diff --git a/dm/gfx/nvn/nvn_mtl.m b/dm/gfx/nvn/nvn_mtl.m
index f73da6dac7..16186f4223 100644
--- a/dm/gfx/nvn/nvn_mtl.m
+++ b/dm/gfx/nvn/nvn_mtl.m
@@ -1868,6 +1868,19 @@ static void ser_note(uint64_t serial) {
 // or another texture takes over their storage (alias_take_over).
 typedef struct { uint64_t pool; int64_t off; nvn_texdesc d; } alias_key;
 static nvn_hmap g_alias = {.klen = sizeof(alias_key)};
+// How many entries of g_by_serial hold each object (key: its address), kept by every writer of g_by_serial (serial_put, serial_del, tex_put, the upgrade's in-place
+// replacement), so alias_take_over asks whether a live texture object still resolves to a shared texture in one lookup. It walked the serial table's whole capacity for
+// each candidate instead: in Ocarina of Time's cutscenes, with about 7,500 serials and thousands of shared textures, that walk was the submit thread's hottest function and
+// frames took up to a second. An entry holds its object (CFBridgingRetain), so a counted address cannot be reused while counted. Under g_smu
+static nvn_hmap g_by_serial_held = {.klen = sizeof(uintptr_t)};
+static void held_add(const void *o, int d) {
+    if (!o) return;
+    uintptr_t k = (uintptr_t)o, n = (uintptr_t)nvn_hm_get(&g_by_serial_held, &k);
+    if (d > 0) nvn_hm_put(&g_by_serial_held, &k, (void *)(n + 1));
+    else if (n > 1) nvn_hm_put(&g_by_serial_held, &k, (void *)(n - 1));
+    else if (n) nvn_hm_del(&g_by_serial_held, &k);   // a slot written past these helpers (a test's table) was never counted: nothing to take away
+}
+static int serial_table_holds(const void *o) { uintptr_t k = (uintptr_t)o; return nvn_hm_get(&g_by_serial_held, &k) != NULL; }
 static nvn_hmap g_alias_seen = {.klen = sizeof(alias_key)};   // key → g_fence + 1 at its last lookup or creation (alias_take_over spares entries looked up lately)
 // Live GPU-only image descriptions, including those not used yet. A dynamic-resolution family declares its largest image before any particular resolution is used.
 // Keeping only descriptions here lets its first allocation cover that family exactly, without allocating every resolution or rounding to a power of two.
@@ -1939,7 +1952,7 @@ static void alias_take_over(const alias_key *nk) {
         if (k.pool != nk->pool || !memcmp(&k, nk, sizeof k) || k.off >= hi) continue;
         nvn_texlayout kl; nvn_tex_layout(&k.d, &kl);
         uint64_t seen = (uint64_t)(uintptr_t)nvn_hm_get(&g_alias_seen, &k);
-        if (k.off + (int64_t)kl.size <= lo || (seen && g_fence + 1 - seen < ALIAS_IDLE) || serial_holds(&g_by_serial, (__bridge id)g_alias.vals[i])) continue;
+        if (k.off + (int64_t)kl.size <= lo || (seen && g_fence + 1 - seen < ALIAS_IDLE) || serial_table_holds(g_alias.vals[i])) continue;
         ks[n++] = k;
     }
     for (uint32_t j = 0; j < n; j++) {
@@ -1953,8 +1966,8 @@ static void alias_take_over(const alias_key *nk) {
     if (n) { g_ser_gen++; g_alias_taken += n; }
 }
 static __unsafe_unretained id serial_get(uint64_t k) { pthread_mutex_lock(&g_smu); __unsafe_unretained id o = hm_get_id(&g_by_serial, &k); pthread_mutex_unlock(&g_smu); return o; }
-static void serial_put(uint64_t k, id o) { pthread_mutex_lock(&g_smu); id old = hm_get_id(&g_by_serial, &k); if (old) { if (alias_has(old)) defer_obj_unref(old); else defer_obj_release(old); } hm_put_id(&g_by_serial, &k, o); ser_note(k); pthread_mutex_unlock(&g_smu); }
-static void serial_del(uint64_t k) { pthread_mutex_lock(&g_smu); id old = hm_get_id(&g_by_serial, &k); if (old) { if (alias_has(old)) defer_obj_unref(old); else defer_obj_release(old); hm_del_id(&g_by_serial, &k); ser_note(k); } pthread_mutex_unlock(&g_smu); }   // nothing to drop: the texture ID caches stay valid
+static void serial_put(uint64_t k, id o) { pthread_mutex_lock(&g_smu); id old = hm_get_id(&g_by_serial, &k); if (old) { held_add((__bridge void *)old, -1); if (alias_has(old)) defer_obj_unref(old); else defer_obj_release(old); } hm_put_id(&g_by_serial, &k, o); held_add((__bridge void *)o, 1); ser_note(k); pthread_mutex_unlock(&g_smu); }
+static void serial_del(uint64_t k) { pthread_mutex_lock(&g_smu); id old = hm_get_id(&g_by_serial, &k); if (old) { held_add((__bridge void *)old, -1); if (alias_has(old)) defer_obj_unref(old); else defer_obj_release(old); hm_del_id(&g_by_serial, &k); ser_note(k); } pthread_mutex_unlock(&g_smu); }   // nothing to drop: the texture ID caches stay valid
 // Texture serials owned by a live NVN texture object (TextureInitialize until Finalize or re-initialization, hle_nvn.c). Command records and texture pool registrations snapshot the
 // texture object (tex_snap in hle_nvn.c, mtl_reg_tex), so replay can resolve a texture the game has already finalized: Breath of the Wild records each frame's transient render targets
 // (the HUD captures, the 640x360 half-resolution buffers) and finalizes their objects before it submits. Those serials went back into g_by_serial and stayed there for good, as no second
@@ -1979,6 +1992,7 @@ static void tex_put(uint64_t k, id o) {   // serials from 2^62 up are external t
     pthread_mutex_lock(&g_smu);
     nvn_hmap *m = k >= (1ull << 62) || nvn_hm_get(&g_tex_live, &k) ? &g_by_serial : &g_ghost;
     id old = hm_get_id(m, &k);
+    if (old && m == &g_by_serial) held_add((__bridge void *)old, -1);
     if (old) obj_drop(old);
     else if (m == &g_ghost) {
         uint64_t *slot = &g_ghost_ring[g_ghost_next++ % NGHOST];
@@ -1986,6 +2000,7 @@ static void tex_put(uint64_t k, id o) {   // serials from 2^62 up are external t
         *slot = k; g_ghost_puts++;
     }
     hm_put_id(m, &k, o); ser_note(k);
+    if (m == &g_by_serial) held_add((__bridge void *)o, 1);
     pthread_mutex_unlock(&g_smu);
 }
 // Host record of a packaged texture: the current MTLTexture, the upload stamp, and the global write generation at the last check
@@ -2195,7 +2210,9 @@ static MTLTextureUsage img_usage(const tex_t *t, MTLPixelFormat pf) {
         for (unsigned j = 0; j < 2; j++) for (uint32_t i = 0; i < tables[j]->cap; i++) {
             if (tables[j]->vals[i] != (__bridge void *)base) continue;
             uint64_t serial; memcpy(&serial, tables[j]->keys + (size_t)i * tables[j]->klen, sizeof serial);
-            void *old = tables[j]->vals[i]; tables[j]->vals[i] = (void *)CFBridgingRetain(upgraded); CFRelease(old); named |= serial == t->serial;
+            void *old = tables[j]->vals[i]; tables[j]->vals[i] = (void *)CFBridgingRetain(upgraded);
+            if (j == 0) { held_add(old, -1); held_add((__bridge void *)upgraded, 1); }   // g_by_serial: the count moves with the name
+            CFRelease(old); named |= serial == t->serial;
         }
         for (uint32_t i = 0; i < g_alias.cap; i++) {
             if (g_alias.vals[i] != (__bridge void *)base) continue;
diff --git a/dm/gfx/tests/Makefile b/dm/gfx/tests/Makefile
index 5288b37c2b..527b8b9dca 100644
--- a/dm/gfx/tests/Makefile
+++ b/dm/gfx/tests/Makefile
@@ -404,6 +404,9 @@ check_t_unwritten_store: RUN_ENV = MTL_DEBUG_LAYER=1
 t_upgrade_alias: t_upgrade_alias.m $(MTL_T_DEP)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_upgrade_alias.m $(MTL_T_SRC) $(FW)
 check_t_upgrade_alias: RUN_ENV = MTL_DEBUG_LAYER=1
+t_alias_takeover: t_alias_takeover.m $(MTL_T_DEP)
+	clang $(CFLAGS) -fobjc-arc -o $@ t_alias_takeover.m $(MTL_T_SRC) $(FW)
+check_t_alias_takeover: RUN_ENV = MTL_DEBUG_LAYER=1
 t_srgb_alias: t_srgb_alias.m $(MTL_T_DEP)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_srgb_alias.m $(MTL_T_SRC) $(FW)
 check_t_srgb_alias: RUN_ENV = MTL_DEBUG_LAYER=1
diff --git a/dm/gfx/tests/t_alias_takeover.m b/dm/gfx/tests/t_alias_takeover.m
new file mode 100644
index 0000000000..5204c45812
--- /dev/null
+++ b/dm/gfx/tests/t_alias_takeover.m
@@ -0,0 +1,75 @@
+// alias_take_over (nvn_mtl.m): when a new shared texture takes over storage, an entry a live texture object still resolves to stays, and one nothing resolves to any more goes once
+// no lookup has found it for ALIAS_IDLE submissions. Whether a live object resolves to it is one lookup in g_by_serial_held, which every writer of g_by_serial keeps (serial_put,
+// serial_del, tex_put and the upgrade's in-place replacement): two serials naming one object count twice and the object stays held until both go. Last, a takeover past 1,000
+// held entries with 20,000 serials in the table, which walked the serial table for each entry before (it prints the time).
+#include "../nvn/nvn_mtl.m"
+void *front_metal_layer(void) { return NULL; }
+module_t *module_of(const void *p) { (void)p; return NULL; }
+module_t *g_mod; int g_nmod;
+void svc_thread_dump(void) {}
+uint64_t now_ns(void) { return bp_now(); }
+static int checks, failed;
+#define CHECK(c, label) do { checks++; if (!(c)) { failed++; fprintf(stderr, "FAIL: %s (line %d)\n", label, __LINE__); } } while (0)
+
+static tex_t tex(uint64_t serial, NVNformat format, int w) {
+    tex_t t = {.magic = MAGIC('T','E','X',0), .serial = serial};
+    t.d = (nvn_texdesc){.target = NVN_TEXTURE_TARGET_2D, .format = format, .width = w, .height = w, .depth = 1, .levels = 1, .samples = 1};
+    nvn_tex_layout(&t.d, &t.L);
+    return t;
+}
+static int shared(id<MTLTexture> t) { pthread_mutex_lock(&g_smu); int r = alias_has(t); pthread_mutex_unlock(&g_smu); return r; }
+static int held(id o) { pthread_mutex_lock(&g_smu); int r = serial_table_holds((__bridge void *)o); pthread_mutex_unlock(&g_smu); return r; }
+static void next_submission(void) { CHECK(mtl_wait(mtl_sub_commit(), 5000000000ull), "submission completes"); mtl_sub_begin(); }
+
+int main(void) { @autoreleasepool {
+    mtl_init(); mtl_sub_begin();
+    tex_t a = tex(91001, NVN_FORMAT_RGBA8, 64), b = tex(91002, NVN_FORMAT_R32F, 64), c = tex(91003, NVN_FORMAT_RG16F, 64);
+    mtl_tex_init(&a); mtl_tex_init(&b); mtl_tex_init(&c);
+    id<MTLTexture> ta = tex_mtl(&a);
+    CHECK(ta && shared(ta) && held(ta), "a live object's shared texture is counted as held");
+    id<MTLTexture> tb = tex_mtl(&b);
+    CHECK(tb && tb != ta && shared(ta) && shared(tb), "another description over the storage leaves the texture a live object resolves to");
+    mtl_tex_release(&a);
+    CHECK(!held(ta) && shared(ta), "the finalized object's texture is no longer held, and stays shared until its storage is taken");
+    for (int i = 0; i <= ALIAS_IDLE; i++) next_submission();
+    id<MTLTexture> tc = tex_mtl(&c);
+    CHECK(tc && !shared(ta), "a texture nothing resolves to goes when another description takes its storage");
+    CHECK(shared(tb) && shared(tc), "the live objects' textures stay");
+
+    // Two serials naming one object: held until both go
+    id o = [NSObject new]; uint64_t s1 = (1ull << 62) + 1, s2 = s1 + 1;
+    tex_put(s1, o); tex_put(s2, o);
+    serial_del(s1);
+    CHECK(held(o), "an object two serials name stays held when one of them goes");
+    serial_del(s2);
+    CHECK(!held(o), "and is not held once both have gone");
+    tex_put(s1, o); serial_put(s1, tb);
+    CHECK(!held(o) && held(tb), "replacing a serial's object moves the count");
+    serial_del(s1);
+
+    {   // The cost: 1,000 held entries over the storage of a 1024 x 1024 texture, 20,000 serials in the table
+        enum { NS = 20000, NA = 1000 };
+        NSMutableArray *keep = [NSMutableArray array];
+        pthread_mutex_lock(&g_smu);
+        for (int i = 0; i < NA; i++) {
+            id e = [NSObject new]; [keep addObject:e];
+            alias_key k; memset(&k, 0, sizeof k); k.off = (int64_t)i * 4096; k.d = tex(0, NVN_FORMAT_RGBA8, 32).d;
+            hm_put_id(&g_alias, &k, e);
+        }
+        pthread_mutex_unlock(&g_smu);
+        for (int i = 0; i < NS; i++) tex_put((1ull << 62) + 100 + i, i < NA ? keep[i] : [NSObject new]);
+        tex_t big = tex(91010, NVN_FORMAT_RGBA8, 1024); alias_key nk; memset(&nk, 0, sizeof nk); nk.d = big.d;
+        pthread_mutex_lock(&g_smu);
+        uint32_t before = g_alias.n; uint64_t t0 = bp_now();
+        alias_take_over(&nk);
+        double ms = (bp_now() - t0) / 1e6; uint32_t after = g_alias.n;
+        pthread_mutex_unlock(&g_smu);
+        printf("a takeover past %d held entries with %u serials in the table: %.3f ms\n", NA, g_by_serial.n, ms);
+        CHECK(after == before, "held entries are kept");
+        CHECK(ms < 1, "the takeover asks each entry in one lookup (under 1 ms; walking the serial table for each took 4.3 ms on an M5 Max)");
+        for (int i = 0; i < NS; i++) serial_del((1ull << 62) + 100 + i);
+    }
+    mtl_tex_release(&b); mtl_tex_release(&c);
+    next_submission(); defer_drain();
+    printf("t_alias_takeover: %d checks, %d failed\n", checks, failed); return failed != 0;
+} }
diff --git a/dm/gfx/tests/t_int_float_view.m b/dm/gfx/tests/t_int_float_view.m
index d9f8e08249..137d1309de 100644
--- a/dm/gfx/tests/t_int_float_view.m
+++ b/dm/gfx/tests/t_int_float_view.m
@@ -26,13 +26,15 @@ int main(void) { @autoreleasepool {
         "vertex float4 vs(uint id [[vertex_id]]) { return float4(id == 1 ? 3.0 : -1.0, id == 2 ? 3.0 : -1.0, 0.0, 1.0); }\n"
         "fragment uint fs(float4 p [[position]], texture2d<float> t [[texture(0)]]) { return as_type<uint>(t.read(uint2(p.xy), 0).x); }\n" options:nil error:&e];
     if (!lib) { printf("shader source: %s\n", e.localizedDescription.UTF8String); return 1; }
-    id<MTLComputePipelineState> cp = [dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"cs"] error:&e];
+    id<MTLFunction> csf = [lib newFunctionWithName:@"cs"];
+    id<MTLComputePipelineState> cp = csf ? [dev newComputePipelineStateWithFunction:csf error:&e] : nil;
     MTLRenderPipelineDescriptor *pd = [MTLRenderPipelineDescriptor new]; pd.vertexFunction = [lib newFunctionWithName:@"vs"]; pd.fragmentFunction = [lib newFunctionWithName:@"fs"]; pd.colorAttachments[0].pixelFormat = MTLPixelFormatR32Uint;
     id<MTLRenderPipelineState> rp = [dev newRenderPipelineStateWithDescriptor:pd error:&e];
     if (!cp || !rp) { printf("pipelines: %s\n", e.localizedDescription.UTF8String); return 1; }
     id<MTLBuffer> co = [dev newBufferWithLength:32 options:MTLResourceStorageModeShared], fo = [dev newBufferWithLength:32 options:MTLResourceStorageModeShared];
     MTLTextureDescriptor *rd = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatR32Uint width:8 height:1 mipmapped:NO]; rd.usage = MTLTextureUsageRenderTarget; rd.storageMode = MTLStorageModePrivate;
     id<MTLTexture> rt = [dev newTextureWithDescriptor:rd];
+    if (!t || !v || !co || !fo || !rt) { printf("textures or buffers not made\n"); return 1; }
     id<MTLCommandBuffer> cb = [q commandBuffer];
     id<MTLComputeCommandEncoder> ce = [cb computeCommandEncoder]; [ce setComputePipelineState:cp]; [ce setTexture:v atIndex:0]; [ce setBuffer:co offset:0 atIndex:0];
     [ce dispatchThreads:MTLSizeMake(8, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)]; [ce endEncoding];
diff --git a/docs/agents/symptoms.md b/docs/agents/symptoms.md
index 22c799e454..179b967d99 100644
--- a/docs/agents/symptoms.md
+++ b/docs/agents/symptoms.md
@@ -573,6 +573,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - ≈13 MB/frame leak. → [Transient render targets shared by storage](../reference/rendering.md#transient-render-targets-shared-by-storage)
 - a model has correct geometry, base colour and normals but appears as a black silhouette; deferred lighting's material-ID depth texture stays zero. → [IMAGE usage does not split texture storage](../reference/rendering.md#image-usage-does-not-split-texture-storage)
 - GPU memory grows with every image shown; the memory line's `shared textures` count keeps rising. → [Shared-storage entries taken over by a new texture](../reference/rendering.md#shared-storage-entries-taken-over-by-a-new-texture)
+- while the game makes many textures (Ocarina of Time's cutscenes), frames of hundreds of milliseconds up to a second, all of it in `submit+present` (the `[longframe]` lines); `sample` of the process names `alias_take_over` the submit thread's hottest function. → [A storage takeover walked the serial table for each entry](../reference/rendering.md#a-storage-takeover-walked-the-serial-table-for-each-entry)
 - GPU memory grows as resolution changes or rooms load, although almost every retained texture still has a live NVN descriptor; reclaiming unused sharing-table entries cannot recover it. → [Live resolution variants share physical image storage](../reference/rendering.md#live-resolution-variants-share-physical-image-storage)
 - the whole picture came out darker and more saturated on the Mac and the iPhone: comparing with a Switch screenshot, every channel at every level equalled the Switch value decoded from sRGB once more (0.55 → 0.26, 0.75 → 0.52). → [Render-target views that change the format](../reference/rendering.md#render-target-views-that-change-the-format)
 - a movie (or another picture the game converts each frame) does not show: the screen keeps the last 3D frame while the decoder delivers every picture. A pass dump shows a pass writing one serial and the composite sampling another serial of the same size, the two Metal formats `RGBA8Unorm` (70) and `RGBA8Unorm_sRGB` (71). → [Textures of an sRGB / linear pair over the same storage](../reference/rendering.md#textures-of-an-srgb--linear-pair-over-the-same-storage)
diff --git a/docs/reference/rendering.md b/docs/reference/rendering.md
index 93a9bbf5c3..2f2720b1dc 100644
--- a/docs/reference/rendering.md
+++ b/docs/reference/rendering.md
@@ -310,6 +310,15 @@ The replay on Metal (`dm/gfx/nvn/nvn_mtl.m`) and the Vulkan and OpenGL layers. E
 - **Found on:** [STEINS;GATE RE:BOOT](../../dm/games/01008BC024FDE000/README.md) (ASTC 4x4 images up to 4096 wide, RGBA8 render targets, R8 masks): from a quick save, 297 entries and 1.29 GB of GPU memory at frame 6,600; 468 entries and 1.58 GB 24,000 frames into a new game. With the fix: 97 entries and 0.93 GB at frame 6,600; 15 minutes of story from the quick save (53,700 frames) stayed at 1.08–1.15 GB of GPU memory and 1.02–1.13 GB resident, and a new game from boot stays at 0.97–1.03 GB from its first story scene to frame 27,300 (with the 16-submission guard).
 - **Check:** Captain Toad: Treasure Tracker's opening course, rendered with agl's transient targets, keeps its 26 entries and takes none over, with the same memory and frame times as before.
 
+### A storage takeover walked the serial table for each entry
+
+- **Symptom:** while the game makes many textures (Ocarina of Time's cutscenes), frames of hundreds of milliseconds up to a second, all of it in `submit+present` (the `[longframe]` lines); `sample` of the process names `alias_take_over` the submit thread's hottest function.
+- **Cause:** for each entry a takeover could remove, `alias_take_over` asked whether a live texture object still resolves to it through `serial_holds(&g_by_serial, …)`, a walk over the serial table's whole capacity. With about 7,500 serials and thousands of shared textures, every new shared texture walked the table once per entry of its pool that overlapped its storage, and the texture pool's walks (`pool_gpu`, which makes a texture for every registered ID of the table's type) make hundreds of textures in one submission.
+- **Fix:** general, in `nvn_mtl.m`: a count per object of the serial-table entries holding it (`g_by_serial_held`), kept by every writer of the serial table (`serial_put`, `serial_del`, `tex_put` and the upgrade's in-place replacement); `alias_take_over` asks it in one lookup.
+- **Check:** `dm/gfx/tests/t_alias_takeover.m`: an entry a live object resolves to stays and one nothing resolves to goes after `ALIAS_IDLE` submissions, two serials naming one object keep it held until both go, and a takeover past 1,000 held entries with 20,000 serials in the table takes 0.025 ms instead of 4.3 ms (M5 Max).
+- **Found on:** The Legend of Zelda: Ocarina of Time (Switch 2, `0400486024DEC000`, update v131072, with PR #1299), Saria's cutscene in Kokiri Forest, M5 Max. In 60 s `sample` captures of the submit thread over the cutscene, `alias_take_over` took 2,277 of the 3,711 samples spent making textures (`tex_mtl_`), 61 %; with the fix 38 of 618, 6 %. Game runs on the wall clock (the game does not run under `DM_DETCLOCK`) put the cutscene's heaviest part in or out of the run, so frame counts moved within their noise.
+- **Open:** the same captures and the memory line show the textures themselves being made again and again: 84,000 to 103,000 storage takeovers in about 200 s of play with 1,400 to 1,900 shared textures alive.
+
 ### Live resolution variants share physical image storage
 
 - **Symptom:** GPU memory grows as resolution changes or rooms load, although almost every retained texture still has a live NVN descriptor; reclaiming unused sharing-table entries cannot recover it.