Skip to content

shader: complete Ampere task and mesh translation - #1

Closed
biosed wants to merge 90 commits into
triasbrata:mainfrom
biosed:codex/native-mesh-complete
Closed

biosed wants to merge 90 commits into
triasbrata:mainfrom
biosed:codex/native-mesh-complete

Conversation

@biosed

@biosed biosed commented Nov 11, 2026

Copy link
Copy Markdown

Implements the remaining native Ampere task/mesh semantics without title data or runtime fallbacks.

  • implement F2IP byte packing and rounding modes
  • let task shaders read their own 0xb800 payload stores
  • carry point/line topology, point size, and topology-specific primitive indices into SPV_EXT_mesh_shader
  • preserve byte-identical output for all 514 previously successful Ocarina task/triangle-mesh modules

Validation:

  • test_dmsh_ampere_mesh.py: 11/11 passed
  • all 520 Ocarina task/mesh jobs translated and passed spirv-val --target-env vulkan1.3
  • 514 prior raw SPIR-V modules byte-identical; additions are five point meshes and task 5d3a096025e375a1

The game package, generated jobs, and translated modules remain outside this GPL source tree.

seedds and others added 30 commits October 31, 2026 09:51
Bind the current Capstone instruction in the operand helper and assert
that a detected engine-region gate has a size register before encoding it.
These make the existing invariants explicit without changing valid rewrite
behavior or widening the Python lint baseline.

Reviewed PR #756. Focused UE-region, lint-ratchet and docs tests pass
(28 tests); the pinned lint check has 392 baseline findings, zero new
findings and zero stale entries. Regression risk is low: instruction
matching and emitted gate encoding are unchanged, so no game regression
run was needed.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Declare the near-zero depth selected by Echoes of Wisdom's NVN calls so its
offline pipeline archive uses main0_z and matches the runtime stamp.

Compare observed runtime depth with each tested profile in NVN smoke runs,
including quick runs without an archive. Detect bundled archive stamp
rejections and prevent mismatched runs from recording a green validation.
Invalidate cached baselines that predate these observations.

Add regression coverage and document the mechanism and validation limits.
Both archives built all 7,495 bases with matched historical host/shader
artifacts: GL was rejected and zero passed 16/16 coverage probes on the
opening route. Castle-entry performance remains unvalidated without a save.

Validation: 77 targeted tests passed, one compiler-dependent check skipped;
Python lint and diff checks passed. Regression risk is low because only a
profile archive input and harness diagnostics change; no multi-game
regress.py check was run.

Refs #758.

Co-authored-by: Marceloalexandregoncalves <62747593+Marceloalexandregoncalves@users.noreply.github.com>
Guest ring allocators update NVNbuffer offset and size directly at +0x18
and +0x20. The host previously read +0x10 and +0x18, so GetAddress kept
the initial offset and GetSize returned the allocation offset.

Reserve the intervening bytes and assert the guest field offsets. All
buffer queries now use the updated fields without a game-specific rule.
Add public-entry-point tests for raw stores, ring wraparound, GPU-only
pool mapping restrictions, builder independence and neighbouring objects.

Fixes #723.
Source: #727.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Record depth_mode zero for SteamWorld Heist II, NieR:Automata, Portal,
DRAGON QUEST III HD-2D Remake, Raging Loop, Subnautica and Below Zero.
Their registered builds select near-zero NVN depth while the omitted key
made offline pipeline builds default to the GL vertex entry.

Confirm the recorded routes against registered module build IDs and
reproduce each game's actual NVN depth call with an isolated headless boot.
Extend the archive-builder regression test to check all eight affected
profiles' stamps and vertex entries, including Echoes of Wisdom. Record
the evidence and validation limits in each game note and the shared reference.

The original 14-profile audit now has no observed depth mismatch. The
shader-library input digests are unchanged; archives must be rebuilt and
apps repacked. Rebuilt archive acceptance and performance remain unmeasured.

Validation: seven isolated depth confirmations passed; 77 targeted tests
passed, one compiler-dependent test skipped; Python lint and diff checks
passed. Regression risk is low: only profile archive metadata, tests and
docs change, so no full multi-game regression run was performed.

Refs #758.
The poll-wait test reads pthread storage after registering its guest thread. Assert that setup invariant before dereferencing the nullable result so a setup failure is explicit and the clang-tidy nullability check passes.
Keep default-valued render functions unspecialized while loading and let the
existing Metal 4 descriptor carry their values into pipeline compilation.
Compute shaders and per-value render variants retain library specialization.

Use one stable fallback record per function so overlapping MTLDevice builds
share a single specialization without blocking unrelated functions. Preserve
the record when the same library entry is registered again.

A gated-library regression fails on the original PR and passes with the fix.
Metal pipeline tests: 90 passed; pending-pipeline tests: 52 passed; docs: 11
passed. The full clang-tidy ratchet passed with no new findings.
Full regressions for Breath of the Wild, Pokemon Scarlet and Tears of the
Kingdom stopped before gameplay because the offline Metal compiler is missing.

