Remove the [DEBUG-flk1] instrumentation

- hle_nvn.c: drop the command memory callback / grant / caller
  logging, the overflow details, the +0x10 write log, the one-off
  NVN table dump, and the DM_DBG_GRANTS, DM_DBG_CALLSITES,
  DM_DBG_CMDMEM_SHRINK and DM_DBG_OLD_USED switches.
- nvn_mtl.m: drop the per-draw NaN probe (DM_DBG_NANPROBE), the
  colour target statistics (DM_DBG_TEXSTAT) and the PASSDUMP
  completion timing.
- Kept as regular diagnostics: the audio callback line every 2000
  callbacks (frames per callback, queued blocks, underruns) and
  DM_DBG_AU_FRAMES (docs/agents/03-bringup.md, crackling on iOS);
  DM_NVN_SHOT up to 4096 frames (the 05-debugging burst recipe).
- Mac 1.2.1 scripted run reaches the opening course with the same
  single red line as before.

dm/gfx/nvn/hle_nvn.c

@@ -1,4 +1,3 @@
-#include <dlfcn.h>
 // NVN 接口层：接管 main 导入的 nvnBootstrapLoader，sdk 里的 NVN 驱动整块不执行。
 // 523 个 NVN 名字一律返回计数跳板（nvn_thunks.S），跳板再转到这里的实现：
 //   · 有语义的（设备 / 内存池 / 贴图与采样器 / 队列 / 窗口 / 同步 / 事件 / 命令缓冲内存）按 NVN 语义实现，对象是真的；
@@ -420,9 +419,7 @@ NVN(nvnCommandBufferInitialize, NVNboolean, NVNcommandBuffer *c, NVNdevice *d) {
     memset(c, 0, sizeof(cmd_t)); CB(c)->magic = MAGIC('C', 'M', 'D', 'B'); CB(c)->dev = d; return 1;
 }
 NVN(nvnCommandBufferFinalize, void, NVNcommandBuffer *c) { CB(c)->magic = 0; }
-NVN(nvnCommandBufferSetMemoryCallback, void, NVNcommandBuffer *c, PFNNVNCOMMANDBUFFERMEMORYCALLBACKPROC cb) { CB(c)->cb = cb;
-    { static void *seen[16]; static int ns; int k = 0; while (k < ns && seen[k] != (void *)cb) k++;   // [DEBUG-flk1] 游戏的命令内存回调地址 | the game's command memory callback addresses
-      if (k == ns && ns < 16) { seen[ns++] = (void *)cb; module_t *m = module_of((void *)cb), *r = module_of(__builtin_return_address(0)); LOG("[DEBUG-flk1] SetMemoryCallback 回调 %s+%#lx（调用者 %s+%#lx）", m ? m->name : "?", (unsigned long)(m ? (uintptr_t)cb - (uintptr_t)m->base : (uintptr_t)cb), r ? r->name : "?", (unsigned long)(r ? (uintptr_t)__builtin_return_address(0) - (uintptr_t)r->base : 0)); } } }
+NVN(nvnCommandBufferSetMemoryCallback, void, NVNcommandBuffer *c, PFNNVNCOMMANDBUFFERMEMORYCALLBACKPROC cb) { CB(c)->cb = cb; }
 NVN(nvnCommandBufferSetMemoryCallbackData, void, NVNcommandBuffer *c, void *d) { CB(c)->cb_data = d; }
 static uint8_t *align8(uint8_t *p) { return (uint8_t *)(((uintptr_t)p + 7) & ~(uintptr_t)7); }
 // 诊断 DM_NVN_CMDSTAT=1：每 300 帧记一行命令录制的用量——每帧录了多少字节 / 多少条、游戏每次给多大一段命令内存、每帧补内存回调几次、宿主补段几次、
