TOTK: the opening cave's stone blocks get their torch light (CPU-filled textures read back from guest memory)

The blocks and steps kept their texture and ambient light but stayed black beside the torch, on the Mac and the iPhone. Their shader (54e8c08998050e79) takes roughness from a 5120x5120 BC4 atlas that was all zero; roughness 0 gives no point light. TOTK fills that atlas on the CPU: it fills 1024x1024 BC4 material textures from the CPU with no NVN flush, reads them back with nvnTextureReadTexels on its own threads before any draw uses them, and copies the blocks into the atlas. tex_dirty_upload only marked such a texture CPU-written when the submit thread first resolved it, so ReadTexels returned the never-uploaded host texture and the atlas got zeros.

- mtl_read_texels (every game): a texture whose content is in guest memory (packaged layout, CPU-written, or with tex_dirty_upload a registered texture written since registration) is read there, on any thread, without resolving or uploading the host texture.
- nvn_ft_poll (games with tex_dirty_upload, now TOTK): at each submission after guest memory changed, registered textures written after their first use become CPU-written and are uploaded at their next use; the atlas itself is filled after it was first bound. At most 128 MB of textures per look, going on where it stopped, through each entry's cached dirty-range handle (searching every pool per check cost 0.4 ms a frame). Textures made as render targets or shader-writable (mtl_tex_gpu_owned) leave the list without an upload.

Found by running MacNativeGames' build of the same game on the same save (their atlas held data and their blocks were lit), then reading the atlas memory inside their process. Mac A/B against the previous commit on the TOTK route, side by side: 30 fps, no drops to interval 3, submit CPU 4.46 vs 4.79 ms. Notes in docs/agents/04-graphics.md, with the still-open block-height rounding bug MacNativeGames reported.

dm/gfx/nvn/hle_nvn.c