Co-authored-by: KillReal1234 <48804399+KillReal1234@users.noreply.github.com>
…738)

Eight color targets and independent depth/stencil store decisions require
ten records. Size the ledger and lookahead buffer from one constant to
prevent the tenth record from corrupting cached pipeline formats.

Retain both physical planes of combined storage, preserve logical aspect
use, normalize live/predicted/historical keys, and reject invalid formats
before either Metal pipeline compiler can assert.

Add asset-free GPU coverage and make its parsing, readback allocations and
format comparisons explicit for the clang-tidy ratchet.

Co-authored-by: Ethan Cha <43095500+ripull125@users.noreply.github.com>
Readback and PNG capture stall measured frames and can change adaptive
pacing. Request default visual proof after the timing window and parse
game speed from the host's current frame-step log format. Preserve
explicit screenshot schedules and document their timing cost.

Link the finding from the Breath of the Wild and Tears of the Kingdom
game notes. The PR's region lint corrections are already on main.

Validated with the 24 benchmark, region and documentation tests.

Co-authored-by: Leandro Piccione <15013313+ComicBit@users.noreply.github.com>
Run the x18 scan and spin-wait analysis in forked pools while keeping
patch emission ordered. Share a worker budget across module rewrites,
honor explicit DM_REWRITE_J settings, and reject invalid counts before
loading modules.

Exit through pool cleanup on SIGTERM so cancelled rewrites reap their
workers. Add regression tests, a rewriter reference entry and its symptom
index entry, and document the completed game validation.

Validation: 173 focused tests, the Python lint/type ratchet and docs checks
passed. Breath of the Wild's full regression passed; all 21 binary and TSV
outputs of its three modules were byte-identical. No completed full-game
regression for other titles in this change.

Co-authored-by: Daniel Johnson <41651107+djdev0010@users.noreply.github.com>
Host-owned shader libraries compiled from source at first use, including
one clear library per attachment slot and output type. Move their MSL to
build inputs, combine the clear variants, and load cached metallibs in
the shared graphics paths.

Compile Mac libraries for MACOS_MIN and recheck the compiler, SDK,
deployment target and sources on every build. Reuse unchanged headers,
but regenerate source-only headers when the Metal compiler becomes
available. Prepare and fingerprint the generated header in the NSP
packer, regression host cache and iOS compile check.

Validation: graphics tests, tool tests, focused build/cache tests, Python
and C lint, documentation checks and iOS source compilation pass. Full
game regression was not completed; gameplay timing remains inconclusive.

Co-authored-by: KillReal1234 <48804399+KillReal1234@users.noreply.github.com>
The settings screen offered pipeline prebuild even when Metal 4 pipelines
were off, although both runtime prebuild paths require them. Add the
existing needs metadata to the shared settings row and document when it
appears. Preserve the platform defaults and saved choices.

Validation: all 17 focused settings and native Swift launcher tests pass.
Regression risk is low: only settings metadata and prose change, so no
game regression runs were needed.

Co-authored-by: v0lumehi <34629096+v0lumehi@users.noreply.github.com>
Emit a shaders_stale event before rebuilding libraries after a chain or
profile input change, with the same explanation in human progress modes.

Check completion stamps separately for the selected targets so partial
current-chain outputs cannot hide old libraries. Count each current shader
key once per selected target instead of summing retained historical copies.
Document the count as scheduled base libraries, including cache restores,
and cover interrupted rebuilds, mixed targets and duplicate old caches.

Co-authored-by: v0lumehi <34629096+v0lumehi@users.noreply.github.com>
A game launched from another fullscreen app could request fullscreen before
its window reached the active Space, causing AppKit to roll the entry back.
Wait for app activation and window Space readiness before making the request,
and allow one repeat when AppKit reports a refused startup attempt.

Track the outstanding startup request and retire its failure handler after
successful entry or player takeover while waiting. A later refused manual
entry can no longer restart startup handling. Keep native fullscreen and the
existing screen refresh-rate matching behavior.

The reporter reproduced the original failure on four of four Hades II
launches and confirmed the readiness fix. The revised callbacks pass all 11
lifecycle harness cases; the original PR fails two. Mac host build, host unit
tests, documentation tests and targeted LLVM 23 clang-tidy pass. Regression
risk is low and confined to fullscreen startup, so no game regression runs
were performed; the revised lifecycle still needs native launcher testing.

Co-authored-by: v0lumehi <34629096+v0lumehi@users.noreply.github.com>
The iOS launcher searched only the app bundle for DLC, even when the base
RomFS was deployed separately. Preserve DM_AOC_DIR overrides, then search
Documents/aoc before the bundle, and exclude external DLC from backups.

Push the complete DLC set from resources.env with replacement semantics.
Directory merges retained removed indices, which the host still counted
and mounted. Send an empty directory for builds without DLC, reject a
missing declared source before transfer, and preserve DLC for image-only
pushes that carry no DLC metadata.