@@ -509,16 +506,7 @@ NVN(nvnCommandBufferAddCommandMemory, void, NVNcommandBuffer *c, const NVNmemory
     const pool_t *pl = (const pool_t *)p;
     if (!pl || n < 64) { if (cs_on()) { pthread_mutex_lock(&g_cs_mu); g_cs.hostseg++; pthread_mutex_unlock(&g_cs_mu); } cmd_new_segment(CB(c), hostseg_take(CB(c)), HOSTSEG_SZ); return; }
     if (!pl->mem) { LOG("【红】AddCommandMemory 池没有 CPU 存储"); return; }
-    {   // [DEBUG-flk1] DM_DBG_GRANTS=1：按（调用者, 长度档）统计游戏给命令内存，每 600 帧打一次 | DM_DBG_GRANTS=1: count the game's command memory grants by (caller, size class), printed every 600 frames
-        static int on = -1; if (on < 0) on = getenv("DM_DBG_GRANTS") != NULL;
-        if (on) { static uintptr_t ra[64]; static size_t sz[64]; static uint64_t cnt[64]; static int nr; static uint64_t last_f; static pthread_mutex_t mu = PTHREAD_MUTEX_INITIALIZER;
-            uintptr_t r = (uintptr_t)__builtin_return_address(0); size_t cls = n >= 0x100000 ? (n + 0xfffff) & ~(size_t)0xfffff : n >= 0x1000 ? (n + 0xfff) & ~(size_t)0xfff : n;
-            pthread_mutex_lock(&mu); int k = 0; while (k < nr && !(ra[k] == r && sz[k] == cls)) k++; if (k == nr && nr < 64) { ra[nr] = r; sz[nr] = cls; nr++; } if (k < 64) cnt[k]++;
-            if (dm_nvn_frame_count() / 600 != last_f / 600) { last_f = dm_nvn_frame_count(); for (int i = 0; i < nr; i++) { module_t *m = module_of((void *)ra[i]); LOG("[DEBUG-flk1] 给命令内存 调用者 %s+%#lx 长度档 %#zx 累计 %llu 次（帧 %llu）", m ? m->name : "?", (unsigned long)(m ? ra[i] - (uintptr_t)m->base : ra[i]), sz[i], (unsigned long long)cnt[i], (unsigned long long)last_f); } }
-            pthread_mutex_unlock(&mu); } }
     hostseg_release((uintptr_t)(pl->mem + off), (uintptr_t)(pl->mem + off + n));
-    { static long shrink = -1; if (shrink < 0) { const char *e = getenv("DM_DBG_CMDMEM_SHRINK"); shrink = e ? atol(e) : 0; }   // [DEBUG-flk1] 把游戏给的每段命令内存截短，逼录制溢出到宿主补段（在 Mac 上复现 iPhone 的溢出） | cut every game command memory chunk short to force overflow into host segments (reproduces the iPhone overflow on the Mac)
-      if (shrink && n > (size_t)shrink + 256) n -= (size_t)shrink; }
     if (cs_on()) { pthread_mutex_lock(&g_cs_mu); g_cs.grants++; g_cs.grant_bytes += n; if (!g_cs.grant_min || n < g_cs.grant_min) g_cs.grant_min = n; if (n > g_cs.grant_max) g_cs.grant_max = n; pthread_mutex_unlock(&g_cs_mu); }
     { static int cblog = -1; if (cblog < 0) cblog = getenv("DM_NVN_CBLOG") != NULL; static uint64_t ng; if (cblog && (++ng <= 60 || ng % 500 == 0)) LOG("[nvn] AddCommandMemory #%llu cmd=%p 池 %p +%#lx 长 %#zx（录制中 %d）", (unsigned long long)ng, (void *)c, (const void *)pl, (long)off, n, CB(c)->recording); }
     cmd_new_segment(CB(c), pl->mem + off, n);
@@ -533,9 +521,6 @@ static uint8_t *cmd_alloc(cmd_t *o, uint32_t size) {
         // The game's callback cannot supply enough (newer-SDK games hand out small chunks sized for real-device commands, or disable the callback and manage memory themselves): continue recording in a host segment instead of aborting
         if (!o->cb || tries == 2) {
             static int once; if (!once++) LOG("[nvn] 命令内存：游戏回调补不够（cb=%p，需 %u 字节），改用宿主补段", (void *)o->cb, size + OP_RESERVE);
-            { static int nlog; if (dm_nvn_frame_count() > 600 && nlog < 30) { nlog++;   // [DEBUG-flk1] 局内溢出明细 | in-game overflow details
-                LOG("[DEBUG-flk1] 溢出 帧 %llu cmd=%p 本次录制已写 %llu 字节，本条要 %u（+留 %d），段内剩 %ld，段尾 %p",
-                    (unsigned long long)dm_nvn_frame_count(), (void *)o, (unsigned long long)o->used, size, OP_RESERVE, (long)(o->end - o->cur), (void *)o->end); } }
             cmd_new_segment(o, hostseg_take(o), HOSTSEG_SZ); break;
         }
         if (cs_on()) { pthread_mutex_lock(&g_cs_mu); g_cs.cbs++; pthread_mutex_unlock(&g_cs_mu); }
@@ -554,22 +539,12 @@ NVN(nvnCommandBufferBeginRecording, void, NVNcommandBuffer *c) {
 }
 NVN(nvnCommandBufferEndRecording, NVNcommandHandle, NVNcommandBuffer *c) {
     cmd_t *o = CB(c);
-    { static int seen; if (o->nvn_flag10 && !seen++) LOG("[DEBUG-flk1] 游戏写了 +0x10：%#llx（cmd=%p）", (unsigned long long)o->nvn_flag10, (void *)o); }   // [DEBUG-flk1]
     uint8_t *p = cmd_alloc(o, 8); put_hdr(p, OP_END, 8);
     o->recording = 0;
     return (NVNcommandHandle)o->begin;
 }
 NVN(nvnCommandBufferIsRecording, NVNboolean, const NVNcommandBuffer *c) { return ((const cmd_t *)c)->recording != 0; }
 NVN(nvnCommandBufferCopyCommands, void, NVNcommandBuffer *c, int n, const NVNcommandHandle *h) {
-    {   // [DEBUG-flk1] DM_DBG_CALLSITES=1：记下调用者（客体返回地址）与次数，每 600 帧打一次 | DM_DBG_CALLSITES=1: record callers (guest return addresses) and counts, printed every 600 frames
-        static int on = -1; if (on < 0) on = getenv("DM_DBG_CALLSITES") != NULL;
-        if (on) { static uintptr_t ra[64]; static uint64_t cnt[64]; static int nra; static pthread_mutex_t mu = PTHREAD_MUTEX_INITIALIZER; static uint64_t last_f;
-            uintptr_t r = (uintptr_t)__builtin_return_address(0); pthread_mutex_lock(&mu);
-            int k = 0; while (k < nra && ra[k] != r) k++; if (k == nra && nra < 64) ra[nra++] = r; if (k < 64) cnt[k] += n;
-            if (dm_nvn_frame_count() / 600 != last_f / 600) { last_f = dm_nvn_frame_count(); for (int i = 0; i < nra; i++) { module_t *m = module_of((void *)ra[i]); LOG("[DEBUG-flk1] CallCommands 调用者 %s+%#lx 累计 %llu 个句柄", m ? m->name : "?", (unsigned long)(m ? ra[i] - (uintptr_t)m->base : ra[i]), (unsigned long long)cnt[i]); } }
-            pthread_mutex_unlock(&mu);
-            { static int tab; if (!tab++) { module_t *m = module_of((void *)r);   // [DEBUG-flk1] 一次性：打出 main+0x35ad800..0x35adc00 的 NVN 函数指针表（解析成宿主符号） | one-off: print the NVN function pointer table at main+0x35ad800..0x35adc00 (resolved to host symbols)
-                if (m) for (uint64_t o = 0x35ad800; o < 0x35adc00; o += 8) { void *fp = *(void **)(m->base + o); Dl_info di; if (fp && dladdr(fp, &di) && di.dli_sname) LOG("[DEBUG-flk1] 表 main+%#llx → %s", (unsigned long long)o, di.dli_sname); } } } } }
     for (int i = 0; i < n; i++) { uint8_t *p = cmd_alloc(CB(c), 16); put_hdr(p, OP_CALL, 16); memcpy(p + 8, &h[i], 8); }
 }
 NVN(nvnCommandBufferSignalEvent, void, NVNcommandBuffer *c, const NVNevent *e, NVNeventSignalMode m, NVNeventSignalLocation loc, int flags, uint32_t v) {
@@ -1394,7 +1369,6 @@ NVN_NOPROTO(nvnWindowSetNumActiveTextures, void, NVNwindow *w, int n) { (void)w;
 // and on 0 (with control memory also 0) treats the recording as empty and marks the handle invalid (from main+0x4bb6c it clears bit 0 of the material entry +0x50); later draws no longer CALL it and inherit the previous object's vertex buffers / textures:
 // on the iPhone a batch of draws on odd frames of the opening course fetched vertices out of bounds, giving huge triangles and NaN (black band, black-edged glow, grey stripes). Without overflow the accumulated value equals cur - begin
 NVN(nvnCommandBufferGetCommandMemoryUsed, size_t, const NVNcommandBuffer *c) { const cmd_t *m = (const cmd_t *)c;
-    { static int old = -1; if (old < 0) old = getenv("DM_DBG_OLD_USED") != NULL; if (old) return m->begin && m->cur >= m->begin ? (size_t)(m->cur - m->begin) : 0; }   // [DEBUG-flk1] 旧算法（对照） | the old computation (for comparison)
     return m->begin ? (size_t)m->used : 0; }
 NVN(nvnCommandBufferGetControlMemoryUsed, size_t, const NVNcommandBuffer *c) { (void)c; return 0; }
 NVN_NOPROTO(nvnCommandBufferSetCommandMemoryCallbackEnabled, void, NVNcommandBuffer *c, uint32_t on) { (void)c; (void)on; }

dm/gfx/nvn/nvn_mtl.m

@@ -532,10 +532,9 @@ static inline void enc_r_blend_color(id<MTL4RenderCommandEncoder> e, float r, fl
 static inline void enc_r_visibility(id<MTL4RenderCommandEncoder> e, MTLVisibilityResultMode m, NSUInteger o) { [e setVisibilityResultMode:m offset:o]; }
 static inline void enc_r_vbuf(id<MTL4RenderCommandEncoder> e, NSUInteger i, uint64_t gpu) { (void)e; [g_vt setAddress:gpu atIndex:i]; }
 static inline void enc_r_args(id<MTL4RenderCommandEncoder> e, MTLRenderStages s, uint64_t gpu, NSUInteger i) { (void)e; [(s == MTLRenderStageVertex ? g_vt : g_ft) setAddress:gpu atIndex:i]; }
-static void np_after_draw(id<MTL4RenderCommandEncoder> e);   // [DEBUG-flk1] 每次绘制后的 NaN 探针 | NaN probe after each draw
-static inline void enc_r_draw(id<MTL4RenderCommandEncoder> e, MTLPrimitiveType p, NSUInteger v, NSUInteger n) { [e drawPrimitives:p vertexStart:v vertexCount:n]; np_after_draw(e); }
-static inline void enc_r_draw_instanced(id<MTL4RenderCommandEncoder> e, MTLPrimitiveType p, NSUInteger v, NSUInteger n, NSUInteger i, NSUInteger b) { [e drawPrimitives:p vertexStart:v vertexCount:n instanceCount:i baseInstance:b]; np_after_draw(e); }
-static inline void enc_r_draw_indexed(id<MTL4RenderCommandEncoder> e, MTLPrimitiveType p, NSUInteger n, MTLIndexType t, uint64_t gpu, NSUInteger len, NSUInteger i, NSInteger bv, NSUInteger bi) { [e drawIndexedPrimitives:p indexCount:n indexType:t indexBuffer:gpu indexBufferLength:len instanceCount:i baseVertex:bv baseInstance:bi]; np_after_draw(e); }
+static inline void enc_r_draw(id<MTL4RenderCommandEncoder> e, MTLPrimitiveType p, NSUInteger v, NSUInteger n) { [e drawPrimitives:p vertexStart:v vertexCount:n]; }
+static inline void enc_r_draw_instanced(id<MTL4RenderCommandEncoder> e, MTLPrimitiveType p, NSUInteger v, NSUInteger n, NSUInteger i, NSUInteger b) { [e drawPrimitives:p vertexStart:v vertexCount:n instanceCount:i baseInstance:b]; }
+static inline void enc_r_draw_indexed(id<MTL4RenderCommandEncoder> e, MTLPrimitiveType p, NSUInteger n, MTLIndexType t, uint64_t gpu, NSUInteger len, NSUInteger i, NSInteger bv, NSUInteger bi) { [e drawIndexedPrimitives:p indexCount:n indexType:t indexBuffer:gpu indexBufferLength:len instanceCount:i baseVertex:bv baseInstance:bi]; }
 static inline id<MTL4ComputeCommandEncoder> enc_c_begin(MTLComputePassDescriptor *d) { (void)d; id<MTL4ComputeCommandEncoder> e = [g_cb computeCommandEncoder]; [e barrierAfterQueueStages:MTLStageAll beforeStages:MTLStageAll visibilityOptions:MTL4VisibilityOptionDevice]; [e setArgumentTable:g_ct]; return e; }
 static inline id<MTLComputeCommandEncoder> enc_c_begin_concurrent(id<MTLCommandBuffer> cb) { return [cb computeCommandEncoderWithDispatchType:MTLDispatchTypeConcurrent]; }
 static inline void enc_c_end(id e) { m4_end(e); }
@@ -745,7 +744,6 @@ int mtl_gpu_span(uint64_t f, uint64_t *t0, uint64_t *t1) {
     return *t1 >= *t0 && *t1;
 }
 static void rp_flush(void);
-static int passdump(void);   // [DEBUG-flk1]
 static const char *cctx_snap(void);   // 这次提交里有什么（出错时打进红行；定义在 mtl_exec 之前） | What is in this submission (printed into the failure line on error; defined before mtl_exec)
 uint64_t mtl_sub_commit(void) {
     defer_fallback();   // 推迟合并通道：提交结束仍有未回放的 P 绘制（正常路径不会到这里） | Deferred merged pass: at the end of the submission there are still P draws not replayed (the normal path never gets here)
@@ -762,7 +760,6 @@ uint64_t mtl_sub_commit(void) {
     MTL4CommitOptions *o = [MTL4CommitOptions new];
     [o addFeedbackHandler:^(id<MTL4CommitFeedback> f) { (void)f; for (dispatch_block_t block in commit_blocks) block(); }];
     uint64_t gpu_f = atomic_load(&g_cur_frame); gpu_tag_begin(gpu_f);
-    int dbg_pd = passdump(); uint64_t dbg_t0 = gpu_ticks();   // [DEBUG-flk1] 窗口内的计数写回与提交→完成时长 | counter write-backs and commit→completion time inside the dump window
     NSString *cctx = @(cctx_snap());
     [o addFeedbackHandler:^(id<MTL4CommitFeedback> cb) {
         gpu_tag_done(gpu_f, cb.GPUStartTime, cb.GPUEndTime);
@@ -776,8 +773,6 @@ uint64_t mtl_sub_commit(void) {
             }
             else { uint64_t t = gpu_ticks(), rec[2] = {ops[i].type == 0 || ops[i].type == 0x10 ? t : 0, t}; memcpy((void *)ops[i].addr, rec, sizeof rec); }
         }
-        if (dbg_pd) { uint64_t t1 = gpu_ticks(); MLOG("[DEBUG-flk1] 完成 F%llu fence=%llu 提交→完成 %.2f ms GPU %.2f ms 写回 %d 项", (unsigned long long)gpu_f, (unsigned long long)fence, (t1 - dbg_t0) * 625.0 / 12 / 1e6, (cb.GPUEndTime - cb.GPUStartTime) * 1e3, n);
-            for (int i = 0; i < n; i++) if (ops[i].kind != 1) { const uint64_t *w = (const uint64_t *)ops[i].addr; MLOG("[DEBUG-flk1]   计数 F%llu kind=%d type=%#x dst=%#llx 值=%llu 时戳=%llu", (unsigned long long)gpu_f, ops[i].kind, ops[i].type, (unsigned long long)ops[i].addr, (unsigned long long)w[0], (unsigned long long)w[1]); } }
         free(ops);
         pthread_mutex_lock(&g_mu); if (fence > g_done) g_done = fence; pthread_cond_broadcast(&g_cv); pthread_mutex_unlock(&g_mu);
         defer_drain();
@@ -4747,7 +4742,6 @@ frz_plain:;
                S.has_vp ? S.vp[0][0] : 0, S.has_vp ? S.vp[0][1] : 0, S.has_vp ? S.vp[0][2] : (float)tw, S.has_vp ? S.vp[0][3] : (float)th);
         }
     }
-    { extern const char *g_np_vs, *g_np_fs; extern int g_np_cnt; g_np_vs = pr->sh[0] ? pr->sh[0]->key : "-"; g_np_fs = pr->sh[1] ? pr->sh[1]->key : "-"; g_np_cnt = dw ? dw->count : di->count; }   // [DEBUG-flk1]
     if (rt_range_on() && !rt_range_pr_ro(pr)) g_pass.range_external=1;
     rt_range_note_depth();
 #ifdef NVN_RT_RANGE_TEST
@@ -5587,13 +5581,11 @@ static void shot(id<MTLTexture> src, const char *path, int screen, const int cro
 // Diagnostic: before the first render pass of the next frame is opened, wait for the GPU to finish, then draw each one (the mip level used by that frame) by full-screen sampling into RGBA8 and save as PNG (float formats are clamped to [0,1])
 static long g_td_frame = -2; static const char *g_td_path; static NSMutableArray *g_td_list; static int g_td_done;
 static void td_dump(void);
-static int g_td_stat_left = -1;   // [DEBUG-flk1] DM_DBG_TEXSTAT=N：连续 N 帧只记统计行（不存图） | DM_DBG_TEXSTAT=N: log stats lines only (no images) for N consecutive frames
 void td_note(MTL4RenderPassDescriptor *pd, NSUInteger nc) {
-    if (g_td_frame == -2) { const char *e = getenv("DM_NVN_TEXDUMP"), *c = e ? strchr(e, ':') : NULL; g_td_frame = c ? atol(e) : -1; g_td_path = c ? c + 1 : NULL; g_td_list = [NSMutableArray array];
-        const char *s = getenv("DM_DBG_TEXSTAT"); g_td_stat_left = s ? atoi(s) : -1; }
+    if (g_td_frame == -2) { const char *e = getenv("DM_NVN_TEXDUMP"), *c = e ? strchr(e, ':') : NULL; g_td_frame = c ? atol(e) : -1; g_td_path = c ? c + 1 : NULL; g_td_list = [NSMutableArray array]; }
     if (g_td_frame < 0 || g_td_done) return;
     long f = (long)atomic_load(&g_cur_frame);
-    if (f > g_td_frame) { td_dump(); if (g_td_stat_left > 0) { g_td_frame = f; g_td_done = 0; } else return; }   // [DEBUG-flk1] 统计模式：接着记本帧 | stats mode: go on recording this frame
+    if (f > g_td_frame) { td_dump(); return; }
     if (f != g_td_frame) return;
     for (NSUInteger i = 0; i < nc && i < 8; i++) {
         id<MTLTexture> t = pd.colorAttachments[i].texture; if (!t) continue;
@@ -5603,109 +5595,10 @@ void td_note(MTL4RenderPassDescriptor *pd, NSUInteger nc) {
         if (!seen) [g_td_list addObject:@[t, @(ser), @(lv)]];
     }
 }
-// [DEBUG-flk1] DM_DBG_NANPROBE=1（配合 DM_NVN_TEXDUMP 的帧）：该帧每次绘制后在同一编码器里插一个全屏探针，用 framebuffer fetch 读颜色附件 0，
-// 在 4×4 网格上按位判 NaN（不受快速数学影响）并原子加到本次绘制的计数槽；不写附件（写掩码全关）、不测深度。帧末与绘制描述（通道、着色器键、顶点数）一起记日志：
-// 计数从 ~0 跳起的第一次绘制就是产生 NaN 的那一个
-// [DEBUG-flk1] DM_DBG_NANPROBE=1 (with the DM_NVN_TEXDUMP frame): after every draw in that frame, insert a full-screen probe into the same encoder that reads color attachment 0 through framebuffer fetch,
-// tests for NaN bitwise on a 4×4 grid (immune to fast math) and atomically adds to this draw's counter slot; it writes nothing (write mask off) and does no depth test. At the end of the frame the counts are logged
-// together with each draw's description (pass, shader keys, vertex count): the first draw where the count jumps from ~0 is the one producing NaN
-const char *g_np_vs = "-", *g_np_fs = "-"; int g_np_cnt;
-#define NP_MAX 4096
-static int g_np_on = -1, g_np_n; static long g_np_frame = -1; static id<MTLBuffer> g_np_buf; static id<MTL4ArgumentTable> g_np_at; static id<MTLDepthStencilState> g_np_dss;
-static NSMutableDictionary<NSString *, id<MTLRenderPipelineState>> *g_np_pso; static id<MTLFunction> g_np_vfn, g_np_ffn; static char (*g_np_desc)[112];
-static void np_after_draw(id<MTL4RenderCommandEncoder> e) {
-    if (g_np_on < 0) g_np_on = getenv("DM_DBG_NANPROBE") != NULL;
-    if (!g_np_on || g_td_frame < 0 || g_td_done || (long)atomic_load(&g_cur_frame) != g_td_frame || !e) return;
-    uint32_t f0 = g_pass.cfmt[0];
-    if (f0 != MTLPixelFormatRG11B10Float && f0 != MTLPixelFormatRGBA16Float && f0 != MTLPixelFormatRGBA32Float && f0 != MTLPixelFormatRG16Float && f0 != MTLPixelFormatR16Float) return;   // 只看能存 NaN 的格式 | only formats that can hold NaN
-    if (g_enc_vis) return;   // 遮挡计数开着：探针会算进游戏的结果，跳过 | occlusion counting on: the probe would count into the game's result, skip
-    if (!g_np_buf) {
-        g_np_buf = [g_dev newBufferWithLength:NP_MAX * 4 options:MTLResourceStorageModeShared]; res_add(g_np_buf); res_commit(); g_np_desc = calloc(NP_MAX, sizeof *g_np_desc);
-        MTL4ArgumentTableDescriptor *ad = [MTL4ArgumentTableDescriptor new]; ad.maxBufferBindCount = 1; NSError *ae = nil; g_np_at = [g_dev newArgumentTableWithDescriptor:ad error:&ae];
-        const char *src = "#include <metal_stdlib>\nusing namespace metal;\nstruct V { float4 p [[position]]; };\n"
-            "vertex V vs(uint id [[vertex_id]]) { float2 uv = float2((id << 1) & 2, id & 2); V o; o.p = float4(uv * float2(2, -2) + float2(-1, 1), 0, 1); return o; }\n"
-            "static bool nanb(float x) { uint u = as_type<uint>(x); return (u & 0x7f800000u) == 0x7f800000u && (u & 0x007fffffu) != 0u; }\n"
-            "fragment float4 fs(V in [[stage_in]], float4 c [[color(0)]], device atomic_uint *n [[buffer(0)]]) { uint2 q = uint2(in.p.xy);\n"
-            "  if (((q.x | q.y) & 3u) == 0u && (nanb(c.x) || nanb(c.y) || nanb(c.z) || nanb(c.w))) atomic_fetch_add_explicit(n, 1u, memory_order_relaxed); return c; }\n";
-        NSError *le = nil; id<MTLLibrary> lib = [g_dev newLibraryWithSource:@(src) options:nil error:&le];
-        if (!lib || !g_np_at) { MLOG("【红】[DEBUG-flk1] NaN 探针建不出：%s", le.localizedDescription.UTF8String); g_np_on = 0; return; }
-        g_np_vfn = [lib newFunctionWithName:@"vs"]; g_np_ffn = [lib newFunctionWithName:@"fs"]; g_np_pso = [NSMutableDictionary dictionary];
-        MTLDepthStencilDescriptor *dd = [MTLDepthStencilDescriptor new]; dd.depthCompareFunction = MTLCompareFunctionAlways; dd.depthWriteEnabled = NO; g_np_dss = [g_dev newDepthStencilStateWithDescriptor:dd];
-    }
-    long f = (long)atomic_load(&g_cur_frame);
-    if (f != g_np_frame) { g_np_frame = f; g_np_n = 0; memset(g_np_buf.contents, 0, NP_MAX * 4); }
-    if (g_np_n >= NP_MAX) return;
-    NSString *k = [NSString stringWithFormat:@"%u,%u,%u,%u,%u,%u,%u,%u|%u|%u|%u", g_pass.cfmt[0], g_pass.cfmt[1], g_pass.cfmt[2], g_pass.cfmt[3], g_pass.cfmt[4], g_pass.cfmt[5], g_pass.cfmt[6], g_pass.cfmt[7], g_pass.dfmt, g_pass.sfmt, g_pass.samples];
-    id<MTLRenderPipelineState> ps = g_np_pso[k];
-    if (!ps) {
-        MTLRenderPipelineDescriptor *pd = [MTLRenderPipelineDescriptor new]; pd.vertexFunction = g_np_vfn; pd.fragmentFunction = g_np_ffn; pd.rasterSampleCount = g_pass.samples ? g_pass.samples : 1;
-        for (int i = 0; i < 8; i++) if (g_pass.cfmt[i]) { pd.colorAttachments[i].pixelFormat = (MTLPixelFormat)g_pass.cfmt[i]; pd.colorAttachments[i].writeMask = MTLColorWriteMaskNone; }
-        if (g_pass.hd) pd.depthAttachmentPixelFormat = (MTLPixelFormat)g_pass.dfmt;
-        if (g_pass.hs) pd.stencilAttachmentPixelFormat = (MTLPixelFormat)g_pass.sfmt;
-        NSError *pe = nil; ps = [g_dev newRenderPipelineStateWithDescriptor:pd error:&pe];
-        if (!ps) { RED_ONCE(67, "[DEBUG-flk1] NaN 探针管线建不出（%s）：%s", k.UTF8String, pe.localizedDescription.UTF8String); return; }
-        g_np_pso[k] = ps;
-    }
-    int slot = g_np_n++;
-    snprintf(g_np_desc[slot], sizeof *g_np_desc, "p%llu vs=%.16s fs=%.16s n=%d fmt=%u %lux%lu", (unsigned long long)g_stat_pass, g_np_vs, g_np_fs, g_np_cnt, f0, (unsigned long)g_pass.tw, (unsigned long)g_pass.th);
-    [g_np_at setAddress:g_np_buf.gpuAddress + (uint64_t)slot * 4 atIndex:0];
-    [e setArgumentTable:g_np_at atStages:MTLRenderStageFragment];
-    [e setRenderPipelineState:ps]; [e setDepthStencilState:g_np_dss];
-    [e setViewport:(MTLViewport){0, 0, (double)g_pass.tw, (double)g_pass.th, 0, 1}]; [e setScissorRect:(MTLScissorRect){0, 0, g_pass.tw, g_pass.th}];
-    [e setCullMode:MTLCullModeNone]; [e setTriangleFillMode:MTLTriangleFillModeFill];
-    [e drawPrimitives:MTLPrimitiveTypeTriangle vertexStart:0 vertexCount:3];
-    [e setArgumentTable:g_ft atStages:MTLRenderStageFragment];   // 还原：游戏的片元参数表；管线 / 深度状态缓存作废，下次绘制重设（视口 / 剪裁 / 剔除等每次绘制本来就重设） | restore: the game's fragment table; invalidate the pipeline / depth caches so the next draw sets them (viewport / scissor / cull etc. are set on every draw anyway)
-    g_enc_pso = nil; g_enc_dss = nil;
-}
-static void np_report(long frame) {
-    if (!g_np_n || g_np_frame != frame) return;
-    const uint32_t *c = g_np_buf.contents; uint32_t prev = 0;
-    MLOG("[DEBUG-flk1] NaN 探针 F%ld 共 %d 次绘制（计数 = 4×4 网格上的 NaN 像素，×16 约为全屏）", frame, g_np_n);
-    for (int i = 0; i < g_np_n; i++) { if (c[i] != prev || i == 0 || i == g_np_n - 1) MLOG("[DEBUG-flk1] 探针 F%ld #%d %s NaN=%u%s", frame, i, g_np_desc[i], c[i], c[i] > prev + 64 ? "  ←跳" : ""); prev = c[i]; }
-    g_np_n = 0;
-}
-// [DEBUG-flk1] 颜色目标统计：各通道均值、NaN / Inf 个数、最大有限值、黑像素占比（RGB 都 < 1/255） | color target stats: per-channel mean, NaN / Inf counts, largest finite value, black share (RGB all < 1/255)
-static float td_f11(uint32_t v, int mb) { uint32_t e = (v >> mb) & 31, m = v & ((1u << mb) - 1); if (e == 31) return m ? NAN : INFINITY; return e ? ldexpf(1 + (float)m / (1 << mb), (int)e - 15) : ldexpf((float)m / (1 << mb), -14); }
-static void td_stat(id<MTLTexture> t, uint64_t ser, NSUInteger lv, long frame) {
-    MTLPixelFormat pf = t.pixelFormat; size_t bpp = pf == MTLPixelFormatRGBA16Float ? 8 : pf == MTLPixelFormatRGBA32Float ? 16 : pf == MTLPixelFormatR16Float || pf == MTLPixelFormatRG8Unorm ? 2 : pf == MTLPixelFormatR8Unorm ? 1 : 4;
-    int known = pf == MTLPixelFormatRGBA16Float || pf == MTLPixelFormatRG11B10Float || pf == MTLPixelFormatRGBA8Unorm || pf == MTLPixelFormatRGBA8Unorm_sRGB || pf == MTLPixelFormatBGRA8Unorm || pf == MTLPixelFormatBGRA8Unorm_sRGB
-             || pf == MTLPixelFormatR16Float || pf == MTLPixelFormatRG16Float || pf == MTLPixelFormatR32Float || pf == MTLPixelFormatRGBA32Float || pf == MTLPixelFormatR8Unorm;
-    NSUInteger w = MAX(t.width >> lv, 1u), h = MAX(t.height >> lv, 1u);
-    if (!known) { MLOG("[DEBUG-flk1] 统计 F%ld ser=%llu lv=%lu 格式 %lu %lux%lu（格式未解）", frame, ser, (unsigned long)lv, (unsigned long)pf, (unsigned long)w, (unsigned long)h); return; }
-    id<MTLBuffer> rb = [g_dev newBufferWithLength:w * h * bpp options:MTLResourceStorageModeShared];
-    id<MTLCommandBuffer> rc = q_cb(); id<MTLBlitCommandEncoder> bb = [rc blitCommandEncoder];
-    [bb copyFromTexture:t sourceSlice:0 sourceLevel:lv sourceOrigin:MTLOriginMake(0, 0, 0) sourceSize:MTLSizeMake(w, h, 1) toBuffer:rb destinationOffset:0 destinationBytesPerRow:w * bpp destinationBytesPerImage:w * h * bpp];
-    [bb endEncoding]; [rc commit]; [rc waitUntilCompleted];
-    const uint8_t *p = rb.contents; double sum[4] = {0}; uint64_t nan = 0, inf = 0, black = 0, n = (uint64_t)w * h; float mx = 0;
-    for (uint64_t i = 0; i < n; i++) {
-        float c[4] = {0, 0, 0, 1};
-        switch (pf) {
-            case MTLPixelFormatRGBA16Float: for (int k = 0; k < 4; k++) c[k] = (float)((const __fp16 *)p)[i * 4 + k]; break;
-            case MTLPixelFormatRGBA32Float: for (int k = 0; k < 4; k++) c[k] = ((const float *)p)[i * 4 + k]; break;
-            case MTLPixelFormatR16Float: c[0] = (float)((const __fp16 *)p)[i]; c[1] = c[2] = 0; break;
-            case MTLPixelFormatRG16Float: c[0] = (float)((const __fp16 *)p)[i * 2]; c[1] = (float)((const __fp16 *)p)[i * 2 + 1]; c[2] = 0; break;
-            case MTLPixelFormatR32Float: c[0] = ((const float *)p)[i]; c[1] = c[2] = 0; break;
-            case MTLPixelFormatRG11B10Float: { uint32_t v = ((const uint32_t *)p)[i]; c[0] = td_f11(v & 0x7ff, 6); c[1] = td_f11((v >> 11) & 0x7ff, 6); c[2] = td_f11(v >> 22, 5); break; }
-            case MTLPixelFormatR8Unorm: c[0] = p[i] / 255.f; c[1] = c[2] = 0; break;
-            default: for (int k = 0; k < 4; k++) c[k] = p[i * 4 + k] / 255.f; if (pf == MTLPixelFormatBGRA8Unorm || pf == MTLPixelFormatBGRA8Unorm_sRGB) { float r = c[0]; c[0] = c[2]; c[2] = r; } break;
-        }
-        int bad = 0;
-        for (int k = 0; k < 4; k++) { if (isnan(c[k])) { nan++; bad = 1; } else if (isinf(c[k])) { inf++; bad = 1; } else { sum[k] += c[k]; if (fabsf(c[k]) > mx) mx = fabsf(c[k]); } }
-        if (!bad && c[0] < 1 / 255.f && c[1] < 1 / 255.f && c[2] < 1 / 255.f) black++;
-    }
-    MLOG("[DEBUG-flk1] 统计 F%ld ser=%llu lv=%lu 格式 %lu %lux%lu 均值 %.4g,%.4g,%.4g,%.4g NaN %llu Inf %llu 最大 %.4g 黑 %.4f", frame, ser, (unsigned long)lv, (unsigned long)pf, (unsigned long)w, (unsigned long)h,
-         sum[0] / n, sum[1] / n, sum[2] / n, sum[3] / n, (unsigned long long)nan, (unsigned long long)inf, mx, (double)black / n);
-}
 static void td_dump(void) {
     g_td_done = 1;
     rp_end(); if (g_blit) { enc_b_end(g_blit); g_blit = nil; }
     mtl_finish_gpu();
-    if (g_td_stat_left > 0) {   // [DEBUG-flk1] 统计模式 | stats mode
-        MLOG("[DEBUG-flk1] 统计 帧 %ld 共 %lu 张颜色目标", g_td_frame, (unsigned long)g_td_list.count);
-        np_report(g_td_frame);
-        for (NSArray *e in g_td_list) { id<MTLTexture> t = e[0]; if (t.textureType == MTLTextureType2D) td_stat(t, [e[1] unsignedLongLongValue], [e[2] unsignedIntegerValue], g_td_frame); }
-        [g_td_list removeAllObjects]; g_td_stat_left--; return;
-    }
     static id<MTLRenderPipelineState> ps; static id<MTLSamplerState> sp;
     if (!ps) {
         static const char *src_msl =
@@ -5969,7 +5862,7 @@ static int present_now(const tex_t *t, double period, uint64_t frame, uint64_t t
     }
     // DM_NVN_SHOT=帧号[,帧号…]:路径（路径里的 %d 换成帧号）
     // DM_NVN_SHOT=frame number[,frame number...]:path (%d in the path is replaced with the frame number)
-    static long shots[4096]; static int nshot = -1; static const char *shot_path;   // [DEBUG-flk1]
+    static long shots[4096]; static int nshot = -1; static const char *shot_path;   // 最多 4096 帧（连拍一段帧给逐帧打分用） | up to 4096 frames (a burst over a frame range, for scoring frame by frame)
     if (nshot < 0) {
         nshot = 0; const char *e = getenv("DM_NVN_SHOT"), *c = e ? strchr(e, ':') : NULL;
         if (c) { shot_path = c + 1; for (const char *q = e; q < c && nshot < 4096; q = strchr(q, ',') ? strchr(q, ',') + 1 : c) shots[nshot++] = atol(q); }

dm/host/hle_audio.c

@@ -61,6 +61,8 @@ static void *event_notify(void *p) {
     return NULL;
 }
 
+// 欠载：回调要数据时待播队列已空的次数 / 补零字节（au_render 每 2000 次回调记一行） | underruns: times the pending queue was empty when the callback needed data / zero-filled bytes (au_render logs one line every 2000 callbacks)
+static _Atomic uint64_t g_au_under_n, g_au_under_b;
 // 从待播队列取 n 字节交给出声端；整块播完就挪到已释放并通知普通线程。
 // Take n bytes from the pending queue and hand them to the output side; once a whole block has played, move it to released and notify the ordinary thread.
 static void pull(aout *a, uint8_t *dst, uint64_t n) {
@@ -69,7 +71,7 @@ static void pull(aout *a, uint8_t *dst, uint64_t n) {
     while (n) {
         uint32_t qt = atomic_load_explicit(&a->q_tail, memory_order_acquire);
         if (qh == qt || atomic_load_explicit(&a->state, memory_order_acquire) != STATE_STARTED) {
-            { extern _Atomic uint64_t g_au_under_n, g_au_under_b; if (qh == qt && atomic_load_explicit(&a->state, memory_order_acquire) == STATE_STARTED) { atomic_fetch_add(&g_au_under_n, 1); atomic_fetch_add(&g_au_under_b, n); } }   // [DEBUG-flk1]
+            if (qh == qt && atomic_load_explicit(&a->state, memory_order_acquire) == STATE_STARTED) { atomic_fetch_add(&g_au_under_n, 1); atomic_fetch_add(&g_au_under_b, n); }
             memset(dst, 0, n); break; }
         gbuf *b = a->q[qh & (QMAX - 1)];
         uint64_t k = b->size - a->pos; if (k > n) k = n;
@@ -91,16 +93,15 @@ static void pull(aout *a, uint8_t *dst, uint64_t n) {
 
 static _Atomic uint64_t g_render_calls;   // 声卡回调累计（生命周期验证：恢复后须继续增长） | Cumulative sound card callbacks (lifecycle verification: must keep growing after resume)
 uint64_t dm_audio_render_calls(void) { return atomic_load_explicit(&g_render_calls, memory_order_relaxed); }
-_Atomic uint64_t g_au_under_n, g_au_under_b;   // [DEBUG-flk1] 欠载：回调要数据时待播队列已空的次数 / 补零字节 | underruns: times the pending queue was empty when the callback needed data / zero-filled bytes
 static OSStatus au_render(void *ctx, AudioUnitRenderActionFlags *f, const AudioTimeStamp *ts, UInt32 bus, UInt32 frames, AudioBufferList *bl) {
     (void)f; (void)ts; (void)bus; (void)frames;
     aout *a = ctx;
     uint64_t nc = atomic_fetch_add_explicit(&g_render_calls, 1, memory_order_relaxed) + 1;
-    {   // [DEBUG-flk1] 每 2000 次回调记一行：每次要多少帧、回调开始时待播多少块、欠载次数 | every 2000 callbacks log one line: frames requested per call, pending blocks at callback start, underruns
+    {   // 每 2000 次回调记一行：每次要多少帧、回调开始时待播多少块、欠载次数（换新设备时看这行；docs/agents/03-bringup.md「iOS 上爆音」） | every 2000 callbacks log one line: frames requested per call, pending blocks at callback start, underruns (check this line on a new device; docs/agents/03-bringup.md "Crackling on iOS only")
         static uint32_t fmin = ~0u, fmax; static uint64_t qsum, u0, ub0; uint32_t qd = atomic_load(&a->q_tail) - atomic_load(&a->q_head);
         if (frames < fmin) fmin = frames; if (frames > fmax) fmax = frames; qsum += qd;
         if (nc % 2000 == 0) { uint64_t un = atomic_load(&g_au_under_n), ubb = atomic_load(&g_au_under_b);
-            LOG("[DEBUG-flk1] 声卡回调 %llu 次：每次 %u–%u 帧，回调时待播均 %.1f 块；这 2000 次里欠载 %llu 次、补零 %.1f ms", (unsigned long long)nc, fmin, fmax, qsum / 2000.0,
+            LOG("audio 声卡回调 %llu 次：每次 %u–%u 帧，回调时待播均 %.1f 块；这 2000 次里欠载 %llu 次、补零 %.1f ms", (unsigned long long)nc, fmin, fmax, qsum / 2000.0,
                 (unsigned long long)(un - u0), (ubb - ub0) / (2.0 * a->ch) / (RATE / 1000.0)); fmin = ~0u; fmax = 0; qsum = 0; u0 = un; ub0 = ubb; } }
     pull(a, bl->mBuffers[0].mData, bl->mBuffers[0].mDataByteSize);
     return noErr;
@@ -121,9 +122,9 @@ static void au_open(aout *a) {
                                      .mChannelsPerFrame = (UInt32)a->ch, .mBitsPerChannel = 16};
     AURenderCallbackStruct cb = {au_render, a};
 #if !TARGET_OS_IPHONE
-    { const char *e = getenv("DM_DBG_AU_FRAMES"); if (e) {   // [DEBUG-flk1] Mac 上把声卡回调改成每次 N 帧（模拟 iOS 缺省约 1024 帧） | on the Mac, make the sound card callback ask N frames each time (emulating the iOS default of about 1024 frames)
+    { const char *e = getenv("DM_DBG_AU_FRAMES"); if (e) {   // 诊断：Mac 上把声卡回调改成每次 N 帧（1024 = iOS 缺省，复现爆音） | diagnostic: on the Mac, make the sound card callback ask N frames each time (1024 = the iOS default, reproduces the crackling)
         UInt32 fr = (UInt32)atoi(e); OSStatus st = AudioUnitSetProperty(a->au, kAudioDevicePropertyBufferFrameSize, kAudioUnitScope_Global, 0, &fr, sizeof fr);
-        LOG("[DEBUG-flk1] 声卡缓冲帧数设为 %u：状态 %d", fr, (int)st); } }
+        LOG("audio 声卡缓冲帧数设为 %u：状态 %d", fr, (int)st); } }
 #endif
     if (AudioUnitSetProperty(a->au, kAudioUnitProperty_StreamFormat, kAudioUnitScope_Input, 0, &f, sizeof f) ||
         AudioUnitSetProperty(a->au, kAudioUnitProperty_SetRenderCallback, kAudioUnitScope_Input, 0, &cb, sizeof cb) ||

docs/agents/03-bringup.md

@@ -79,7 +79,7 @@ Each entry: symptom → root cause → fix (where the fix lives). Check these fi
 - **The CPU processor drops the final mix.** `dsp::SendCommandBuffer`, `dsp::Wait` and `dsp::Signal` are stubs in the application build, and DeviceSink (command 0x13) is a no-op in `CommandListProcessor::Process`. Take the output after `Process` from the mix buffers the DeviceSink command names (layout in `hle_audren.c`).
 - **Generation needs the renderer's lock.** `AudioRenderSystem::Update` holds the SDK mutex at `AudioRenderSystem + 0x1f0`; `GenerateCommand` does not take it, so the host does (the auto-mode system thread does the same in `SendCommandToDsp`).
 - **At least one rendered frame between two updates.** An effect's "needs init" parameter state is set only in the update where the game flags it new; the next update copies the game's value over it. Mario updates twice in a row at start-up; with a real audio device the first frame came after both, and the light limiter later crashed on a null state pointer (`ApplyLightLimiterEffect`, `sdk+0x228494`). The `AudioRenderSystem::Update` hook renders a catch-up frame first when nothing was rendered since the last update. `DM_AUDREN_TRACE=1` logs every update with the number of frames rendered since the previous one.
-- **Crackling on iOS only** → RemoteIO asks for as many frames per callback as the audio session's I/O buffer duration, about 21 ms by default (1024 frames at 48 kHz), while the renderer keeps 4 blocks × 5 ms = 20 ms in flight. Every callback came up 1.3 ms short and was padded with zeros, about 47 gaps per second. Reproduced on the Mac by forcing 1024-frame callbacks (1999 underruns in 2000 callbacks). Fix: `AVAudioSession` `setPreferredIOBufferDuration:0.005` (and 48 kHz) before the game opens audio and again on return to the foreground (`front_ios_audio_session` in `front_ios.m`); the device then calls back with 240 frames and nothing underruns. Log the callback sizes and the queued blocks at each callback when checking a new device.
+- **Crackling on iOS only** → RemoteIO asks for as many frames per callback as the audio session's I/O buffer duration, about 21 ms by default (1024 frames at 48 kHz), while the renderer keeps 4 blocks × 5 ms = 20 ms in flight. Every callback came up 1.3 ms short and was padded with zeros, about 47 gaps per second. Reproduced on the Mac by forcing 1024-frame callbacks (`DM_DBG_AU_FRAMES=1024` with `DM_AUDIO=1`: 1999 underruns in 2000 callbacks). Fix: `AVAudioSession` `setPreferredIOBufferDuration:0.005` (and 48 kHz) before the game opens audio and again on return to the foreground (`front_ios_audio_session` in `front_ios.m`); the device then calls back with 240 frames and nothing underruns. When checking a new device, read the host's `audio 声卡回调` line, logged every 2000 callbacks: frames per callback, queued blocks, underruns.
 - **Checking audio without listening**: `DM_AUDIO_DUMP=file` writes the PCM the host plays (16-bit, 48 kHz, stereo). RMS per second, clipping share, spectral flatness (music ≈ 0.1–0.5, white noise ≈ 0.85) and the strongest FFT peaks (they should land on musical notes) tell real music from noise or silence.
 
 ### NVN object sizes (memory corruption)