@@ -277,7 +277,7 @@ NVN(nvnTextureBuilderGetStorageAlignment, size_t, const NVNtextureBuilder *b) {
 // textures with content 0, and the G-buffer's base colour came out black. Textures that can hold such data (not render targets, not linear, no packaged layout, in a pool
 // the CPU can write) are registered by address range; a flush that covers part of one marks it CPU-written (mtl_tex_cpu_written: uploaded from guest memory, re-uploaded
 // when its pages change). Off by default: in another game a flushed range could cover an image a shader writes, whose GPU content the upload would replace
-typedef struct { const uint8_t *mem; uint64_t size; tex_t t; uint64_t stamp; } flush_tex;   // stamp: dirty-page generation at registration (tex_dirty_upload)
+typedef struct { const uint8_t *mem; uint64_t size; tex_t t; uint64_t stamp; const void *span; } flush_tex;   // stamp: dirty-page generation at registration (tex_dirty_upload); span: its tracked range (nvn_dirty_span), so a check needs no search over every pool
 static flush_tex *g_ft; static size_t g_nft, g_cft; static pthread_mutex_t g_ft_mu = PTHREAD_MUTEX_INITIALIZER;
 static void ft_add(const tex_t *t) {
     // Buffer textures are views onto the pool buffer the GPU reads directly (tex_mtl in nvn_mtl.m); an upload would push Age of Calamity's 1,048,576-texel R32F one through tex_desc, past Metal's 2D limit
@@ -286,7 +286,7 @@ static void ft_add(const tex_t *t) {
     pthread_mutex_lock(&g_ft_mu);
     if (g_nft == g_cft && !(g_ft = realloc(g_ft, (g_cft = g_cft ? 2 * g_cft : 1024) * sizeof *g_ft))) die("flush_tex 分配失败");
     uint64_t st = nvn_dirty_now(); if (DM_GAME_TEX_DIRTY_UPLOAD) nvn_dirty_protect(t->mem, t->L.size);   // see nvn_ft_written
-    g_ft[g_nft++] = (flush_tex){t->mem, t->L.size, *t, st};
+    g_ft[g_nft++] = (flush_tex){t->mem, t->L.size, *t, st, nvn_dirty_span(t->mem, t->L.size)};
     pthread_mutex_unlock(&g_ft_mu);
 }
 static void ft_del(uint64_t serial) {
@@ -302,12 +302,38 @@ int nvn_ft_written(const tex_t *t) {
     if (!DM_GAME_TEX_DIRTY_UPLOAD || !t->serial) return 0;
     int r = 0; pthread_mutex_lock(&g_ft_mu);
     for (size_t i = 0; i < g_nft; i++) if (g_ft[i].t.serial == t->serial) {
-        if (nvn_dirty_since(g_ft[i].mem, g_ft[i].size, g_ft[i].stamp)) { r = 1; g_ft[i] = g_ft[--g_nft]; }
+        if (nvn_dirty_since_span(g_ft[i].span, g_ft[i].mem, g_ft[i].size, g_ft[i].stamp)) { r = 1; g_ft[i] = g_ft[--g_nft]; }
         break; }
     pthread_mutex_unlock(&g_ft_mu);
     if (r) mtl_tex_cpu_written(t);
     return r;
 }
+// Later writes (tex_dirty_upload): a texture first resolved before the CPU filled it stays a GPU-only texture, and the first-use check above never runs for it again.
+// Once per submission, when guest memory has changed since the last look, the registered textures whose pages were written since registration become CPU-written
+// (tex_mtl_ in nvn_mtl.m then drops the GPU-only texture and uploads it). Tears of the Kingdom fills its 5120x5120 BC4 roughness atlas from the CPU after first binding it,
+// as tiles stream in. A texture the GPU writes (render target, shader-written image) may share pages with data the CPU writes: it leaves the list without an upload, which would
+// replace the GPU's content
+// Each poll looks at no more than FT_POLL_BYTES of registered textures, going on where the last one stopped; any guest write starts another pass over the whole list.
+// Checking every texture at every submission cost TOTK 0.4 ms of submit CPU per frame on the Mac; this way a texture is picked up a few submissions later at most
+enum { FT_POLL_BYTES = 128 << 20 };
+void nvn_ft_poll(void) {
+    if (!DM_GAME_TEX_DIRTY_UPLOAD) return;
+    static uint64_t gen; static size_t cur, left;
+    static tex_t *hit; static size_t chit; size_t nhit = 0;
+    uint64_t now = nvn_dirty_now(), scanned = 0;
+    pthread_mutex_lock(&g_ft_mu);
+    if (now != gen) { gen = now; left = g_nft; }   // guest memory changed: every entry gets one more look
+    while (left && g_nft && scanned < FT_POLL_BYTES) {
+        if (cur >= g_nft) cur = 0;
+        scanned += g_ft[cur].size; left--;
+        if (nvn_dirty_since_span(g_ft[cur].span, g_ft[cur].mem, g_ft[cur].size, g_ft[cur].stamp)) {
+            if (nhit == chit && !(hit = realloc(hit, (chit = chit ? 2 * chit : 64) * sizeof *hit))) die("nvn_ft_poll 分配失败");
+            hit[nhit++] = g_ft[cur].t; g_ft[cur] = g_ft[--g_nft];
+        } else cur++;
+    }
+    pthread_mutex_unlock(&g_ft_mu);
+    for (size_t i = 0; i < nhit; i++) if (!mtl_tex_gpu_owned(hit[i].serial)) mtl_tex_cpu_written(&hit[i]);   // outside g_ft_mu: the ownership check takes the replay's serial table lock
+}
 static void ft_flush(const uint8_t *a, uint64_t n) {
     if (!DM_GAME_TEX_FLUSH_UPLOAD || !a || !n) return;
     pthread_mutex_lock(&g_ft_mu);
@@ -955,6 +981,7 @@ NVN(nvnQueueSubmitCommands, void, NVNqueue *q, int n, const NVNcommandHandle *h)
     uint64_t t0 = now_ns(), c0 = clock_gettime_nsec_np(CLOCK_THREAD_CPUTIME_ID), i0, y0, s0 = thr_sys_ns(); thr_counts(&i0, &y0);
     ((queue_t *)q)->submits++;
     mtl_sub_begin();
+    nvn_ft_poll();
     mtl_sub_index((const uint64_t *)h, n);
     for (int i = 0; i < n; i++) { mtl_exec_rest((const uint64_t *)h + i + 1, n - i - 1); mtl_exec((const uint8_t *)h[i]); }
     mtl_exec_rest(NULL, 0);

dm/gfx/nvn/nvn_mtl.m

@@ -1460,6 +1460,12 @@ static id<MTLTexture> tex_linear(const tex_t *t) {
 // CPU-written block-linear textures (WriteTexels): the game gave no packaged layout and changes the content from the CPU each frame (Engage's movie frames: a Y plane R8 + a UV plane RG8, WriteTexels every frame).
 // Only packaged textures used to be uploaded from guest memory, so such a texture became a GPU-only texture whose content stayed 0. Their serials are recorded and they resolve through the packaged path (upload + re-upload on dirty pages)
 static pthread_mutex_t g_cpu_tex_mu = PTHREAD_MUTEX_INITIALIZER; static CFMutableSetRef g_cpu_tex;
+// The GPU owns the texture's content: its host texture was made to be rendered into or written by shaders (nvn_ft_poll in hle_nvn.c never re-uploads such a texture from guest memory)
+int mtl_tex_gpu_owned(uint64_t serial) {
+    __unsafe_unretained id o = tex_get(serial);
+    if (!o || ![o conformsToProtocol:@protocol(MTLTexture)]) return 0;
+    return (((id<MTLTexture>)o).usage & (MTLTextureUsageRenderTarget | MTLTextureUsageShaderWrite)) != 0;
+}
 void mtl_tex_cpu_written(const tex_t *t) {
     uint64_t serial = t->serial;
     if (!serial) return;
@@ -5735,7 +5741,12 @@ void mtl_read_texels(const tex_t *t, unsigned level, unsigned layer0, const NVNc
     int bw, bh, bpb = nvn_fmt_bpb(t->d.format, &bw, &bh);
     size_t wb = (size_t)((r->width + bw - 1) / bw) * bpb, hb = (r->height + bh - 1) / bh, x0 = (size_t)(r->xoffset / bw) * bpb, y0 = r->yoffset / bh;
     int nz = r->depth ? r->depth : 1, is3d = t->d.target == 2;
-    if (t->has_layout && t->mem) {
+    // A texture the CPU filled has its content in guest memory: read it there, on whatever thread asks, without resolving (and maybe uploading) the host texture.
+    // That includes a tex_dirty_upload texture the CPU wrote since registration but the submit thread has not resolved yet: Tears of the Kingdom fills 1024x1024 BC4 material
+    // textures from the CPU, reads them back with ReadTexels on its own thread before any draw uses them, and copies the blocks into its 5120x5120 roughness atlas. Read from the
+    // never-uploaded host texture they were zeros, the atlas stayed empty, and the opening cave's stone blocks got roughness 0 and no torch light
+    int nvn_ft_written(const tex_t *);
+    if ((t->has_layout || tex_cpu_written(t->serial) || (DM_GAME_TEX_DIRTY_UPLOAD && nvn_ft_written(t))) && t->mem) {
         uint32_t lw = MAX(t->d.width >> level, 1u), lh = MAX(t->d.height >> level, 1u), lz = is3d ? MAX(t->d.depth >> level, 1u) : 1;
         size_t lp = (size_t)((lw + bw - 1) / bw) * bpb, lr = (lh + bh - 1) / bh;
         uint8_t *lin = malloc(lp * lr * lz);

dm/gfx/nvn/nvn_obj.h

@@ -180,6 +180,7 @@ void mtl_read_texels(const tex_t *t, unsigned level, unsigned layer0, const NVNc
 void mtl_reg_tex(int id, const tex_t *t, const view_t *v);
 // textures written by WriteTexels (recorded by serial): the replay layer treats them like packaged textures, uploaded from guest memory and re-uploaded on dirty pages
 void mtl_tex_cpu_written(const tex_t *t);
+int mtl_tex_gpu_owned(uint64_t serial);                                  // its host texture is a render target or shader-writable: GPU content, never re-uploaded from guest memory
 // A program's per-stage shader key (hle_nvn.c): the first 8 bytes of sha256(first 0x880 bytes of the control section (its own length when shorter) + code area). Computed at SetShaders (NVN lets the caller free the control section once it returns, and Unity does);
 // a stage whose code area was not yet in any memory pool then is computed the first time it is asked for. Returns 0 = no such stage
 int prog_key(const prog_t *p, unsigned stage, uint8_t key[8]);

docs/agents/04-graphics.md

@@ -171,11 +171,8 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog
 - **How the tile lighting was traced**: the G-buffer class byte selects per-class half-resolution passes through a depth buffer of class values (`gl_FragDepth = mix(0.5, 1, nibble-swapped class / 255)`, then full-screen or per-tile quads with depth `EQUAL`), so `DM_NVN_DEPTHOVR` (depth test forced to `ALWAYS`) with `HILITE` tells "rejected by depth" from "never rasterized". `TEXSNAP` copies used to run before the pass they were meant to show (classic-queue command buffer committed ahead of the Metal 4 one); they now go through the replay's copy encoder.
 - **Clears inside a pass, with a partial channel mask or with a partial scissor** were dropped (red 「通道中途 / 部分掩码 / 部分剪裁的清色未实现」, 「通道中途的剪裁深度模板清除未实现」). NVN clears obey the scissor and the clear's own channel or stencil write mask, and can come between draws, while a Metal load action clears a whole attachment at the start of a pass. Such clears are now drawn as one triangle over the pass in the open render encoder (`clear_color_tri` / `clear_ds_tri` in `nvn_mtl.m`, adapted from MacNativeGames, MIT). The pipeline has the pass's formats, the clear's mask as the write mask of its slot and no writes to the other slots. Depth is `ALWAYS` (and written for a depth clear), stencil `REPLACE` through the clear's write mask, the scissor is clamped as for draws, and visibility counting is off for the triangle. The value comes from the frame arena through the triangle's own argument tables, and the draws' tables are put back after it. A clear outside a pass opens the pass as a draw would. Clears a load action can express keep it, and two cases are still cleared whole, logged once as 「深度模板清除带剪裁或部分模板掩码，按整张清」: a depth / stencil clear with a scissor outside a pass, and a partial stencil write mask. TOTK issues the first (a depth clear scissored to 8×8 of a 64×64 target), and so does Wonder. With `render_cond` TOTK's scissored colour clears are all skipped; without it they are 50×50 tiles of a 500×400 target cleared to opaque black, about five a frame. Metroid Dread logged the red line once while loading on its smoke route.
 - **8-bit indices written earlier in the same submission** were read stale (red 「8 位索引缓冲本提交里有 GPU 侧写入：按当前内容展开」). Metal has no 8-bit index type, so replay widens 8-bit indices to 16-bit on the CPU. When a copy, update or fill encoded earlier in the submission covers them (`pend_hits`), the CPU reads the bytes from before that write. Those draws are now widened on the GPU (`wide_run`): a compute pass reads the indices from a 16-byte-aligned address, with the byte shift as a constant, and writes 16-bit ones into the frame arena. As in `ixc_run`, the first such draw widens every such draw after it up to the next buffer write outside a render pass, so only a draw met inside an open render pass with no widening yet would split it. In TOTK about one draw per widening, up to a dozen widenings a frame on the title and load screens, no pass split (log line `8 位索引 GPU 展开`). Draws whose indices no write covers keep the CPU widening.
-- **Open: the stone blocks and steps of the opening cave get no torch light (Mac and iPhone).** On the Switch they are grey stone lit warm on the side facing Zelda's torch; here they keep their stone texture and the ambient light but stay black beside the torch, while the cave walls next to them are lit. Fast loop: the in-cave autosave (`run/save` in the work folder, keys `f2500:A,f2800:A,f3100:A,f3600:A`, shot at frame 4500 with the torch on its stand right of Link; the scene timing varies, check the flame is in the shot). `DM_DETCLOCK` cannot be used: TOTK stops at frame 1, waiting for guest time that only moves on presents. What is known so far (2026-10-19):
-  - The blocks are drawn by fragment shader `54e8c08998050e79` (two draws in the terrain G-buffer pass). Colour target 3 holds the normal in red and green and, in blue, the normal's sign bits plus a roughness term (`(blue & 252) / 252`, read by the tile-light shaders such as `c23ef1f842fcfba5` in the pass that adds point lights into the RG11B10 light buffer). The blocks write roughness 0 there: they sample it from a 5120×5120 BC4 atlas (texture slot 25, `ser=86` in that run) that is all zero, unless bit 25 of their material uniform `c7[8].y` is set, which it is not.
-  - Nothing tracked writes that atlas: its guest memory holds 6 non-zero bytes, no copy, render target, image or texture over the same memory touches it. Filling it by hand (BC4 value 200) did not light the blocks, so the zero roughness is not the cause.
-  - Not yet ruled out: the torch's shadow map (casters such as the blocks fully in their own shadow), the light volumes' half-resolution depth test, and the blocks' normal map. MacNativeGames (github.com/ly05010419/MacNativeGames, MIT), a TOTK port built on an older copy of this replay layer, reports a texture-layout bug that is also ours: NVN chooses the automatic block height by truncating `height_in_blocks * 1.5 / 8`, while `nvn_tex.c` rounds the division up, so some sizes (a 160×160 BC3 texture: H3 instead of H2) are read with the wrong layout and show padding as black texels. The fix (`pick_tile(ah * 3 / 16)`) changes every game's layouts at those sizes and waits for a check across the profiles. The same project renders into BC textures through a block-sized auxiliary texture (here those draws are dropped: red 「渲染目标是压缩格式贴图」, a 32×32 BC4 target in TOTK) and gives compressible storage to the newest format when a texture of another format takes it.
-  - The diagnostic patch of this session (tagged `[DEBUG-c4]` logging, a later-write poll for `tex_dirty_upload`, a CPU-side read-back for CPU-written textures, the block-height change) is kept outside the repository in `_agent/power/totk_blocks_wip.patch` on the T7.
+- **Fixed: the stone blocks and steps of the opening cave got no torch light (Mac and iPhone).** On the Switch they are grey stone lit warm on the side facing the torch; here they kept their texture and the ambient light but stayed black. Their fragment shader `54e8c08998050e79` takes its roughness from a 5120×5120 BC4 atlas (roughness = the square root of its red, packed into the upper six bits of blue in colour target 3, which the tile-light shaders read); the atlas was all zero, roughness 0 gives no point light. TOTK fills the atlas on the CPU: it fills 1024×1024 BC4 material textures from the CPU with no NVN flush, reads them back with `nvnTextureReadTexels` on its own threads before any draw uses them, and copies the blocks into the atlas. `tex_dirty_upload` only marked such a texture CPU-written when the submit thread first resolved it, so `ReadTexels` returned the never-uploaded host texture: zeros, and the atlas got zeros. Now `mtl_read_texels` reads any texture whose content is in guest memory there (a packaged layout, a CPU-written texture, or a `tex_dirty_upload` texture written since registration), on any thread, without an upload; and `nvn_ft_poll` (`hle_nvn.c`, at each submission after guest memory changed, at most 128 MB of textures per look, going on where it stopped) marks registered textures written after their first use, because the atlas itself is filled after it was first bound. Textures the GPU renders into or writes (`mtl_tex_gpu_owned`) leave that list without an upload. Found by running MacNativeGames' build of the same game on the same save: their atlas held data and their blocks were lit (their `ReadTexels` reads guest memory). How it was traced: `DM_NVN_TEXDUMP` (the blocks' roughness bits in colour target 3), the shader's MSL, `DM_NVN_TEXRAW` on the atlas and `DM_NVN_TEXINFO` (guest bytes), then lldb reading the atlas memory inside their process. Fast loop: the in-cave autosave (`run/save`, keys `f2500:A,f2800:A,f3100:A,f3600:A`, shot at frame 4500); `DM_DETCLOCK` cannot be used (TOTK stops at frame 1 waiting for guest time).
+- **Texture layout (open):** MacNativeGames found that NVN chooses the automatic block height by truncating `height_in_blocks * 1.5 / 8`, while `nvn_tex.c` rounds the division up, so some sizes (a 160×160 BC3 texture: H3 instead of H2) are read with the wrong layout. The fix (`pick_tile(ah * 3 / 16)`) changes every game's layouts at those sizes and waits for a check across the profiles.
 
 ## 7. Diagnostic switches (host)