Add synthetic transfer regression tests and retain the author's Foundation
startup tests. Document the mechanism and regenerate the symptom index.
Validation: 59 tests passed, one disk-image test skipped; 185 iOS sources
compiled; Python lint has no new findings. Regression risk is low: only
iOS discovery and transfer change, so no game regression run was needed.
Device transfers remain untested.

Co-authored-by: v0lumehi <34629096+v0lumehi@users.noreply.github.com>
The fixture could miss backoff when CI descheduling broke every short
polling run. Control its classification clock while keeping real SVC
wait and wake paths, and check threshold, gap and reset boundaries.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Flush put stb_vorbis into Ogg resynchronization, which retains trailing
bytes for the next call. The host passes complete packet pages and drops
those unconsumed bytes, so decoding never recovered after a seek or loop.
Close the decoder on Flush and reopen it from retained headers at the next
packet, resetting overlap and synthetic Ogg state together.

Keep the synthetic repeated-Flush regression and the documented finding.
The regression fails with the old implementation and passes with this
change; the movie VP8/VP9, AVC, Vorbis and documentation checks pass.
Confirmation on the affected iPhone remains pending.

Co-authored-by: Oliver Ames <7366664+oliverames@users.noreply.github.com>
A 32-bit checked index copied with and xI,xS,#0xffffffff was treated
as a masked index requiring another matching mask on the checked value.
Recognize this zero-extension pattern after every existing jump-table
recognizer, preserving previously selected table extents.

Verify the guarded path through the load, add and branch. Reject register
clobbers, calls, system calls, traps, undecodable gaps and entries bypassing
the check; validate the table extent and decoded targets.

Add focused acceptance and rejection tests, including complete synthetic
rewrites for the new five-entry form and preservation of an older
eight-entry match. Correct the reference's preservation claim and update
the generated symptom index.

Validation: 70 jump-recognition, patch and documentation tests pass with
the pinned dependencies; the Python lint/type ratchet has no new findings.
Regression risk is low: the new recognizer runs only after all earlier
forms fail, so no game regression runs were needed. Mario Tennis Aces
gameplay remains untested.

Co-authored-by: Iheb nachet <45464932+iheb69nachet@users.noreply.github.com>
Publish bounded callback-statistics snapshots for the existing watchdog so
stdio locking cannot block PCM delivery or guest buffer-release events.
Preserve unread reports and merged-window counts, and copy log arguments
before releasing the output-registry lock or freeing an output.

Cover blocked diagnostic output, independent release events, backlog
accounting and reports that outlive their audio output.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
bench.py did not pass the build's pso/mac archive to the host, so a
benchmark could spend its entire measuring window rebuilding pipeline
bases before the first frame. Default DM_PSO_BUNDLE to that directory
when present, preserving inherited and explicit overrides as the Mac
launcher does.

Check the archive handoff in the CLI test and isolate its inherited
DM_PSO_BUNDLE setting so supported overrides cannot break the test.

Validation: six focused tests pass with the setting unset, off, and a
custom path (18 tests); Python lint has zero new findings. One initial
run produced no frames; subsequent runs passed. Game regression runs
were skipped because the guarded tool default and test isolation have
low regression risk.

Co-authored-by: Tom Ramsden <13732780+gnosticgeek@users.noreply.github.com>
Register update 1.0.2 with its shader tables, gameplay note and gallery.
Both games reach exploration and resume autosaves on the Mac; the full
Mac shader build and signed app are complete.

Accept full-width bounds in ADR-relative jump tables and track copied
index provenance. Include paired pixel dependencies in partial shader
builds and prevent incomplete vertex outputs from entering the cache.

Validation: 98 focused tests; all 69,924 shader records; offline pipeline
archive; strict Mac app signature and resource inventory. Moderate shared
tool risk: build/interface regressions pass for ANIMAL WELL, Bastion and
SteamWorld Dig across three NVN engines. Rewriter outputs for Tears of
the Kingdom, Mario Wonder and ANIMAL WELL are identical to baseline;
runtime comparisons for the first two were inconclusive on a busy Mac.

Gameplay testing stopped at the user's request. Combat, later areas,
a final-library replay and iPhone builds remain open.
The iOS archive default targeted only g18p, so another GPU family needed
an explicit build override. Read optional pso_gpu targets in both packers
behind the existing pso_mtl4 opt-in gate, retaining explicit overrides
and the previous defaults when the key is absent.

Keep game settings unchanged. Record the contributor's Wonder archive
probe in its game note and clarify archive selection and device validation.

Validation: 50 targeted tests run, one Metal compiler test skipped; all
121 profile headers unchanged; shell syntax valid. Pinned Python lint
reports the same pre-existing reportArgumentType at
dm/tools/tests/test_shader_toolchain_cache.py:134 on fresh origin/main
and this commit (393 findings, 392 in its baseline); no additional finding.
Regression risk is low: no shipped profile uses pso_gpu and the host and
shader chain are unchanged.

Co-authored-by: Oliver Ames <7366664+oliverames@users.noreply.github.com>
The GL packaging path copies shader_map.tsv and pso_manifest.tsv, but
Well Dweller's profile omitted both and --prepare-xcode failed at the
copy step. Add empty tables matching the established GL profile format.

