The game finalizes texture objects on its own thread ahead of the replay, which still draws the frame the game recorded
before that; on the console the memory still holds the texels those draws read. Kakariko moves about 30 of its
textures to new addresses every four or five frames, and each move finalized the old descriptors. The replay's draws
of the frame before then met finalized CPU-written slots, and the stale rule gave them the stand-in even when their
own entry was still there with its content. The village's haze dropped out of a frame a few times a minute, and the
haze sprites that met a stand-in left hard-edged clear patches. A play session in Kakariko logged up to 4,805 stand-in
lookups in 60 frames, 1,826 in the 60 frames around one haze dropout, and 612 in the 60 frames after a run of
patched frames.
Such a slot now resolves to its own entry while that entry was uploaded or drawn by the GPU, and was neither superseded
by another descriptor at its address nor overwritten. The lookup samples the image as it is: it neither takes the
entry back nor reads its storage, so freed storage is still never read (the access violation the stale rule was made
for). NX2NAVI_TEX_STALE_KEPT=0 gives the stand-in again. The perf log's finalize line counts these lookups.
In a replay of a recorded Kakariko walk, most 60-frame windows went from stand-ins to the kept image (up to 328 a
window, stand-ins 0). In play, stand-ins remain in short bursts where the entry was already gone (up to about 2,000 in
60 frames, against the kept image in most windows). The tester still sees flicker that the 640x360 frame recorder
does not catch; that is not settled by this.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Once the released entries pass RELEASED_MAX, the sweep dropped an entry that comes back intact as soon as its 120
presents were up. Kakariko releases the ASTC textures of its hazy air and takes them back every few hundred frames, so
about 30 of them went each time, and the frame that needed one again sampled its stand-in while it was made and
uploaded anew. Such an entry now waits out GAP_TIMES the longest it stayed released before a take-back, as an entry
whose content lives only in its image does. NX2NAVI_TEX_KEEP_TAKEN=0 goes back to the old rule, and
NX2NAVI_TEX_RELEASED_MAX (a test) lowers the count from which the sweep starts, so a short run reaches it.
The haze still dropped out with this alone: most of those frames came from finalized slots, which the next commit
deals with.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The watch's thread takes the kernel lock every two seconds to look at every guest thread; a release does not need it running, and wait_watch_s = 10 turns it on again for a hang report.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Hyrule Field, standing, 3 workers: with 1, 2 or 4 groups out the frame rate stayed within one frame per second of
the 50 the mode without pipelining makes, over three batches, though the replay is shorter. A frame there is the
replay and then about 6.6 ms of GPU work after it, which the next frame waits for; the design comment in vk_exec.c
now says so. Sending a part's finished command buffers while later groups record was tried and gained nothing.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
With NX2NAVI_REPLAY_PIPE=1 (off by default) the submitting thread hands a group's chunks to the workers, takes chunks
of it itself until none is left to take, and goes on with the ops after the group in a new command buffer after the
chunks' places in the part (vk_sync_append_later), also into the next group's front pass.
Up to four groups are handed out and not in their part yet; their chunks wait in one queue in the order of the stream.
The submitting thread waits for them only where the order needs it: for room, before a part goes to the GPU, before
copies and texture or buffer clears, where DLAA's run or the contact shadows' hook acts, and where a lookup would make
a texture entry or take back a retired one (the op goes again afterwards, as a chunk's does in turn).
What the chunks read stays as the stream had it while the submitting thread runs ahead: uniform snapshot drops count
from their place in the stream, a chunk reads only the inline uniform updates before its group, the HDR present's
source goes by stream position, the texture pool set syncs under the chunks' lock, each group keeps its own render
conditions, occlusion segments and closed query slots, and the queries after a group's last report go into the next
segment made after it. A group's first chunk runs its pass hooks after the group before began its pass; turns wait
for every earlier chunk, of earlier groups too. Locks the chunks take, the submitting thread takes while it is behind.
Draws keep the sampled-texture descriptors they resolved per draw context and reuse them without the texture lock
while their slot is unchanged; the workers waited for that lock at nearly every draw. Texture lookups that change
nothing no longer wait for their turn. The perf line says where the submitting thread waited, what the chunks took,
and which texture lookups waited; NX2NAVI_GPU_TIMELINE also shows the chunks and the waits.
Off by default because it does not pay yet: in Hyrule Field with 3 workers it shortens the replay (11.5 against 12.1
ms a frame) but not the frame (49.2 against 49.6 fps in the same batch). The GPU gets a group's commands only once its
last chunk is recorded, so the main thread's wait for the GPU after the frame's first submission grew by 0.8 ms.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
NX2NAVI_VRAM_BUDGET_MB changes only the number the decisions compare with. The ballast takes that much device-local
memory right after the allocator is made, in dedicated buffers of 256 MiB, and holds it until the allocator goes, so the
driver's own behaviour at its limit (slow allocations, memory moved to system RAM, refusals) shows as on an 8 GB card.
Checked on an RTX 5090: 24000 of 24000 MiB taken and held; through eight reloads of Kakariko the use went past the card
(32,218 MiB at the end) and Kakariko fell from 118-151 fps in the first four rounds to 10-24 fps in the last four. Texture
detail was still chosen as full: the choice reads the driver's budget, which the ballast does not lower.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The game's sound engine takes a voice per sound from the renderer's voice count, so sounds that are never let go
would show as a count that climbs from area to area until new sounds are refused. The line gives the voices in use,
those playing, the renderer's voice count and the most a census found so far.
Used to check whether the sound engine leaks voices above 60 fps: through eight reloads of Kakariko the most was 30 of
96 at a 60 fps limit and 31 at 144, and every return to the title brought the count back to 12 at both rates.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Maxwell samples a texture by its own descriptor, so a 2D texture under a 3D sample gives its texel at (u, v). Vulkan
has no 3D view of a 2D image, and get_view found no view for such a binding: the draw sampled a stand-in. Such a
texture (plain colour, one sample, up to 1024 x 1024 texels) now gets a 3D image of one slice with all of its levels,
made in the upload batch and copied again when the texture was uploaded or drawn into since, or another layer is
asked for (vk_texture.c vk_texture_flat3d_locked). NX2NAVI_TEX_FLAT3D=0 gives the stand-in as before.
Seen in Kakariko: two of the game's 4 x 4 RGBA8 textures are sampled as 3D textures, and each got its copy; the runs
through eight reloads at 60 and 144 fps ended without errors.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The day's update (main+0x3cdf70) adds trunc(2 x frames per update) game seconds a frame. With the set ostep's step,
the display interval in 60 Hz frames, that sum lost its fraction every frame: 1 a frame from 61 to 119 fps, 0 above
120 (the clock stood still), and at 60 fps a step a little under one frame dropped a second now and then. Its convert
now calls the patch page's +0xf60 (runtime/frame_rate.c frame_rate_day_host), which keeps the fraction per counter and
hands out the whole seconds. NX2NAVI_TRACE=perf logs, every 10 s, the seconds given beside those the cut would have
given.
Measured in Kakariko, 10 s windows: at a 60 fps limit 998 game seconds against 881 from the cut alone; at 144 fps
1,201 against 219 to 972 (the steady 1,200 a 10 s window should have).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
frames.tsv gains step_frames, game_step_frames and guest_presents. The game writes its step at its own present, a frame or two before the replay presents that frame, so the present count is there to line the two up.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Each texture made (with how its content started), uploaded, uploaded again, released, taken back or destroyed is one line with its frame, address, format, size, layers and flags, so a one-frame flicker in the thumbnails can be matched to the textures of that frame. events.tsv is written first: the thumbnails take about five seconds, and a game closed meanwhile kept no frames.tsv.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Mesh draws ended a group and stayed on the submitting thread: 119 a frame in Hyrule Field at about 13 us each, and the
groups around them were cut short. They are recorded in the chunks now: the entries filled ahead for the indirect ones
are the draw context's, the ring index is taken atomically and the ring is made before the chunks record, each chunk
counts the ops that break a run of them on its own, and the lazy lookups of the mesh and shading-rate entry points store
their pointer before the flag that says it is there. With the mesh fallback, and while NX2NAVI_GPU_PASSES copies the
workgroup counts, they stay on the submitting thread.
A chunk's closed occlusion query slots go into a list of its own and into their segments once the group is done, in
chunk order, instead of under a lock per query. The per-object motion check decides alone per draw context (its
variant memo is per context) and takes its lock only to list a draw. The occlusion query ring has 16384 slots: the
chunks' blocks are sized to their draws, and with 4096 the ring came round to slots not reset yet.
The perf line splits the submitting thread's time with workers on: ops outside groups by kind, the front pass, its own
chunks and the wait for the others, and which ops ended the groups.
Hyrule Field, RX 7900 XT, 3440x1440, FSR 4 Native AA, same batch: standing 51.7 fps with 3 workers against 43.6 with
the workers off (replay 11.2 against 15.1 ms); along the field route 48.9 against 44.7. Screenshots standing and along
the route match the replay without workers.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The chunks of a parallel group took query slots from the ring as they recorded, so the slots of the group's chunks
interleaved. The completion thread reads and resets slots in runs that follow the ring, and with interleaved slots it
made four times the result calls and a reset per slot: the wait after a frame's first submission ended 3.5 ms after
the GPU had run it instead of 0.6 ms, and the replay time the workers saved went into that wait (44 fps with 3 workers
as without). The front pass now gives each chunk a block of slots, as many as its draws that count samples and render,
in the order of the stream; a segment's slots are taken in ring order. Standing in Hyrule Field, 3 workers: 49.8 fps
against 43.7 with the workers off.
NX2NAVI_GPU_TIMELINE=<frame> writes four frames from that present on to the log: each replay, each timed part's submit
and GPU run, the waits for the GPU and the workers' groups, so a run shows whether the GPU waits for the replay, a wait,
or the CPU behind one.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Samples-passed resets and reports ended most groups in Hyrule Field (the game reports twice without a reset there, so
counting stays on): the front pass now replays them in the order of the stream, makes each occlusion segment and
queues its completion callback in that order, and each worker follows the counting from its chunk's first op and adds
the queries it closes to the segment the stream gives them. Per-object motion takes a lock instead of waiting for the
earlier chunks: its list's order does not change the picture. The locks the chunks share are taken only inside a
chunk, as outside groups the submitting thread alone replays; taken always, they cost the replay about 1 ms a frame.
The perf line counts the draws by what ended their group.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Stage 2 of NX2NAVI_REPLAY_WORKERS (default 0, the replay as before). A run of draws is cut into groups: a group ends
where the render targets change after its first draw or clear, and before an inline uniform update, a mesh draw and a
samples-passed counter op, which the submitting thread replays itself in the order of the stream. A group of 64 draws
or more is first walked by a front pass on the submitting thread: it applies the state and binds, settles the CPU side
of the timestamp reports and decides the render conditions in the order of the stream, and copies the replay state at
the first op of each chunk. The workers and the submitting thread then take the chunks in order and record each into a
command buffer of its own, which vk_sync_append puts into the part in the order of the stream.
The frame's pass hooks (contact shadows, DLSS) run in a group's first chunk, which the others start after. What still
depends on the order between chunks waits until every earlier chunk of the group is done: a texture lookup that would
make or take back an entry returns nothing and the chunk replays the op again in turn, and so do the contact shadows'
pass, per-object motion and the HDR pass's source. The tables the chunks share take locks (the shader table, the
descriptor lists, mirror use and GPU write notes, uniform snapshot drops, target frames, the DLAA camera votes and the
previous-frame motion notes); each chunk keeps its own occlusion query open, its own feedback memo and no pass marks.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Reading the thread-local pointer in the functions that run per binding and per descriptor cost about 0.3 ms of
replay a frame in Hyrule Field (45.0 against 45.8 fps). VkExec.dc now names the context; the draw and look-ahead
entry points fill it from the calling thread's on first use, the functions that hold the state read it from there and
the few helpers without it take it as a parameter. A replay worker's copy of the state carries the worker's context.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Stage 1 of recording a submission on several threads. A submission's draws come in runs between ops that have to
stay in the replay's order on its own thread (dispatches, copies, texture and buffer clears, events, the game's
barriers, calls). With NX2NAVI_REPLAY_WORKERS=1 a run of 16 draws or more is recorded by a worker thread: it takes a
copy of the replay state, records the run into a command buffer of its own pool with a draw context of its own, ends
it behind a full barrier and hands the state back. vk_sync_append puts that command buffer into the part between the
replay's own pieces before and after it, so the GPU runs everything in the op stream's order. The submitting thread
waits for each run, so every side effect keeps the replay's order; this stage checks that a run recorded apart from
a copy of the state gives the same pictures. The default, 0, replays as before.
walk's handling of the ops that do not steer it moves into op_one, which the worker uses too. NX2NAVI_REPLAY_RUNS=1
logs every 600 frames how the draws fall into such runs and which ops break them. The draw replay's hot functions load
their recording context once instead of at every use.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
vk_exec_draw.c kept what a draw, bind or dispatch carries to the next one in file statics: the state the command
buffer holds, the descriptor signatures and skips, the per-binding reuse, the shader and pipeline memos, the shader
lookup's caches and the submission's transient resources. They now live in one DrawCtx, and every thread starts with
the replay's own, so the replay behaves as before. A later stage can give each recording thread a context of its own.
vk_exec.c's inline commands (vk_exec_inline_begin) find the replay state the calling thread records with through a
thread-local pointer too, instead of the file's single state and a flag.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The BC3 encoder's colour half is a four-colour BC1 block (c0 above c1, or every index 0 when they are equal), so a BC1 image of an opaque texture decodes to exactly the texels the BC3 one did. The encoder writes that half alone with its specialization constant BC1 set; its push constants are as before, so the BC3 pipeline is unchanged.
Every ASTC texture the CPU provides is made BC3, and its first upload checks its blocks while it copies them into staging: when all decode with alpha 255 (tex_astc.c astc_blocks_opaque, from the block header for all but the endpoint modes with alpha, which are decoded), the entry gets a BC1 image in place of the BC3 one before the decode. Later uploads into a BC1 image check again and go back to BC3 for content with alpha. A large texture's copy and check run on the ASTC decode workers, so its first upload copies faster than before. A check of its own where the texture is made cost 1.7 s of the render thread over a seven-minute teleport tour.
On that tour (Kakariko Village, the Graveyard, Death Mountain Trail, Kokiri Forest, Hyrule Field, RX 7900 XT), against BC3 alone in the same batch: the ASTC textures in Hyrule Field at the end 2,568 instead of 3,104 MiB, video memory at its peak 11,814 instead of 12,374 MiB, process commit at its peak 25,539 instead of 26,139 MiB; the first 7,200 frames in 93.7 instead of 101.3 s and the run's slowest 1 % at 24.1 instead of 21.1 fps. --probe-astc checks the header test against a decode of every generated block, the loop, its copy and the worker pool against it, and each BC1 block against the colour half of its BC3 block. NX2NAVI_ASTC_BC1=0 (astc_bc1 = 0) keeps every ASTC texture BC3; the sampling probe NX2NAVI_ASTC_BC_PROBE it replaces is gone.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Where the lazy write-back is off (AMD's driver, outside Windows), what a binding written by a shader or a copy holds goes
back to guest memory at the end of every submission. With NX2NAVI_WB_SKIP_MB=n a written range of n MiB or more goes
back for its first 64 KiB, which carries the few words the CPU reads at the start of such buffers, and the rest stays
in the mirror, whose pool is then never evicted. A CPU read further in would see older bytes, and a CPU write further
in would have the next sync copy that page's older guest bytes over the GPU's, so it is a test switch, default 0.
Measured with NX2NAVI_WB_LAZY=0 on an RTX 5090, 75 s runs with warm caches, off against NX2NAVI_WB_SKIP_MB=1:
Kakariko walk GPU median 6.60 against 4.20 ms a frame, 98 against 132 fps; Hyrule Field walk GPU 10.09 against
4.88 ms, 76 against 132 fps, 1% low 61 against 98 fps. 78 to 94 MiB a frame stayed in the mirrors. The last frames
of both runs in each area show the same picture, and the logs have no errors.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The prologue of a part copies the pages the CPU wrote since a mirror's last sync from guest memory into the mirror when
the GPU runs it, a frame or more after the game's submit when the GPU is behind: the GPU then reads what the game
wrote there meanwhile, values meant for a later frame. With NX2NAVI_MIRROR_CAPTURE_KB=n the CPU copies up to n KiB of
those pages per part into a host buffer of the submission when the part is handed over, and the prologue copies them
from there. A pool keeps the GPU-time copy while a write-back of its mirror into guest memory is still on the GPU (its
fence, now kept per pool), when it has a twin (the device address table's writes reach guest memory unrecorded), and
for a range with a page that cannot be read without a fault (memory.c mem_copy_readable, under the freeze lock, so a
lazily armed range stays with the GPU). With the default of every mirror twinned (NX2NAVI_GPA_TWIN_MAX_KB=0) it
copies nothing; it is a test switch for the replay's late reads, default 0.
Measured on an RTX 5090 with NX2NAVI_GPA_TWIN_MAX_KB=64 (mirrors of pools over 64 KiB without a twin), 75 s runs
with warm caches, NX2NAVI_MIRROR_CAPTURE_KB 0 against 8192: Kakariko 0.95 MiB a frame copied in 149 ranges for
0.21 ms of the replay thread, 116.7 against 118.4 fps (interval median 6.26 against 6.46 ms); Hyrule Field walk
1.38 MiB in 224 ranges, 0.37 ms, 110.2 against 107.2 fps, 1% low 83.3 against 78.1. 0.08 MiB a frame stayed with the
GPU-time copy. A first version that took the freeze lock and asked the system for each range's pages cost 2.9 to 4.3
ms a frame. A Kakariko run with NX2NAVI_REPLAY_ASYNC=1 and the copy on played through without errors.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A staging buffer, a transient buffer of the replay or a descriptor pool that the driver refuses for memory used to end
the game at once. Now the buffers kept for reuse go back to the driver first: the staging cache, the staging arena's
free chunks, the write-back snapshot pool and the headroom's spare blocks (no new spare block for 30 s after that), and
the allocation is tried once more; only a second refusal ends the game. On integrated GPUs the device-local heap is the
machine's RAM and fills before the commit checks see anything: staging buffers of 11 KiB to 16 MiB refused there ended
the game in Hyrule Field and on the first scene of play (#465, #567).
NX2NAVI_STAGING_FAIL_AT=n (diagnostics) takes the n-th staging buffer as refused once, to test the path.
Check: a Kakariko run with NX2NAVI_STAGING_FAIL_AT=300 logged "staging buffer of 524288 bytes refused (-2): 348.2 MiB of kept buffers and spare
blocks given back, made at the second try" and played on to the end of its script. The transient buffer and descriptor
pool paths call the same give-back and were not forced.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A place's task and mesh programs all arrive in its first frames, and each one's draws are left out until its compute
pipelines are built, so with one thread the foliage of a first visit filled in one program at a time.
NX2NAVI_MESH_FALLBACK_WORKERS (auto: one thread for each 4 logical processors, 1 to 4) sets how many threads make
them, and the log says how long the queue took to empty the first time.
Measured with the fallback forced on, no pipeline cache and new exe names (cold driver cache), the adult save's
first 75 s: with 1 worker the first 5 programs were ready after 420 ms, 3,266 draws were skipped while a pipeline or
program was pending, 1% low 16.1 fps and 145 frames over 50 ms; with 4 workers 256 ms, 956 draws, 55.7 fps and 35.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A mesh pipeline whose fragment stage reads a location the mesh program never writes now gets a fragment module where
that input is a private variable set to (0, 0, 0, 1) at entry, or to (s, t, 0, 1) of gl_PointCoord when the mesh
program draws points. The fragment shaders recorded with two of the game's point-drawing mesh programs read location
5, which those programs never write; Vulkan leaves such an input undefined.
Checks: tests/fs_fill_inputs_check.py passes; 237 recorded fragment modules rewritten offline pass spirv-val; a
60 s Kakariko run with the change shows the area as before with no Vulkan errors.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
An entry the GPU drew into since its last write-back holds texels its storage lacks; the fence write-back
(vk_texture_write_back_small) sends such targets back when the CPU waits, about once a frame. An entry dropped before
that, for a budget, a commit or video memory squeeze, its release or a replacement at its address, went without, and a
texture the game then made again from that storage started from what the storage held before the GPU drew. drop_tex
now writes such an entry back first when the fence write-back would take it (wb_wanted) and the CPU has not rewritten
the storage since. An entry another target drew over (its storage is the newer one's) and drops at shutdown are left
as they are. The commit drop pause counts the drop either way (comment beside commit_drop_note). The perf line of
texture origins counts the write-backs.
Checked: the runtime compiles; --probe-vkmem passes (164 GPU checks, 0 validation errors; it needs a writable TEMP).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
AMD's Windows driver answers VK_ERROR_UNKNOWN for some pipelines its compiler gives up on, and the draws that need
such a pipeline were left out for good. create_graphics now tries once more with VK_PIPELINE_CREATE_DISABLE_OPTIMIZATION_BIT
(slower shaders for that pipeline instead of a missing draw) and logs the first such retry. NX2NAVI_VK_NOOPT gains a
value, mesh: only pipelines with task and mesh stages, the foliage's, are built without the optimiser, whose
optimisation is a suspect in device losses some minutes into a forest on AMD GPUs; with it a report can be narrowed
without slowing every pipeline.
Checked: the runtime compiles. A 60 s run with VK_NOOPT=mesh on an RTX GPU logs the mesh-only line, has no error and
the same GPU time as the default (4.09 ms a frame against 4.09), 134.5 fps. The retry cannot be triggered on this
GPU.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A game that stops with nothing in its log leaves no clue to which threads wait for which. A thread of its own now
looks every 2 s for guest threads blocked that long on a lock (ArbitrateLock) or on another thread's end
(WaitSynchronization on a thread handle) and logs each such wait once: the waiter, the lock's holder or the thread
waited for, and whether that one is itself waiting and on what. Waits on condition variables, events and addresses
are left out, since idle worker threads wait on them for good. At most 24 lines a run. NX2NAVI_WAIT_WATCH_S sets the
threshold (10; 0 off). The wait's start is now kept for every blocking wait, not only under the sched trace.
Checked: the runtime compiles and tests/sync_wake_check.py passes. A 90 s run with WAIT_WATCH_S=1 logged one chain
during an area load ("thread 0x102d has waited 3 s for the lock ... held by thread 0x105a, itself waiting 2.6 s"),
and nothing in steady play; with the default of 10 s, earlier runs log nothing.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The game asks for 1x or 2x on its ground, road and wall textures, so they blur a few metres off at the shallow angles
a third-person camera sees them at. aniso now defaults to 16 (the menu's Texture filtering row and the ini still offer
game, 2, 4, 8 and 16); as before it raises only filtered, mipmapped textures that tile, and shadow maps, data maps and
screen-sized buffers keep the game's setting. A player's saved choice is kept.
Measured: two 60 s runs of an adult save on an RTX GPU at 2560x1440 without DLSS, one with aniso = game and one with
16: the GPU time a frame went from 3.92 to 4.06 ms (+0.14 ms), the frame rate stayed CPU-bound (142.7 and 148.3 fps
mean, within the runs' spread).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Ocarina's last post-process pass stores RGB9E5 texels through an integer view into its R11G11B10F scene target, and
the stand-in for the untranslatable upscaler decodes those raw words. Now and then the same pass writes them into the
RGB9E5 image at that address instead (1 of 529 frames sampled in Hyrule Field on an RX 9070 XT, #522): the stand-in
noted only R11G11B10F images, took the R11G11B10F image there, which held an older picture, and the frame showed a black
scene under an intact HUD. Writes through an integer view into either image are now noted with their order, the
stand-in takes the one written last, and a native RGB9E5 image is decoded as it is, without the copy that turns raw
words into RGB9E5.
Checked: the runtime compiles; a 60 s run without DLSS (the stand-in in use) logs "RGB9E5 fallback decode on", ran at
142.7 fps with no error. The native write itself did not occur in that run; it is rare and AMD-reported.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The boot's prewarm and the shared list's preparation started a builder on every core at once, so the driver began
that many graphics pipelines at the same moment, its first ones of the run among them. NVIDIA's 621.14 was reported to
fault in a sizable share of starts that way; a driver's first graphics pipeline sets up its compiler. warm_builder
now takes a one-job mode, and both callers build one queued pipeline on the starting thread before they start the
others.
Checked on an RTX GPU with --probe-prewarm in a copy of a dev folder (its prewarm skip marker removed), twice each with
the baseline and this build, alternating: no fault in four runs. With the 20 s budget the first pass built 2,203 of
3,226 recorded pipelines before and 2,135 after; the second pass built all 3,226 in both. A fault rate of a few
percent cannot be shown in four runs.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The loader added the addend for ABS64 and GLOB_DAT and always dropped it for JUMP_SLOT, though the AArch64 ELF ABI
gives S + A for all three. A JUMP_SLOT that binds a guest function now gets its addend too; one that binds a host
function (an override or an observer, where an offset into the trampoline means nothing) drops it and logs the
relocation, so a module that ever carries one shows up in the log instead of jumping to the wrong place.
Checked: the runtime compiles. The game's three modules hold 11,698 JUMP_SLOT relocations, none with a nonzero addend
(read from the 1.1.1 ExeFS), so nothing changes for this game.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
nvn_tex.c multiplied a level's width by its samples and its blocks by their bytes in 32 bits, and
nvn_tex_linear_stride did the same: a descriptor with a width past 2^30 (garbled, or not set up yet) wrapped to a small
size, and the texture got a small allocation for a large image. Block counts, row bytes and GOB counts are now 64-bit
(mip, blk, pick_tile and adj_tile take 64-bit sizes), and the linear stride saturates at 2^32 - 1 instead of wrapping.
Real textures lay out as before.
Checked: --probe-nvn-tex gains two cases (an RGBA8 texture 2^30 + 64 texels wide, and its linear stride). Before, the
first failed: 2,048 bytes against 2^35 + 2,048. After, all 135 cases pass, the 40 layout goldens of the original SDK
among them; --probe-nvn shows the same four failures as the baseline.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A device that reports an axis held from the start (a virtual pad nobody holds, a trigger or pedal mapped onto a stick,
a stick resting far off its middle) turned the camera or moved Link on its own, and since pad_active counted it as
pressed, it could take the gamepad role from the controller in the player's hands whenever that one was idle, and keep
it. input.c now reads every axis through pad_axis: 0 until the axis has been within a quarter of its range of rest
(a trigger's rest is 0), then the device's value. The game's sticks and triggers, the menu's stick readout and the
choice of gamepad all use it. A stick the player holds at the start counts once it has been let go.
Checked: the runtime compiles; new tests/pad_rest_check.py (pad_axis cut from input.c, SDL stood in for) passes: a
held axis reads 0 for 100 polls, a released one reads the device, and axes and gamepads keep separate states;
tests/mouse_stick_check.py still passes. Not run with a real device.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
F11 switched fullscreen whatever modifier was held, so an overlay or capture tool's Ctrl+F11 or Shift+F11 shortcut also
switched the game's window. The Win32 window now ignores F11 with Ctrl or Shift held (Alt+F11 already arrives as a
system key it does not handle), and the SDL window with any modifier. docs/controls.md says so.
Checked: the runtime compiles; gfx_window_sdl.c passes a syntax check against the SDL3 headers. Not run.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The SM86 lowering keeps the uniform registers in the general registers after the highest one the program uses
(HighestGpr), with one spare so pairs stay aligned. The scan read the register fields as single registers, so a
128-bit load or store whose upper registers the program names nowhere else (LDG, LDS, LDL and STG, STS, STL with size
6) left Rd+2 and Rd+3 to UR0 and UR1, and the load overwrote them. The scan now counts Rd..Rd+3 (Rd..Rd+1 for 64 bits)
of these forms. An overestimate only moves the uniform registers up.
Checked: no translation changes in the game's 520 task and mesh programs, its 4,238 NVN2 compute programs, or the
2,450 vertex and fragment programs whose stored input id starts with 0, all from an installed shader cache, and all
pass spirv-val (vulkan1.3): the case does not occur in what this game version has shown so far.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
cache_die moved the process to the logical processors of the largest L3 cache whenever the L3 sizes differed. That is
right for a two-die X3D desktop part with 8 cores a die, but it also held the game to the 4 large cores of a mobile
part whose 8 small cores share a smaller L3, and to the 6 cores of a 12-core X3D part, where the recording thread's
reserved core then left the game 5. The die is now used only when it has 8 or more physical cores (perf.c
cpu_placement logs when one is passed over). And while the process runs on that die alone, RECORDER_CORE=auto
reserves no core for the recording thread, since the game's threads already have the die to themselves; =1 still
reserves one.
Checked: all runtime files compile. This PC has a single-die processor (one L3), where neither rule applies: a run
still reserves the recording core (logical processors 14 and 15 of 16) and logs no cache die.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
On Linux with RADV (and the Steam Deck) the game points MESA_SHADER_CACHE_DIR at mesa-cache beside the exe, a folder
nothing else writes, but left Mesa's limit of 1 GB. The stage libraries the library warm builds in play are kept only
in the driver's cache, and a preparation of every shader variant is larger than that limit, so what did not fit was
compiled again in play. The limit is now 8 GB in that folder unless MESA_SHADER_CACHE_MAX_SIZE is set, as NVIDIA's
cache beside the game already has.
Checked: all runtime files compile. Not run on Linux.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
NX2NAVI_PREPARE_PIPELINES=mesh, the default, builds only the list's mesh pipelines, because a vertex pipeline is
fast-linked from its stage libraries in a few milliseconds and its vertex module depends on the output union this
install has recorded. Neither holds where the library is off (vk_caps_gpl: AMD's Windows driver for RDNA 2 and older,
or gpl = 0): a vertex pipeline a draw misses is compiled whole while that draw is left out, and its vertex outputs come
from its own fragment shader, since independent layouts need the library, so the listed pipeline is the one the draw
asks for. There mesh now builds every record, and the summary line says so.
Checked: all runtime files compile. Not run: no install on this PC holds the shared list, and the case is a GPU without
the library (NX2NAVI_GPL=0 would show it on any GPU).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The loops the translator writes for narrow storage buffer stores and float atomic adds end only when their
compare-and-swap succeeds. With robust buffer access on, an atomic on a word outside its buffer may return anything,
so the compare can fail for ever and the GPU hangs until the device is lost, as reported right after the first
cutscene on an Intel Arc GPU (#610). New translator patch 0041 counts the attempts in a private word (a local cannot
count round a loop in that pass, which runs on SSA form) and leaves the loop after 4,096; contention on a word needs
about as many attempts as invocations race for it. Shared memory loops stay as they are: a shared word in range is
always settled, and the mesh programs run that loop for every primitive index byte.
Checked: all 4,238 NVN2 compute programs an installed shader cache stores translate and pass spirv-val (vulkan1.3); 45
contain a compare-and-swap, 11 of them in a storage loop, which grew by 256 bytes each and now compare the count with
4096; the other 34 are unchanged. tools/xlate_smoke.py passes.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Besides the per-frame update that the set ostep patches, the graphics system's update (main+0x3b4ea90) divides 60.0 by
the engine's nominal rate and stores the result as a global (main+0x4d41fc8): the 60 Hz frames of an update. Its one
reader is the buttons' hold-to-repeat timer (main+0x3483680), which adds it to each held button's time and repeats the
button once that passes the button's delay and period. Above 60 fps the global stayed 1.0 a rendered frame, so a held
button repeated fps / 60 times too often. Two more ostep sites load frames per update from the patch page there, so the
two steps are the same number. NX2NAVI_SPEED_PROBE=1 adds the global to the fps line.
Checked: the generator's check passes for all 34 main sites against the 1.1.1 module, and the generated C differs from
before only in the patched function. Two 60 s runs of an adult save at about 135 fps, uncapped: before, the global
step read 1.000 frames while the game's step was 0.42-0.50; after, it read 0.43-0.46, inside the game's step range;
frame rate unchanged (132-138 fps in both).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
An NVN2 compute program's control section gives its local memory in bytes a thread at +0x68, beside the shared memory
size at +0x64 that the job already carried. Without it the translator declared its default of 0x1000 bytes: one of the
game's compute programs needs 4,752 bytes a thread, and its translation declared a 1,024-word local array that it
indexed up to word 1,187, outside the array. The job now carries "local_mem" from +0x68, and the lowering starts a
compute program's spill area after its own local memory (at least 0x1000 bytes) instead of always at 0x1000, where a
larger program's own memory would have overlapped it.
Checked: that program's translation now declares 1,188 words (indices up to 1,187) and passes spirv-val; two programs
of 3,456 bytes a thread translate as before. All 266 compute programs whose stored input id starts with 0 in an
installed shader cache translate and pass spirv-val (vulkan1.3), and tools/xlate_smoke.py passes.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The game's scene-loading thread makes a few NVN calls with a null object while an area loads or Kakariko's time of day
changes: a memory pool's flags (and the other pool getters), a texture builder's storage size or alignment, and the
mapping of a buffer that has no pool, into which it then writes 0x3d0 bytes. Each read the object at address 0 plus a
field's offset, an access violation. The pool getters, nvnMemoryPoolInvalidateMappedRange and the two builder queries
now answer 0 for a null object (nvn_no_object), and nvnBufferMap of a buffer with no pool returns a block of scratch
memory (64 KiB, or the buffer's size when it records a larger one) instead of 0. The first such calls are logged with
their caller. These fit the crash reports in Castle Town and in Kakariko before Death Mountain (#605, #602), which have
not been confirmed to be this.
Checked: --probe-nvn gains three checks (pool getters of no pool, builder queries of no builder, a 0x3d0-byte write
through the mapping of a buffer with no pool) and they pass; its four other failures are the same on the baseline.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A mesh program's header lists the components of each output vector it has; the fragment stage may read the others.
The native mesh entry function never wrote them, so they read as whatever the output held, and the fallback's vertex
stage wrote 0 for all of them, w included. Maxwell gives an unwritten component (0, 0, 0, 1), as the pipeline code
already pads a missing vertex output. Both paths now write that default: MeshSpirv.cs after copying each vertex's lines,
mesh_fb.c meshfb_vs with w as the bits of 1.0f whatever the type.
Checked: all 520 task and mesh programs of an installed shader cache translate and pass spirv-val (vulkan1.3); 500
mesh modules gained about 730 bytes each on average, so many of the game's mesh programs have partly listed vectors.
--probe-meshfb passes (fallback pictures equal the mesh pipeline's), and its module part converts all 520 new
translations with 0 failures, all 1,020 modules pass spirv-val and their compute pipelines build.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
On a GPU without mesh shaders the fallback replaces the translated module's entry function with its own, which cleared
only the primitive count and then copied the whole slot (vertex lines, count, index bytes) out of workgroup memory to
be drawn. Lines and index bytes the program did not write kept what an earlier workgroup left there and were drawn as
vertices and primitives. The fallback's entry function now sets every word of the slot to zero in a loop over the
workgroup's threads before the payload copy and the barrier, as the translator's own entry function does.
Checked: --probe-meshfb passes before and after (the fallback's pictures match the mesh pipeline's texel for texel on
an RTX GPU); its module part converts all 520 task and mesh programs of an installed shader cache with 0 failures, all
1,020 resulting modules pass spirv-val (vulkan1.3) and their 520 compute pipelines build on the device.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Vulkan leaves workgroup memory undefined, and a mesh program need not write every vertex it has room for or every
index byte. The mesh entry function copies all of them into the outputs, so what an earlier workgroup left in shared
memory became vertices and primitives: on AMD, long thin triangles across the foliage's shadow map. The entry function
now sets the vertex lines, the primitive count and the index bytes to zero in a loop over the workgroup's threads
before the payload copy and the barrier. Task programs keep clearing only their mesh count.
Checked by translating all 520 task and mesh programs an installed shader cache stores with the translator before and
after the change: every one translates and passes spirv-val (vulkan1.3) both times; each mesh module grows by about
88 bytes, task modules are unchanged. tools/xlate_smoke.py passes.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
