From 05ef7bd08bb670f133197c670ff1928b0dfff756 Mon Sep 17 00:00:00 2001 From: baromas Date: Sat, 7 Nov 2026 22:53:08 +0300 Subject: [PATCH] perf: bypass matched OoT movie reconstruction in diagnostic mode Port the guarded bilinear movie path to the existing writable-image and direct dispatch interfaces of the MD5Visual branch. Preserve the original reconstruction by default. The synthetic probe passed 72 checks with Metal API validation. --- THIRD_PARTY_NOTICES.md | 2 + games/totk/Makefile | 6 +- games/totk/runtime/gpu_path_probe.m | 3 +- games/totk/runtime/loader.c | 2 +- .../third_party/d2r_gfx/gfx/nvn/nvn_mtl.h | 1 + .../third_party/d2r_gfx/gfx/nvn/nvn_mtl.m | 24 ++- .../d2r_gfx/gfx/nvn/oot_movie_direct.inc | 161 ++++++++++++++++++ .../d2r_gfx/gfx/nvn/oot_movie_probe.inc | 119 +++++++++++++ 8 files changed, 313 insertions(+), 5 deletions(-) create mode 100644 games/totk/third_party/d2r_gfx/gfx/nvn/oot_movie_direct.inc create mode 100644 games/totk/third_party/d2r_gfx/gfx/nvn/oot_movie_probe.inc diff --git a/THIRD_PARTY_NOTICES.md b/THIRD_PARTY_NOTICES.md index 935f6d207..81369146e 100644 --- a/THIRD_PARTY_NOTICES.md +++ b/THIRD_PARTY_NOTICES.md @@ -36,3 +36,5 @@ Fetched or linked at build time (not in this repository): | [ALVR](https://github.com/alvr-org/ALVR) client v20.14.0, downloaded (SHA-256 checked) and installed on the headset by `zeldaVR/install_vr.sh` | shows the VR picture on the Meta Quest 3 | MIT | | [Noto Sans CJK SC](https://github.com/notofonts/noto-cjk) (Source Han Sans) DemiLight at `165c01b`, downloaded and checked by `windows_games/ra2/scripts/languages.py` | the Simplified Chinese glyphs it draws into the player's game font | SIL OFL 1.1 | | [OpenCC](https://github.com/BYVoid/OpenCC), run (not included) by `windows_games/ra2/scripts/languages.py` | Traditional to Simplified Chinese | Apache-2.0 | + +OoT movie diagnostic: command-sequence matching adapted from [nxnative PR #1331](https://github.com/seedds/nxnative/pull/1331), `76c20edba480db753ce7092cd2677f6c4a8daa15`, MIT (`games/totk/third_party/licenses/nxnative.txt`). diff --git a/games/totk/Makefile b/games/totk/Makefile index 2bd8c8798..c1b750db7 100644 --- a/games/totk/Makefile +++ b/games/totk/Makefile @@ -147,7 +147,7 @@ verify-boot: runtime -GFX_HEADERS := $(wildcard third_party/d2r_gfx/gfx/nvn/*.h third_party/d2r_gfx/gfx/nvn/include/*.h third_party/d2r_gfx/host/*.h zeldaVR/gfx/*.h zeldaVR/runtime/*.h) +GFX_HEADERS := third_party/d2r_gfx/gfx/nvn/oot_movie_direct.inc third_party/d2r_gfx/gfx/nvn/oot_movie_probe.inc $(wildcard third_party/d2r_gfx/gfx/nvn/*.h third_party/d2r_gfx/gfx/nvn/include/*.h third_party/d2r_gfx/host/*.h zeldaVR/gfx/*.h zeldaVR/runtime/*.h) build/gfx/%.o: % $(GFX_HEADERS) Makefile @mkdir -p "$(dir $@)" @@ -515,6 +515,10 @@ build/totk-native-oot: $(wildcard runtime/*.[ch] runtime/*.m runtime/*.S) $(GFX_ runtime-oot binary-oot: SHADER_LIST ?= ../oot/assets/shader_list.tsv binary-oot: build/totk-native-oot build/pso_manifest.tsv build/pso_title.tsv runtime-oot: binary-oot game-data-oot + +.PHONY: probe-gpu-oot-movie +probe-gpu-oot-movie: build/totk-native-oot + MTL_DEBUG_LAYER=1 MTL_DEBUG_LAYER_ERROR_MODE=assert TOTK_HEADLESS=1 TOTK_OOT_MOVIE_DIRECT=1 build/totk-native-oot rewritten/oot $@ # Production Opus-wrapper regression checks; no game assets or audio device needed. OPUS_TEST_WORK ?= /tmp/totk-native-opus-codec OPUS_PREFIX ?= /opt/homebrew/opt/opus diff --git a/games/totk/runtime/gpu_path_probe.m b/games/totk/runtime/gpu_path_probe.m index f048e3584..8c00409e0 100644 --- a/games/totk/runtime/gpu_path_probe.m +++ b/games/totk/runtime/gpu_path_probe.m @@ -7,8 +7,9 @@ #include void runtime_gpu_path_probe(const char *which) { - int checked = !strcmp(which, "probe-gpu-writeback") ? mtl_probe_writeback() : !strcmp(which, "probe-gpu-sampler-border") ? mtl_probe_sampler_border() : !strcmp(which, "probe-gpu-pool-retire") ? mtl_probe_pool_retire() : !strcmp(which, "probe-gpu-tex-slots") ? mtl_probe_tex_slots() : !strcmp(which, "probe-gpu-arg-epoch") ? mtl_probe_arg_epoch() : !strcmp(which, "probe-gpu-anomaly") ? mtl_probe_anomaly() : !strcmp(which, "probe-gpu-vertex-streams") ? mtl_probe_vertex_streams() : !strcmp(which, "probe-gpu-vin") ? mtl_probe_vertex_input_types() : !strcmp(which, "probe-gpu-fout") ? mtl_probe_fragment_output_types() : !strcmp(which, "probe-gpu-pipeline-archive") ? mtl_probe_pipeline_archive() : !strcmp(which, "probe-gpu-layer-merge") ? mtl_probe_layer_merge() : !strcmp(which, "probe-gpu-cube-merge") ? mtl_probe_cube_merge() : !strcmp(which, "probe-gpu-texture-retire") ? mtl_probe_texture_retire() : !strcmp(which, "probe-gpu-texels-region") ? mtl_probe_texels_region() : !strcmp(which, "probe-gpu-linear-rt") ? mtl_probe_linear_rt_upload() : !strcmp(which, "probe-gpu-storage-view") ? mtl_probe_storage_view() : !strcmp(which, "probe-gpu-compressed-alias") ? mtl_probe_compressed_alias() : !strcmp(which, "probe-gpu-upload-batch") ? mtl_probe_upload_batch() : !strcmp(which, "probe-gpu-ds-clear") ? mtl_probe_ds_clear() : !strcmp(which, "probe-gpu-ub-rename") ? mtl_probe_ub_rename() : !strcmp(which, "probe-gpu-ubo-snapshot") ? mtl_probe_ubo_snapshot() : !strcmp(which, "probe-gpu-drv-c0") ? mtl_probe_drv_c0() : !strcmp(which, "probe-gpu-copy-alias") ? mtl_probe_copy_alias() : !strcmp(which, "probe-gpu-dispatch-chain") ? mtl_probe_dispatch_chain() : !strcmp(which, "probe-gpu-srgb-view") ? mtl_probe_srgb_view() : !strcmp(which, "probe-gpu-depth-output") ? mtl_probe_depth_output() : !strcmp(which, "probe-gpu-widen") ? mtl_probe_widen() : !strcmp(which, "probe-gpu-midclear") ? mtl_probe_midclear() : !strcmp(which, "probe-gpu-winding") ? mtl_probe_winding() : !strcmp(which, "probe-gpu-tile") ? mtl_probe_tile() : mtl_probe_blockrt(); + int checked = !strcmp(which, "probe-gpu-oot-movie") ? mtl_probe_oot_movie() : !strcmp(which, "probe-gpu-writeback") ? mtl_probe_writeback() : !strcmp(which, "probe-gpu-sampler-border") ? mtl_probe_sampler_border() : !strcmp(which, "probe-gpu-pool-retire") ? mtl_probe_pool_retire() : !strcmp(which, "probe-gpu-tex-slots") ? mtl_probe_tex_slots() : !strcmp(which, "probe-gpu-arg-epoch") ? mtl_probe_arg_epoch() : !strcmp(which, "probe-gpu-anomaly") ? mtl_probe_anomaly() : !strcmp(which, "probe-gpu-vertex-streams") ? mtl_probe_vertex_streams() : !strcmp(which, "probe-gpu-vin") ? mtl_probe_vertex_input_types() : !strcmp(which, "probe-gpu-fout") ? mtl_probe_fragment_output_types() : !strcmp(which, "probe-gpu-pipeline-archive") ? mtl_probe_pipeline_archive() : !strcmp(which, "probe-gpu-layer-merge") ? mtl_probe_layer_merge() : !strcmp(which, "probe-gpu-cube-merge") ? mtl_probe_cube_merge() : !strcmp(which, "probe-gpu-texture-retire") ? mtl_probe_texture_retire() : !strcmp(which, "probe-gpu-texels-region") ? mtl_probe_texels_region() : !strcmp(which, "probe-gpu-linear-rt") ? mtl_probe_linear_rt_upload() : !strcmp(which, "probe-gpu-storage-view") ? mtl_probe_storage_view() : !strcmp(which, "probe-gpu-compressed-alias") ? mtl_probe_compressed_alias() : !strcmp(which, "probe-gpu-upload-batch") ? mtl_probe_upload_batch() : !strcmp(which, "probe-gpu-ds-clear") ? mtl_probe_ds_clear() : !strcmp(which, "probe-gpu-ub-rename") ? mtl_probe_ub_rename() : !strcmp(which, "probe-gpu-ubo-snapshot") ? mtl_probe_ubo_snapshot() : !strcmp(which, "probe-gpu-drv-c0") ? mtl_probe_drv_c0() : !strcmp(which, "probe-gpu-copy-alias") ? mtl_probe_copy_alias() : !strcmp(which, "probe-gpu-dispatch-chain") ? mtl_probe_dispatch_chain() : !strcmp(which, "probe-gpu-srgb-view") ? mtl_probe_srgb_view() : !strcmp(which, "probe-gpu-depth-output") ? mtl_probe_depth_output() : !strcmp(which, "probe-gpu-widen") ? mtl_probe_widen() : !strcmp(which, "probe-gpu-midclear") ? mtl_probe_midclear() : !strcmp(which, "probe-gpu-winding") ? mtl_probe_winding() : !strcmp(which, "probe-gpu-tile") ? mtl_probe_tile() : mtl_probe_blockrt(); printf("{\"status\":\"pass\",\"probe\":\"%s\",\"checked\":%d,\"scope\":\"%s\"}\n", which, checked, + !strcmp(which, "probe-gpu-oot-movie") ? "OoT matched movie reconstruction, fallback, RGBA/sRGB and output ownership" : !strcmp(which, "probe-gpu-writeback") ? "terrain exact snapshots and BotW tiny GPU render completion write-back, CPU overwrite protection and padding" : !strcmp(which, "probe-gpu-sampler-border") ? "production VS/FS sampler border colour switching, mixed addressing and GPU readback" : !strcmp(which, "probe-gpu-texels-region") ? "production texel GPU region read/write, pitch guards, packed/block formats, mip/layer/volume, overlap fallback and CPU/Metal mirror visibility" : diff --git a/games/totk/runtime/loader.c b/games/totk/runtime/loader.c index b9ad206e3..d064242e0 100644 --- a/games/totk/runtime/loader.c +++ b/games/totk/runtime/loader.c @@ -641,5 +641,5 @@ int main(int argc,char **argv) { #endif verify_identity(!strcmp(argv[2],"probe-identity-negative")); for(unsigned i=0;i seed-shaders (scripts/ns2_shader_corpus.py)");runtime_shader_seed(argv[3],argc>4?argv[4]:NULL);}else if(!strcmp(argv[2],"probe-swkbd"))runtime_swkbd_probe();else if(!strcmp(argv[2],"probe-save"))runtime_save_probe();else if(!strcmp(argv[2],"boot")){if(getenv("TOTK_VERIFY_STARTUP")||getenv("TOTK_HEADLESS"))boot();else runtime_front_run(boot);}else runtime_fail("unknown mode");return 0; + if(!strcmp(argv[2],"probe-gpu-oot-movie")||!strcmp(argv[2],"probe-gpu-writeback")||!strcmp(argv[2],"probe-gpu-sampler-border")||!strcmp(argv[2],"probe-gpu-pool-retire")||!strcmp(argv[2],"probe-gpu-tex-slots")||!strcmp(argv[2],"probe-gpu-arg-epoch")||!strcmp(argv[2],"probe-gpu-anomaly")||!strcmp(argv[2],"probe-gpu-vertex-streams")||!strcmp(argv[2],"probe-gpu-vin")||!strcmp(argv[2],"probe-gpu-fout")||!strcmp(argv[2],"probe-gpu-pipeline-archive")||!strcmp(argv[2],"probe-gpu-layer-merge")||!strcmp(argv[2],"probe-gpu-cube-merge")||!strcmp(argv[2],"probe-gpu-texture-retire")||!strcmp(argv[2],"probe-gpu-texels-region")||!strcmp(argv[2],"probe-gpu-linear-rt")||!strcmp(argv[2],"probe-gpu-storage-view")||!strcmp(argv[2],"probe-gpu-compressed-alias")||!strcmp(argv[2],"probe-gpu-upload-batch")||!strcmp(argv[2],"probe-gpu-ds-clear")||!strcmp(argv[2],"probe-gpu-ub-rename")||!strcmp(argv[2],"probe-gpu-ubo-snapshot")||!strcmp(argv[2],"probe-gpu-drv-c0")||!strcmp(argv[2],"probe-gpu-dispatch-chain")||!strcmp(argv[2],"probe-gpu-copy-alias")||!strcmp(argv[2],"probe-gpu-srgb-view")||!strcmp(argv[2],"probe-gpu-depth-output"))runtime_gpu_path_probe(argv[2]);else if(!strcmp(argv[2],"probe-memory"))runtime_memory_probe();else if(!strcmp(argv[2],"probe-gfx"))runtime_gfx_probe(0);else if(!strcmp(argv[2],"probe-vpool"))runtime_vpool_probe();else if(!strcmp(argv[2],"probe-gpu-widen")||!strcmp(argv[2],"probe-gpu-blockrt")||!strcmp(argv[2],"probe-gpu-midclear")||!strcmp(argv[2],"probe-gpu-winding")||!strcmp(argv[2],"probe-gpu-tile"))runtime_gpu_path_probe(argv[2]);else if(!strcmp(argv[2],"probe-audio"))runtime_audio_probe();else if(!strcmp(argv[2],"probe-gfx-negative"))runtime_gfx_probe(1);else if(!strcmp(argv[2],"probe"))probe();else if(!strcmp(argv[2],"probe-fs"))runtime_fs_probe();else if(!strcmp(argv[2],"seed-shaders")){if(argc<4)runtime_fail("seed-shaders needs the corpus: seed-shaders (scripts/ns2_shader_corpus.py)");runtime_shader_seed(argv[3],argc>4?argv[4]:NULL);}else if(!strcmp(argv[2],"probe-swkbd"))runtime_swkbd_probe();else if(!strcmp(argv[2],"probe-save"))runtime_save_probe();else if(!strcmp(argv[2],"boot")){if(getenv("TOTK_VERIFY_STARTUP")||getenv("TOTK_HEADLESS"))boot();else runtime_front_run(boot);}else runtime_fail("unknown mode");return 0; } diff --git a/games/totk/third_party/d2r_gfx/gfx/nvn/nvn_mtl.h b/games/totk/third_party/d2r_gfx/gfx/nvn/nvn_mtl.h index 02dd43427..61d33e038 100644 --- a/games/totk/third_party/d2r_gfx/gfx/nvn/nvn_mtl.h +++ b/games/totk/third_party/d2r_gfx/gfx/nvn/nvn_mtl.h @@ -56,6 +56,7 @@ int mtl_probe_srgb_view(void); int mtl_probe_storage_view(void); int mtl_probe_compressed_alias(void); int mtl_probe_linear_rt_upload(void); +int mtl_probe_oot_movie(void); int mtl_probe_texels_region(void); int mtl_probe_texture_retire(void); int mtl_probe_ds_clear(void); diff --git a/games/totk/third_party/d2r_gfx/gfx/nvn/nvn_mtl.m b/games/totk/third_party/d2r_gfx/gfx/nvn/nvn_mtl.m index 7ebd0348f..e4005f10d 100644 --- a/games/totk/third_party/d2r_gfx/gfx/nvn/nvn_mtl.m +++ b/games/totk/third_party/d2r_gfx/gfx/nvn/nvn_mtl.m @@ -1212,7 +1212,12 @@ static void b2_scene_shot(int eye) { // diagnostic: the eye's own last pass ta else MLOG("[vr] b2shot: this eye pass drew no presented texture"); } } void mtl_eye_capture(int eye) { g_b2_capture = eye; } +static id g_oot_movie_decoded; +static uint64_t g_oot_movie_sub; +static const uint8_t *g_oot_movie_skip_end; void mtl_sub_begin(void) { + g_oot_movie_decoded = nil; + g_oot_movie_skip_end = NULL; g_eye_region = 3; // stereo B2: a submission starts outside eye marks g_fc_begin0 = clock_gettime_nsec_np(CLOCK_UPTIME_RAW); /* The submission's autorelease pool, drained in mtl_sub_commit, opens before anything else: the guest's submit @@ -14934,7 +14939,8 @@ stages with device visibility (blit()), as the end / start barriers of two encod if (g_disp_n % 20000 == 0) MLOG("[nvn] 计算派发 至 %llu 次：其中 %llu 次（%.0f%%）接在已开的计算编码器里，没有另开编码器", (unsigned long long)g_disp_n, (unsigned long long)g_disp_shared, 100.0 * (double)g_disp_shared / (double)g_disp_n); return g_blit; } -static void dispatch(const p_dispatch *dp) { +#include "oot_movie_direct.inc" +static void dispatch_at(const p_dispatch *dp, const uint8_t *next) { zskip_clear(); const prog_t *cp = S.cprog ? S.cprog : S.prog; // the compute program (S.cprog); a probe that sets S.prog only, as before if (!cp) { RED_ONCE(27, "派发前没有绑定程序"); return; } @@ -14945,6 +14951,9 @@ static void dispatch(const p_dispatch *dp) { if (cpr->missing) { static uint64_t n; if (!n++) MLOG("[shader-miss] 计算着色器还没翻译：跳过这个派发（以后同类不再提示）"); return; } NvShader *cs = cpr->sh[5]; if (!cs) { RED_ONCE(28, "派发的程序没有计算着色器（或未映射）"); return; } + int movie = oot_movie_enabled() && dp->x > 0 && dp->y > 0 && dp->z == 1 ? oot_movie_role(cs->key) : 0; + if (movie == 1) g_oot_movie_decoded = nil; + if (movie == 2 && oot_movie_replace(cs)) return; /* TOTK_SKIP_CS=all | [,...] (diagnostic): those compute dispatches are not run (bisecting a GPU hang to its shader); TOTK_SKIP_CS_LOG=1 names each compute shader once as it is first dispatched. */ { static char list[16][17]; static int nl = -1, all, logk; static NSMutableSet *seen; @@ -15028,6 +15037,10 @@ first waits for the encode stage to finish that pass (rs_sync, ~5% of the replay if (atomic_load(&g_draw_log) > 0) MLOG("[uw] F%llu pass %llu dispatch cs %.16s grid %dx%dx%d", (unsigned long long)atomic_load(&g_cur_frame), g_stat_pass, cs->key, dp->x, dp->y, dp->z); enc_c_dispatch(ce, MTLSizeMake(dp->x, dp->y, dp->z), MTLSizeMake(tg[0], tg[1], tg[2])); if (!reuse) enc_c_end(ce); + if (movie == 1) { + oot_movie_remember(cs); + g_oot_movie_skip_end = oot_movie_try_network(next); + } bright_lut(cs); // 3D 画面加亮：游戏算完调色表之后重写 | 3D picture brightening: rewrite after the game finishes computing the color lookup table /* TOTK_EXPO_LOG= (diagnostic): after each dispatch of the exposure kernel (ae725063), copy floats 96..127 of its storage buffer 0 (statistics; exposure at 108..111) into a shared ring and write them with the frame number once @@ -15045,6 +15058,7 @@ first waits for the encode stage to finish that pass (rs_sync, ~5% of the replay } // GPU：贴图 → 缓冲（行距取 SetCopyRowStride，0 = 紧排） // GPU: texture -> buffer (row stride taken from SetCopyRowStride, 0 = tightly packed) +static void dispatch(const p_dispatch *dp) { dispatch_at(dp, NULL); } static void copy_t2b(const p_copy_b2t *c) { { static long on = -1; static int left = 24; if (on < 0) { const char *e = getenv("TOTK_T2T_LOG"); on = e ? atol(e) : 0; if (e && !on) on = 1; } // with TOTK_T2T_LOG: buffer <-> texture copies too if (on > 0 && (long)atomic_load(&g_cur_frame) >= on && left > 0 && c->tex) { left--; @@ -15766,7 +15780,11 @@ void mtl_exec(const uint8_t *p) { case OP_COPY_T2B: copy_t2b(a); break; case OP_COPY_T2T: copy_t2t(a); break; case OP_CLEAR_TEX: clear_tex(a); break; - case OP_DISPATCH: dispatch(a); break; + case OP_DISPATCH: + if (g_oot_movie_skip_end) { + if (p == g_oot_movie_skip_end) g_oot_movie_skip_end = NULL; + } else dispatch_at(a, p + h.size); + break; case OP_SET_TEX_POOL: { const dpool_t *tp = ((const p_set_pool *)a)->pool; extern int nvn_tex_pool_base(const void *pool); // hle_nvn.c: IDs past this pool's part of the table read 0 S.tex_pool_n = tp ? nvn_tex_pool_base(tp) + tp->n : 0; break; } case OP_BARRIER: case OP_SET_SAMP_POOL: break; @@ -18171,6 +18189,8 @@ int mtl_probe_compressed_alias(void) { /* CPU-visible linear-render-target textures contain ordinary pool data until the GPU writes them. Exercise both water-height formats, cold sampling, a CPU texel update, and subsequent GPU rendering. */ +#include "oot_movie_probe.inc" + int mtl_probe_linear_rt_upload(void) { mtl_init();int checked=0; for(int mode=0;mode<2;mode++) { diff --git a/games/totk/third_party/d2r_gfx/gfx/nvn/oot_movie_direct.inc b/games/totk/third_party/d2r_gfx/gfx/nvn/oot_movie_direct.inc new file mode 100644 index 000000000..68b630ab8 --- /dev/null +++ b/games/totk/third_party/d2r_gfx/gfx/nvn/oot_movie_direct.inc @@ -0,0 +1,161 @@ +/* Diagnostic only: isolate decoded movie output from NDVE reconstruction. + Based on the matching/lifetime strategy in nxnative PR #1331, 76c20edba480db753ce7092cd2677f6c4a8daa15 + (MIT, third_party/licenses/nxnative.txt). This renderer uses an ordered compute scale instead of a render copy. + This changes reconstruction quality, not a production correctness fix. */ +static const char oot_movie_convert_key[] = "49dff010b9832251bb18e000c71a1ace470fd79dcf444a581132dd9f347becf7"; +static const char oot_movie_final_key[] = "4f41ccf24f163e0dd5f7347190607a892a42575928fa6441245c43b7302d1440"; +static int oot_movie_enabled(void) { +#if defined(GAME_OOT) + static int enabled = -1; + if (enabled < 0) { const char *e = getenv("TOTK_OOT_MOVIE_DIRECT"); enabled = e && !strcmp(e, "1"); } + return enabled; +#else + return 0; +#endif +} +static int oot_movie_key_is(const char *key, const char *hash) { + return !strncmp(key, hash, 64) && (!key[64] || key[64] == '-'); +} +static int oot_movie_role(const char *key) { + if (oot_movie_key_is(key, oot_movie_convert_key)) return 1; + if (oot_movie_key_is(key, oot_movie_final_key)) return 2; + return 0; +} +static id oot_movie_image_handle(NvShader *cs, unsigned slot, uint64_t handle, uint64_t *serial, int write) { + int binding = 0; + for (int i = 0; i < cs->rf.nb; i++) { + const nvn_rbind *b = &cs->rf.b[i]; + if (b->kind == NVN_RB_IMG && b->reg == slot && b->rw && b->dtt == MTLTextureType2D) binding = 1; + } + if (!binding) return nil; + int tid = (int)(handle & 0xfffff); + id t = nil; + registry_lock(); + if (tid > 0 && tid < g_ntreg && g_treg[tid].ok) { + if (write) tex_ensure_writable(tid); + t = tex_by_id(tid); + if (serial) *serial = g_treg[tid].t.serial; + } + registry_unlock(); + if (t.textureType != MTLTextureType2D || t.sampleCount != 1 || + (t.pixelFormat != MTLPixelFormatRGBA8Unorm && t.pixelFormat != MTLPixelFormatRGBA8Unorm_sRGB) || + !swizzle_eq(t.swizzle, MTLTextureSwizzleChannelsDefault)) return nil; + return t; +} +static id oot_movie_image(NvShader *cs, unsigned slot, uint64_t *serial, int write) { + return oot_movie_image_handle(cs, slot, S.img[5][slot], serial, write); +} +static int oot_movie_scale(id in, id out) { + if (in.width != 1920 || in.height != 1080 || out.width != 2560 || out.height != 1440 || in == out || + in.textureType != MTLTextureType2D || out.textureType != MTLTextureType2D || + in.pixelFormat != MTLPixelFormatRGBA8Unorm || + (out.pixelFormat != MTLPixelFormatRGBA8Unorm && out.pixelFormat != MTLPixelFormatRGBA8Unorm_sRGB) || + !(in.usage & MTLTextureUsageShaderRead) || !(out.usage & MTLTextureUsageShaderWrite)) return 0; + static id ps; + static id table; + if (!ps || !table) { + NSError *error = nil; + NSString *source = @"#include \nusing namespace metal;\n" + "kernel void scale_movie(texture2d src [[texture(0)]], texture2d dst [[texture(1)]], uint2 p [[thread_position_in_grid]]) {" + " if(p.x>=dst.get_width()||p.y>=dst.get_height()) return;" + " constexpr sampler s(coord::normalized,address::clamp_to_edge,filter::linear);" + " float2 uv=(float2(p)+0.5)/float2(dst.get_width(),dst.get_height()); dst.write(src.sample(s,uv,level(0)),p); }"; + id lib = [g_dev newLibraryWithSource:source options:nil error:&error]; + id fn = lib ? [lib newFunctionWithName:@"scale_movie"] : nil; + ps = fn ? [g_dev newComputePipelineStateWithFunction:fn error:&error] : nil; + MTL4ArgumentTableDescriptor *ad = [MTL4ArgumentTableDescriptor new]; ad.maxTextureBindCount = 2; + table = [g_dev newArgumentTableWithDescriptor:ad error:&error]; + if (!ps || !table) { MLOG("[oot-movie] diagnostic preparation failed: %s", error.localizedDescription.UTF8String); return 0; } + } + rp_end(); blit_close(); + id ce = (id)enc_c_begin(NULL); + [ce setArgumentTable:table]; [table setTexture:in.gpuResourceID atIndex:0]; [table setTexture:out.gpuResourceID atIndex:1]; + enc_c_pso(ce, ps); enc_c_dispatch(ce, MTLSizeMake(320, 180, 1), MTLSizeMake(8, 8, 1)); enc_c_end(ce); + args_dirty_all(); return 1; +} +static void oot_movie_remember(NvShader *cs) { + g_oot_movie_decoded = oot_movie_image(cs, 0, NULL, 0); + if (g_oot_movie_decoded.width != 1920 || g_oot_movie_decoded.height != 1080 || + g_oot_movie_decoded.pixelFormat != MTLPixelFormatRGBA8Unorm) g_oot_movie_decoded = nil; + g_oot_movie_sub = g_sub_serial; +} +static const char *const oot_movie_network[] = { + "ba079f68ae288dbabe8eeaae96b3a4847c520f652a9e282c58c7662f4a8cefed", + "a8b4fb48fee0ccf87d23bd3dc0e92eca70e3fe11a3226c60de0c0846930cf30a", + "e2eebd7ba18fe35f0124f8ae33abc500260eeb514d01793fad49933bfb9f412b", + "e6be82d835fbe4d32e6d77fa7e653fdd80a7b5fa13049ff93e44db5fc93c163e", + "9442573ee4ace011f83321aaa40b4ded45cc5a07e267fa555de38ee70115c040", + "0a56002cdc7e1fcee7ee193bf042d2846ebccd24ca72457438c0c712c978943e", + "0a56002cdc7e1fcee7ee193bf042d2846ebccd24ca72457438c0c712c978943e", + "d15370a382630d4e7f27fa0eec59333cfa0b2751b7472b3436abbe58d4d9edce", + "e6be82d835fbe4d32e6d77fa7e653fdd80a7b5fa13049ff93e44db5fc93c163e", + "7f7ff0f36b54d90f71b63124eea41818cb9ffd387812d1e046d7d19a4b9230cc", + "a8b4fb48fee0ccf87d23bd3dc0e92eca70e3fe11a3226c60de0c0846930cf30a", + "5a7d69b067a7f011a4817b4357cffc2f50a2af35ea5c79bb3826aa722c431d11", + oot_movie_final_key, +}; + +/* Match the whole same-list sequence before changing GPU work. State and + barriers still replay normally through mtl_exec. */ +static const uint8_t *oot_movie_try_network(const uint8_t *p) { + if (!p || !g_oot_movie_decoded || g_oot_movie_sub != g_sub_serial) return NULL; + NvShader *cs = nil; + uint64_t dst = S.img[5][1]; + unsigned index = 0; + for (unsigned count = 0; count < 128; count++) { + op_hdr h; memcpy(&h, p, sizeof h); + if (h.size < sizeof h || (h.size & 7)) return NULL; + const void *a = p + sizeof h; + switch (h.op) { + case OP_BIND_PROGRAM: { + const p_bind_program *b = a; + if (h.size < sizeof h + sizeof *b || b->stages != NVN_SHADER_STAGE_COMPUTE_BIT || !b->prog) + return NULL; + cs = prog_get(b->prog)->sh[5]; + break; + } + case OP_BIND_IMG: { + const p_bind_handles *b = a; + if (h.size < sizeof h + offsetof(p_bind_handles, h) || b->stage != 5 || b->src || b->first < 0 || + b->first >= 32 || b->n <= 0 || b->n > 32 - b->first || + h.size < sizeof h + offsetof(p_bind_handles, h) + sizeof b->h[0] * (unsigned)b->n) + return NULL; + if (b->first <= 1 && b->first + b->n > 1) dst = b->h[1 - b->first]; + break; + } + case OP_BIND_SSBO: case OP_BIND_UBO: case OP_BIND_TEX: case OP_BIND_SAMP: + case OP_PUSH_CONST: case OP_BARRIER: break; + case OP_DISPATCH: { + const p_dispatch *d = a; + if (h.size < sizeof h + sizeof *d || d->x <= 0 || d->y <= 0 || d->z != 1 || + !cs || !oot_movie_key_is(cs->key, oot_movie_network[index])) + return NULL; + if (++index != sizeof oot_movie_network / sizeof *oot_movie_network) break; + uint64_t serial = 0; + id out = oot_movie_image_handle(cs, 1, dst, &serial, 1); + if (!oot_movie_scale(g_oot_movie_decoded, out)) return NULL; + gpu_owned_mark(serial); zskip_del(serial); + g_oot_movie_decoded = nil; + static unsigned replaced; + if (++replaced <= 3 || replaced % 600 == 0) + MLOG("[oot-movie] direct scale replaced %u matched reconstruction sequences", replaced); + return p; + } + default: return NULL; + } + p += h.size; + } + return NULL; +} +static int oot_movie_replace(NvShader *cs) { + id in = g_oot_movie_decoded; g_oot_movie_decoded = nil; + if (!in || g_oot_movie_sub != g_sub_serial) return 0; + rp_end(); uint64_t serial = 0; + id out = oot_movie_image(cs, 1, &serial, 1); + if (!oot_movie_scale(in, out)) return 0; + gpu_owned_mark(serial); zskip_del(serial); + static unsigned count; + if (++count <= 3 || count % 600 == 0) + MLOG("[oot-movie] EXPERIMENT: direct bilinear 1920x1080 -> 2560x1440; replaced %u final dispatches, output serial %llu", count, (unsigned long long)serial); + return 1; +} diff --git a/games/totk/third_party/d2r_gfx/gfx/nvn/oot_movie_probe.inc b/games/totk/third_party/d2r_gfx/gfx/nvn/oot_movie_probe.inc new file mode 100644 index 000000000..a142e67d9 --- /dev/null +++ b/games/totk/third_party/d2r_gfx/gfx/nvn/oot_movie_probe.inc @@ -0,0 +1,119 @@ +/* Synthetic GPU data only. Production dispatch matching, resource lifetime, fallback and bilinear scale. */ +int mtl_probe_oot_movie(void) { + mtl_init(); int checked = 0; + if (!oot_movie_enabled()) MDIE("run the OoT probe with TOTK_OOT_MOVIE_DIRECT=1"); + if (oot_movie_role("49dff010b9832251_b") || oot_movie_role("unknown") || + oot_movie_role(oot_movie_convert_key) != 1 || oot_movie_role(oot_movie_final_key) != 2) + MDIE("movie key matching"); + checked += 4; + NSError *error = nil; + id lib = [g_dev newLibraryWithSource:@"#include \nusing namespace metal;\nstruct A { texture2d out [[id(0)]]; }; kernel void fallback(constant A &a [[buffer(0)]],uint2 p [[thread_position_in_grid]]) {if(p.xrf.nb=1;convert->rf.b[0]=(nvn_rbind){.kind=NVN_RB_IMG,.reg=0,.rw=1,.dtt=MTLTextureType2D}; + final->rf.nb=1;final->rf.b[0]=(nvn_rbind){.kind=NVN_RB_IMG,.reg=1,.aid=0,.rw=1,.dtt=MTLTextureType2D}; + final->fn=[lib newFunctionWithName:@"fallback"];final->ae=[final->fn newArgumentEncoderWithBufferIndex:0];final->arg_len=final->ae.encodedLength; + final->rf.tg[0]=8;final->rf.tg[1]=8;final->rf.tg[2]=1;final->stage=5; + snprintf(final->key,sizeof final->key,"%s-probe_b",oot_movie_final_key); + NvProg *np=[NvProg new];np->sh[5]=final;prog_t program={.serial=0xfffffa00};serial_put(program.serial,np); + enum { IW=1920,IH=1080,OW=2560,OH=1440 }; + uint8_t *pixels=malloc(IW*IH*4);if(!pixels)MDIE("movie fixture allocation"); + for(int y=0;y rb=probe_buf(256*5);op_hdr end={OP_END,8}; + const unsigned xy[5][2]={{0,0},{OW-1,OH-1},{OW/2,OH/2},{17,29},{OW-31,OH-43}}; + for(int mode=0;mode<2;mode++) { + tex_t t[2]={{0}};pool_t pools[2]={{0}};void *mem[2];size_t size[2]; + for(int j=0;j<2;j++) { + t[j].serial=0xfffffa10+mode*4+j; + t[j].d=(nvn_texdesc){.width=j?OW:IW,.height=j?OH:IH,.depth=1,.levels=1,.target=1,.format=j&&mode?0x38:0x25, + .swizzle_r=2,.swizzle_g=3,.swizzle_b=4,.swizzle_a=5}; + nvn_tex_layout(&t[j].d,&t[j].L);size[j]=(t[j].L.size+getpagesize()-1)&~((size_t)getpagesize()-1); + mem[j]=mmap(NULL,size[j],PROT_READ|PROT_WRITE,MAP_PRIVATE|MAP_ANON,-1,0);if(mem[j]==MAP_FAILED)MDIE("movie fixture mmap"); + if(!j&&nvn_swizzle(&t[j].d,0,0,pixels,IW*4,mem[j]))MDIE("movie fixture swizzle"); + pools[j]=(pool_t){.mem=mem[j],.size=size[j],.flags=0x24};t[j].mem=mem[j];t[j].pool=&pools[j]; + mtl_pool_add(mem[j],size[j]);nvn_dirty_add_range(mem[j],size[j]);mtl_reg_tex(j+1,&t[j],NULL); + } + mtl_sub_begin();mtl_exec((const uint8_t *)&end);S.img[5][0]=1;S.img[5][1]=2;S.cprog=&program; + oot_movie_remember(convert);if(!g_oot_movie_decoded)MDIE("movie conversion binding"); + id in=g_oot_movie_decoded; + if(oot_movie_scale(in,in))MDIE("movie accepted wrong destination");checked++; + convert->rf.b[0].rw=0;if(oot_movie_image(convert,0,NULL,0))MDIE("movie accepted read-only conversion reflection");convert->rf.b[0].rw=1;checked++; + S.img[5][1]=0;if(oot_movie_replace(final))MDIE("movie accepted missing output");S.img[5][1]=2;checked++; + /* A stale image cannot survive into another submission. */ + oot_movie_remember(convert);mtl_sub_commit();mtl_finish();mtl_sub_begin();mtl_exec((const uint8_t *)&end); + if(g_oot_movie_decoded||oot_movie_replace(final))MDIE("movie stale submission accepted");checked++; + S.cprog=&program;S.img[5][0]=1;S.img[5][1]=2; + /* No conversion: the production dispatch must execute the magenta fixture kernel. */ + p_dispatch dp={OW/8,OH/8,1,0};dispatch(&dp); + registry_lock();id out=tex_by_id(2);registry_unlock(); + enc_b_copy_texture_buffer(blit(),out,0,0,MTLOriginMake(0,0,0),MTLSizeMake(1,1,1),rb,0,256,256); + mtl_sub_commit();mtl_finish();const uint8_t *b=rb.contents; + if(b[0]!=255||b[1]!=0||b[2]!=255||b[3]!=255)MDIE("movie fallback dispatch lost");checked+=4; + mtl_sub_begin();mtl_exec((const uint8_t *)&end);S.cprog=&program;S.img[5][0]=1;S.img[5][1]=2; + oot_movie_remember(convert); + if (!mode) { + dispatch(&dp); + } else { + enum { N = sizeof oot_movie_network / sizeof *oot_movie_network }; + uint8_t control[0x880] = {[0x714] = 5}; + prog_t programs[N - 1] = {0}; + struct { op_hdr bind; p_bind_program program; op_hdr run; p_dispatch dispatch; } ops[N]; + struct { uint8_t bytes[sizeof ops]; op_hdr end; } packet; + packet.end = end; + program.nstage = 1; + program.sd[0].control = control; + program.sd[0].data = (uintptr_t)control; + for (unsigned i = 0; i < N; ++i) { + prog_t *p = &program; + if (i != N - 1) { + p = &programs[i]; + *p = program; + p->serial = 0xfffffb00 + i; + NvShader *shader = [NvShader new]; + snprintf(shader->key, sizeof shader->key, "%s-probe_b", oot_movie_network[i]); + NvProg *resolved = [NvProg new]; + resolved->sh[5] = shader; + serial_put(p->serial, resolved); + } + ops[i].bind = (op_hdr){OP_BIND_PROGRAM, sizeof(op_hdr) + sizeof(p_bind_program)}; + ops[i].program = (p_bind_program){p, NVN_SHADER_STAGE_COMPUTE_BIT}; + ops[i].run = (op_hdr){OP_DISPATCH, sizeof(op_hdr) + sizeof(p_dispatch)}; + ops[i].dispatch = dp; + } + ops[1].bind.op = OP_FENCE_SYNC; + memcpy(packet.bytes, ops, sizeof ops); + if (oot_movie_try_network(packet.bytes)) MDIE("movie accepted a side-effect command"); + ops[1].bind.op = OP_BIND_PROGRAM; + memcpy(packet.bytes, ops, sizeof ops); + S.img[5][1] = 0; + if (oot_movie_try_network(packet.bytes)) MDIE("movie accepted a missing sequence output"); + S.img[5][1] = 2; + g_oot_movie_skip_end = oot_movie_try_network(packet.bytes); + if (g_oot_movie_skip_end != packet.bytes + (N - 1) * sizeof ops[0] + sizeof(op_hdr) + sizeof(p_bind_program)) + MDIE("movie sequence was not fully recognized"); + mtl_exec(packet.bytes); + if (g_oot_movie_skip_end) MDIE("movie sequence skip state leaked"); + checked += 4; + for (unsigned i = 0; i < N - 1; ++i) mtl_prog_release(&programs[i]); + } + if(g_oot_movie_decoded||!gpu_owned(t[1].serial)||S.cprog!=&program||S.img[5][1]!=2)MDIE("movie state/ownership");checked+=4; + registry_lock();out=tex_by_id(2);registry_unlock(); + for(int i=0;i<5;i++)enc_b_copy_texture_buffer(blit(),out,0,0,MTLOriginMake(xy[i][0],xy[i][1],0),MTLSizeMake(1,1,1),rb,i*256,256,256); + mtl_sub_commit();mtl_finish(); + for(int i=0;i<5;i++) { + float sx=fmaxf(0,fminf(IW-1,(xy[i][0]+.5f)*IW/OW-.5f)),sy=fmaxf(0,fminf(IH-1,(xy[i][1]+.5f)*IH/OH-.5f)); + unsigned x0=(unsigned)sx,y0=(unsigned)sy,x1=MIN(x0+1,IW-1),y1=MIN(y0+1,IH-1);float tx=sx-x0,ty=sy-y0; + for(int c=0;c<4;c++) { + float v=((1-ty)*((1-tx)*pixels[(y0*IW+x0)*4+c]+tx*pixels[(y0*IW+x1)*4+c])+ty*((1-tx)*pixels[(y1*IW+x0)*4+c]+tx*pixels[(y1*IW+x1)*4+c]))/255.f; + if(mode&&c<3)v=v<=.0031308f?v*12.92f:1.055f*powf(v,1/2.4f)-.055f; + int got=((uint8_t *)rb.contents)[i*256+c],want=lroundf(v*255); + if(abs(got-want)>2)MDIE("movie scale mode%d sample%d channel%d got%d want%d",mode,i,c,got,want);checked++; + } + } + S.cprog=NULL;S.img[5][0]=S.img[5][1]=0; + for(int j=0;j<2;j++){mtl_tex_release(&t[j]);mtl_pool_del(mem[j]);nvn_dirty_remove_range(mem[j]);mtl_finish();munmap(mem[j],size[j]);} + } + mtl_prog_release(&program);res_remove(rb);free(pixels);return checked; +}