Verified the packaging failure on the baseline and the successful copy
with these tables. All six PR CI checks passed; the contributor reports
boot checks on an iPhone 17 Pro Max and an iPad Pro 11-inch (M4).

Fixes #824.

Co-authored-by: Phil Wu <194177+runty@users.noreply.github.com>
A guest SyncWait can run post_run on a long-lived thread whose outer
autorelease pool never drains. ARC can retain and autorelease the
receiver of MTLBuffer.contents even though explicit readback ownership
is balanced, keeping completed buffers alive indefinitely.

Drain those temporary lifetimes per completed record, after releasing
the post-op mutex, while preserving guest writes and retirement order.
The gated real-GPU test covers guest and callback ownership, mirrored
bytes, event visibility, late feedback and buffer retirement. It passes
with both encoders and fails at the lifetime assertion on unfixed main.

Co-authored-by: Leandro Piccione <15013313+ComicBit@users.noreply.github.com>
Validate the complete command list before histogram discovery or PCM
output. Initialize the histogram and stop appends before snprintf's
untruncated length can advance beyond the buffer.

The regression fixture covers malformed-list rejection and stereo
recovery across SDK 4/5/6/16/21, plus oversized histogram formatting.
The previous code reproduces both bounds violations under ASan/UBSan;
the fixed fixture passes plain and sanitized with patterned stack locals.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Always close an opened metadata stream, preserve the first error across
close and cleanup, and rename only after writing and closing succeed.

The production-writer failure-injection fixture fails eight assertions
before the fix and passes all 41 afterward. The 54 existing save-sync
checks also pass, both plain and under ASan/UBSan.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Check the folder-backed directory handle allocation before dereferencing
it. On failure, close the opened DIR, preserve errno and return through
the existing filesystem error path without publishing a handle.

The real-handler fixture reproduces the prior null write under
ASan/UBSan. All 19 checks pass after the fix, including descriptor
closure, unchanged output handles and subsequent successful enumeration.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Record Metal 32023.922.4 for both platforms and keep CI, compiler checks,
player instructions and the generated shader chain in agreement.

Exclude only sleeping zero-CPU orphan MSBuild reuse workers from the
regression busy-machine check, with active/parented/build-command controls.
This fixes an idle build server blocking validation without relaxing the
compiler pin or measurement thresholds.

Earlier-source validation covers five shader profiles without new build
or interface failures, the native graphics suite and unsigned iOS builds.
Gameplay timing remains inconclusive with concurrent work on the Mac.
Separate possible image writes from raster writes in attachment tracking.
Suppress the tile store for loaded color attachments written only through
images, while preserving raster/clear and depth/stencil store decisions.
Capture physical output facts before merged-pipeline mapping is reset.

Include GPU readback controls for direct and deferred encoders, sticky
raster writes, clears, specialization and the optional store-saving switch.
Based on Jacob-Tate's diagnosis and initial separation in #187.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Match ast.Module.body's list[ast.stmt] type on Python 3.15. The selected
function nodes and fixture execution are unchanged; this avoids the list
invariance diagnostic without widening the Python lint baseline.

The 19 shader-cache tests pass and the pinned lint tools report no new
or stale findings.
NVN clears all layers of the bound render-target view. A load-action
clear in a non-layered Metal pass covered only its first layer, leaving
the remaining shadow-cube faces and array layers unchanged.

Extend the existing standalone clear pass to cover the view's full layer
range before loading it in the draw pass. Extend texraw diagnostics to
save every array slice and cube face, and test the production replay.

Integrate with #730 by marking a raster clear only when the draw pass's
actual load action clears. A clear already completed in a standalone
pass must not keep a stale tile store over later shader image writes.
The literal GPU regression fails four checks without this tracking fix.

Before rebasing, the full Metal graphics suite, LLVM 23 lint, all 11
documentation tests and Breath of the Wild's full outdoors route passed.
Those comparisons used the documented Xcode 27.2 compiler override on
both revisions against the old compiler lock. After rebasing onto the
Xcode 27.2 pin and #730, focused lint and documentation checks pass, as
do 120 image-store checks and 30 layered-clear checks under Metal
validation. The maintainer stopped the remaining Tears of the Kingdom
and Pokemon Scarlet regression runs; document that partial coverage.

Co-authored-by: Riski Aditiawarman <62930684+keanu-stitch8@users.noreply.github.com>
seedds and others added 27 commits November 1, 2026 21:07
…are, mesh geometry, linear run checks, drains only for later writes, SUATOM

- Ampere HMMA.16816 (F16 accumulator) in compute shaders goes through simdgroup_matrix: the translator emits a helper call
  (TensorPass, HelperFunctionManager), dmsh_build.py rewrites its body; agl's upscaler 2.3-6.5 -> 1.1-1.6 ms a frame.
- Geometry runs check dependencies against the run's merged write and read intervals instead of every earlier job (O(n) not
  O(n^2)); DK's interval switches 30 -> 12 on its scripted route.
- After a drain, CPU reads of GPU-written data wait only for writes encoded since it (sbw_note / pend_add feed a list): the
  level opening went from 1.3 s over three frames to one frame of 0.2-0.3 s.
- Mesh geometry: a threadgroup takes as many triangles as fit Metal's 32 KB mesh output (dmsh_mesh.prims_for, entries named
  dm_mesh_p16), a failed mesh pipeline sends its pair back to the compute path, and DK's geo_mesh.tsv lists its one pair:
  translation 9.9 -> 7.2 ms, GPU 11.8 -> 9.0 ms a frame; two per-draw lookups cached (6.6 ms).
- SUATOM (0x593) decoded: all 22,852 of DK's shaders translate.
- A program's shader libraries load in the background when the game makes it (mtl_prog_prefetch, DM_NVN_PREFETCH=0 off).
- Shift + arrow keys drive the right stick on the Mac; t_nvn's timestamp check follows the profile's render_cond.

Chain 9de03be216c6. Tests: the graphics suite, the geometry suites, all Python tests and the docs test pass. A play session:
frame p95 under 18.4 ms in every window, one interval switch. The regression run was skipped at the user's request.
…cache compares a byte key, and two tests guard their indices

- dmsh_build.py hmma_msl: a signature that does not match raises ValueError instead of calling .group on None (pyright).
- nvn_geometry.inc geo_color_omit: the cache key is a byte array (the target record's bytes, serials, mask, clears), not a padded
  struct compared with memcmp (clang-tidy bugprone-suspicious-memory-comparison); unequal padding only costs a recompute.
- t_cheats.c: the dry-run check tests fl >= 0 before g_cheat[fl] (clang-analyzer ArrayBound).
- test_split_heaps.py: heaps_pak_files' result is asserted not None before it is indexed (pyright).

Chain b7cd656abd7c. lint_python: CI's six new findings gone (three left locally come from this machine's venv having unicorn and
numpy); lint_c tidy with LLVM 23.1.4: 272 findings, the baseline's 272. Python tests (115 files), host tests, geometry suites pass.
…rom any thread, the prewarm starts once outside it

The background shader-library loading (mtl_prog_prefetch) made shmap_load a dispatch_once around the whole load, including
warm_start at its end. pso_prebuild, which warm_start runs, calls shmap_load itself; that used to return at once (the map was
set), and inside the dispatch_once it re-entered it: libdispatch aborted Tears of the Kingdom and Pokémon Scarlet at boot
(_dispatch_once_wait.cold). Donkey Kong Bananza runs no prebuild, so it never showed.

shmap_read_once is the dispatch_once around reading the map, the library directory and its tables; shmap_load reads through it
and then starts the prewarm once, its flag set before warm_start runs, so a call from inside it returns as before. The prefetch
reads the map only, so it never starts the prewarm on a game's thread.

Regression (regress.py check, baseline 742c9fb, main before 4de40e6): Pokémon Scarlet PASS (29.9 fps, p99 33.97 vs 33.92 ms);
Tears of the Kingdom CHANGED, nothing worse (59.8 vs 59.9 fps, p99 17.89 vs 17.77 ms; one shader changed: 588e242a52f5f7aa now
reads memory through the pool table, the general global-memory fix of this series). The run before this fix crashed both.
…olved before using it

The boot fix (a1b9a2a) changed how clang's analyzer walks prog_get -> shmap_load in this test (it compiles nvn_mtl.m in), and two
findings appeared in it: a null dereference of the program's fragment stage (now checked: the test fails cleanly when the program
does not resolve its two stages) and a RetainCount leak report on the program prog_get returns, a false positive the file already
silences elsewhere (g_by_serial owns it through serial_put's CFBridgingRetain): // NOLINT with that reason.

lint_c tidy (LLVM 23.1.4): 272 findings, the baseline's 272. t_image_attachment_store: 120 checks pass.
….keys

build_ipa_from_nsp.py takes a dump folder (base/, updates/<version>/, dlc/) in place of NSPs and calls pack_dump.py for the step extraction does, then goes on with its other steps. A dump's files are not encrypted, so it needs no keys. It packs only a game with a profile, as the survey reads NSPs, and has no control data: the name comes from the profile and the app has no icon. pack_dump.py raises an error the script can report instead of exiting, and keeps code modules that are unchanged.

README: How to play points to the Switch 2 section, which now uses this script.
…rm starts are kept

Hollow Knight (0100633007D48000, update 1.5.12620, SDK 20.5.21, Unity 6000.0.61f1 IL2CPP with FMOD on NVN): a profile
with its note, games table row and gallery screenshot, and no game-specific rewriter cut or host change. Played on the
Mac: the title, a new game, the opening cinematic, then King's Pass with its fights, deaths and the shade, and a save
loaded after a restart. 60 fps with a window (p95 16.7 ms; 30 while the cinematic plays, as the game asks); the alias
region's high-water mark +1.75 GB of 4.38 GB. Open: the iPhone, and everything past King's Pass.

The profile sets depth_mode zero, since the game sets near = 0 at startup; with the default the prebuild's pipeline
bases are for the wrong vertex entry and go unused. pso_manifest.tsv holds 110 lines recorded along that route. The
shader tables and the other keys are those of PR #735, whose tables a separate run of the NSP script reproduced byte
for byte.

warm_enqueue makes its own list: warm_start made it, so a library the prefetch finished before the first submission, or
in a run with no prewarm list, was dropped and its first draw loaded it again. dm/gfx/tests/t_shmap_warm.m covers that
and a1b9a2a's boot abort, needing no game: the map's read leaves the prewarm unstarted, and the first submission gets
through the prebuild's second shmap_load. The same sequence trapped with exit 133 before a1b9a2a. docs/reference has
the entry that commit did not add.

Checks: test_docs, 114 Python test files and make -C dm/gfx/tests run pass. test_host_msl.py fails, the same 6 failures
and 1 error on a clean tree of main's code. C lint was not run locally (no clang-tidy).
Regression (regress.py check, baseline origin/main, DM_WINDOW=0; moderate reach: 3 games chosen, stopped by the
maintainer after the first): Blasphemous PASS, its baseline crashing with the boot abort, on a build with an earlier
form of the boot fix; Super Mario Bros. Wonder and Kirby have no result. Hollow Knight ran with a window on this code.

Co-authored-by: joaquinsandoval2307 <279042628+joaquinsandoval2307@users.noreply.github.com>
… image marks them and the host names them (#810)

Issue #810's crash in nn::util::RelocationTable::Relocate came from the reporter's own files: their Event/Demo154_2.sbeventpack has our length but another SHA-256, from a base NSP that fails the hash check, extracted with --allow-damaged. Two general gaps made it look like a host bug.

nx_extract: LibHac 0.19's CachedStorage.GetBlock takes a block off its list before reading into it and does not put it back when the read fails, so after 4 failed reads the next read that misses the cache throws "The LinkedList is empty", which is no hash error. nx_extract stopped with exit code 134 and no list, with or without --allow-damaged. Each hash error now gets a new checked filesystem, on the image, folder and manifest paths. The compare with a base RomFS folder is inside the same catch.

Image and host: nothing in a run showed that a build held damaged files. The last byte of a DMROMFS1 index entry is now flags, 1 on a file that failed the check; every reader refuses another value or a flag on a directory. The host logs [RED] with the count at the mount and with the file's name at its first open. A clean image is byte-identical to before, so its identity does not change. Images extracted earlier and folder extraction are not marked.

Checked on a copy of SteamWorld Dig with 40 flipped bytes: the old build crashes; the new one lists 36 files and exits 3, or exits 0 with --allow-damaged. The flagged entries are exactly the 35 listed RomFS files (the 36th is ExeFS's sdk, which --allow-damaged still leaves out, as before), and all 31 files that differ from a clean extraction are flagged. The clean image is the same from the old and the new extractor. A boot on the damaged image logged the count and the files it opened. test_romfs_read.py puts a real RomFS behind a real CachedStorage on all three paths; t_fs.c checks the count, the single report and the refused flags. Not checked: an iOS build or a device. #810 stays open until the reporter's extraction from a clean NSP loads the scene.

The reporter diagnosed the lost cache blocks and carried a local patch that reopens the filesystem.
…ize set after the store; movies into BGRA8 textures

Update 1.1.1, Unreal Engine 4.27, SDK 12.3.2: profile, note, prewarm list and gallery entry. Tested on the Mac through the prologue and turns of the first battle at 60 fps; the end of a battle and the iPhone are untested.

ue_region.py did not recognise a reservation whose x1 is set from the constant after the store of the size, so the engine asked for its whole 62 GB and aborted in a 3.75 GB block. The rewriter now makes that instruction mov x1, xS. Of the 131 main files on the drive, only this game's result changes.

The host wrote movie pictures in RGBA order into any client texture, swapping red and blue in BGRA8 and BGRX8 ones; they now go through a copy with the two swapped.
…#855)

The notifier cleared notify_armed with a release-only store, then read
notify_pending. A producer could set pending and see the old armed value, so
it did not post, while the notifier read the old pending and went back to
sleep: PCM played and the buffer was released, but the guest's event never
came. Apple-silicon code is stlr then ldapr, which permits this.

Disarm with an acquire-release exchange. If the producer's exchange came
first, the notifier acquires its pending store; otherwise the producer posts,
or finds the notifier already re-armed with its own post.

t_audio_stats gains 100,000 individual handoffs. 400 runs each, old notifier
against the fix: plain 8 failed against 0, ASan/UBSan 74 failed against 0.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
The window stayed on screen, not responding, until the game ended: after
run() returns, launcher.c blocks in waitpid, so no run-loop pass commits
the order-out or runs its fade. finish turns the window's animation off,
and dm_mac_launcher flushes the order-out before it returns. Neither
alone is enough.

Closes #688

Co-authored-by: Suraj Mandal <11650177+surajmandalcell@users.noreply.github.com>
…es PROT_NONE and MADV_FREE, which kept the pages resident

On Darwin mprotect(PROT_NONE) only drops the page table entries, and madvise(MADV_FREE) does nothing to pages that are not writable, though it returns 0.
UnmapPhysicalMemory and a heap shrink now put a MAP_FIXED inaccessible mapping over the range. After an unmap the range's exposed bits clear, so mapping it again needs no memset.

Mac A/B, one run each: Scarlet peak footprint 4,540 -> 3,775 MiB, guest 1,809 -> 1,372 MiB. TotK footprint flat, guest 1,887 -> 1,685 MiB.
regress.py check flagged TotK's 57 s screenshot: both pictures render correctly and differ by where the moving scene is; baselines differ as much. iPhone and Z-A not measured.
…ngths and bounded inflation (#659)

bnd2_resources sliced stored ranges without checking them and inflated without a limit, so a truncated stored block was accepted and a compressed one could outgrow its record. Now the header and table must lie inside the file, a block's data after the table and inside the file, and its declared size (alignment bits masked) must match what is stored or inflated. Inflation stops one byte past the declared size, and a damaged stream names its resource. A bundle that fails is named and its shaders are left out, as with any unreadable file.

On Need for Speed: Hot Pursuit Remastered 1.0.3 the output is byte for byte that of the old reader over all 3,370 bundles and 518,261 blocks, with nothing rejected. The note also records a contributor's iPhone and Mac tests.

Co-authored-by: yonatani94 <49233340+yonatani94@users.noreply.github.com>
Update 1.1.0 (v65536), SDK 19.3.16, Unity 2022.3.62f2 IL2CPP with FMOD on NVN. A profile, its tables, a note and a gallery image: no host, chain or rewriter change. The SDK-wide node area cut applies, and this sdk aligns regions to 64 KB or 2 MB itself.

The layout is Cult of the Lamb's (a 1 GB pool in the 5 GB block), and depth_mode is zero, as the game sets it. Played on the Mac: the first-run screens, the title, a new game, the opening movie, the dock dialogue, Paris Streets and the sewer's first room with a fight, a death and a save resumed after a restart. 60 fps with a window, the alias region's high-water mark +3.5 GB of 4.38 GB.

The registered profile was run from a fresh save to the pier with walking, dashing and healing; the fight, the death and the save resume ran on the provisional profile, which has the same shader tables. Not checked: iOS (no IPA built), sound by ear, a real controller, anything past the sewer's first room.
Guest releases now replace backing with mmap and fall back to mprotect. Update both fixtures to inject failures at both operations and check the new heap-shrink call sequence. Keep fatal-path state checks and cover successful fallback releases and zeroed alias recommits.

Validated: full host unit suite, full ASan/UBSan suite, targeted memory TSan and dirty-tracking tests, and LLVM 23.1.4 clang-tidy with no baseline changes. Regression risk is low: only test fixtures changed, so game regression runs were skipped.
…p GPU skinning (#857)

Hello Kitty Party Land (Unity 6000.3.11f1) drew its characters as huge stretched
triangles. Its ComputeShader objects (class 72, type hash 43a0948a...) have no
type tree from a bundle and no layout in unityfs.py, so the split left out their
kernels, Unity's built-in Internal-Skinning among them. The run then logged "the
dispatched program has no compute shader" and the characters drew from vertex
buffers nothing wrote.

COMPUTE_6000_3 is the 2022.3 layout with an int resType after each resource's
texDimension and threadGroupSize a staticvector again. It is keyed by type hash and
used only for stripped files no bundle lends a tree to.

All 159 nodes and alignment flags match AssetRipper's 6000.3.0b4, 6000.3.0f1,
6000.3.11f1 and 6000.3.20f1 dumps, compared here. PR #857 and PR #791 report the same
hash on different games; those game results were not repeated here, since no game
with this hash is available. The fixture from PR #791 reads to its exact end and into
the split, and fails with the 2022.3 layout. Silver Pines (6000.3.15f1) already has
these kernels in its table without this layout.

PR #791's Shader layout is left out: its hash is not confirmed.

Co-authored-by: James <12531452+capibawa@users.noreply.github.com>
Co-authored-by: Tom Ramsden <13732780+gnosticgeek@users.noreply.github.com>
A movie that loops without a Flush or end of stream restarts its PTS at the
next IDR. The reorder queue sorted by PTS alone, so the previous loop's tail
sorted behind the new pictures and used up their lookahead: later loops left
in decode order.

Tag each IDR-delimited coded sequence when its packet is submitted, sort by
sequence then PTS, and release a completed sequence's tail while the current
sequence keeps its depth. Flush still discards, end of stream still drains,
and timestamps are unchanged.

t_movie_avc feeds synthetic Annex B and length-prefixed packets through the
production parser, callback refcons and pixel copy, substituting only
decompression; it fails on the previous decoder. Separately, x264 B-frame
clips looped three times through real VideoToolbox had 123 of 180 pictures in
the wrong position before and none after. Gameplay was not run on this code.

Fixes #889

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Co-authored-by: zatteogit <64840342+zatteogit@users.noreply.github.com>
…unchanged (#895)

The offline archive's key held the chain label, the profile stamp, the
builder's source and the libraries' absolute paths, so a chain update
recompiled an archive from the same inputs. The key is now what metal-tt
compiles from: the pipeline script without library locations, the referenced
libraries' names and bytes, the architecture, the platform, the minimum OS and
metal-tt's version. The reflection is still checked against the current chain
first, and prebuilt.tsv and the other metadata are written on every run.

Measured on Pokemon Legends: Z-A after a comment-only chain edit: 462.78 s
with the old key, 10.10 s with the new one. An update that changes a library
in the archive still rebuilds it. Existing keys rebuild once.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
The split's key held every graphics Python tool, every profile file and the
whole decoder host, whose compiled DMSH_CHAIN label changes with each chain,
so a translation-only update extracted an unchanged package again. The key now
follows shader_split.py's own imports at any depth and only the profile
settings and tables the split reads or is checked against, plus the modules it
reads and the capture folders' files.

A decoder host records its build inputs with only the chain label normalized,
bound to its executable's bytes: any other host source, the compiler, the SDK
or the generated profile header still splits again, and an unknown or replaced
host counts by its bytes. Rewritten modules are hashed by content. A second
key over the package and the supplemental pairs rejects damaged output.
Existing split caches warm up once.

Measured on Pokemon Legends: Z-A after a comment-only chain edit: 276.04 s
with the old key, 0.64 s with the new one.

Co-authored-by: Eliran Zach <4012041+thehobbit85@users.noreply.github.com>
Update 1.4.1 (v720896), SDK 18.3.5, Unity 2022.3.44f1 IL2CPP with CRIWARE on NVN. A profile, its tables, a note and a gallery image: no host, chain or rewriter change. The SDK-wide node area and region alignment cuts apply.

The layout is Cult of the Lamb's (a 1 GB pool in the 5 GB block), and depth_mode is zero, as the game sets it. Played on the Mac: the first-run screens, the title and its attract battles, a new game, the first map with a battle by hand, the rest of it and the start of a second map under Auto, a manual save, Continue on the autosave after a restart, and the start of the Anniversary Expansion. 60 fps, no red line, no crash and no memory growth over 40-minute runs.

The alias high-water mark reaches +4.27 GB of 4.38 GB with 1.82 GB of live regions, which is not exhaustion: VammManager::AllocateAddressRegion searches from the highest region down, so freed holes wait until the top no longer fits. With the alias region cut to 3 GB the end of the table stayed at +2.99 GB for about 10 minutes with no failure. docs/reference/memory.md has the entry.

The fresh-save route ran on the provisional profile, which has the same shaders; the registered profile ran the attract loop and Auto battles from a save. Not checked: the 19 CRI USM movies (none played on any route), iOS (no IPA built), sound by ear, a real controller, anything past the start of the second map.
…change (#893)

Symptom: a movie with a Vorbis track starts clean, then crackles and stutters, with no underrun in the log (#768). DM_AUDIO_DUMP shows exact zeros in whole 512-sample blocks of the game's mix.

Cause: stb_vorbis ends a packet's samples at its window's right_start, which depends on the next window, not at the window's centre. With blocksizes 256 / 2048 a long window before a short one gives 1,472 samples and the short one 128, where the specification has 1,024 and 576. The total is the same, but the outputs do not match the time of the next packet.

Fix: md_vorbis keeps the decoded samples and each packet leaves with (previous blocksize + its blocksize) / 4 of them, with the blocksizes read from the packets' mode numbers. Samples kept after the last packet lie past the centre of its window, where the specification has none, so they are dropped. A Flush drops them too.

Checked: t_movie_vpx, with a synthetic stream of both window sizes that fails with the old lengths; t_codec_video; the movie test under ASan and UBSan; CI's clang-tidy on the two files. Smash 13.0.5 on the Mac, both opening movies, the same build with the old and the new lengths: 411 silent blocks in short runs before, none after, and no more large sample jumps than before. Not listened to in this check; the PR's author listened to it. The crash in #768 is not fixed by this.

Co-authored-by: Eduard Rotaru <189956204+nobz226@users.noreply.github.com>
…auncher)

The Android port of nxnative: an Android host (libnxhost.so per game) that
runs the game's own code natively, NVN / NVN2 replayed on Vulkan (mesh and
task programs included), the launcher app with on-device preparation of the
user's own dump (code rewrite and shader translation on the phone), the
Readiness check, Lossless Scaling frame generation and foldable layouts.

The Android build is GPL-3.0-or-later as a whole: see android/NOTICE.md.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Ampere opcode 0x923 mode 0xb800 reads the task payload written through opcode 0x927 mode 0xb800. Keep this view distinct from the mesh stage's task-input view and cover it with a hand-encoded TaskEXT fixture whose loaded word controls OpEmitMeshTasksEXT.
@biosed

biosed commented Nov 11, 2026

Copy link
Copy Markdown
Author

Superseded by the mergeable, explicitly based review at biosed#1. Current upstream main no longer contains the Switch 2 native foundation, so rebasing only these three translator fixes onto it would not build.

@biosed biosed closed this Nov 11, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants