MagicBeanDev/Hookshot
Archived
Watch
6
0
Fork
You've already forked Hookshot
 0
Code Issues 125 Pull requests 205 Activity

Working fallback mesh shaders (AMD RDNA1) and various performance and accuracy fixes #1023

Closed
rakefire-impala wants to merge 18 commits from working_fallback_mesh_shaders_RDNA1 into main
pull from: working_fallback_mesh_shaders_RDNA1
merge into: MagicBeanDev:main
Conversation 7 Commits 18 Files changed 34 +2523 -89
rakefire-impala commented 2026-11-09 06:41:58 +00:00 (Migrated from github.com)

Mesh shader fallback for GPUs without mesh shaders (RDNA1 and older), and fixes for AMD

This brings back #640 and makes it work. On an AMD RX 5500M (RDNA1, no VK_EXT_mesh_shader), the foliage, the floor grass, the ivy and the door of Link's house now draw. It also fixes what is likely the cause of the RTX 4060 Ti device loss that got #640 reverted (806c566), two AMD rendering bugs, and several texture problems on low-VRAM cards.

The full write-up, covering what was broken, how it was fixed and how it behaves, is in docs/mesh_fallback_rdna1.md.

Important

Updated since the first version (10 new commits, 9-18 below):

  • Pausing and resuming after alt-tab is fast again. While the game was paused or in the background, a texture budget that kept shrinking evicted about 1,000 textures (2,851 MB down to 1,137 MB) in a few seconds, and unpausing had to make them all again. That no longer happens.
  • No more black glyphs instead of the hearts, and no broken minimap view cone on 4 GB cards. Same cause as above.
  • The "radioactive" glowing stems and flowers are fixed. This was a shader translator bug on every GPU, not the fallback. Upstream fixed it in PR #759, which is cherry-picked here.
  • The black stepped patch on the Deku Tree's floor is gone, and walking stutters less, from two upstream v2.0.6 AMD texture fixes, cherry-picked.
  • New debugging tools: a live switch file to hide, list, read back or inspect fallback draws, with no restart.

Screenshots

Both are from an AMD RX 5500M laptop (4 GB of VRAM, no mesh shaders), with the fallback on Auto, 50% resolution, the handheld profile and frame generation 2x.

Kokiri Forest: the foliage, ground cover, trees and textures all draw, and the minimap is correct.

Image

House: the floor grass, the ivy, the door and the lighting are correct. Without the fallback, the floor here is flat moss with no grass and the door does not render. On this card, before the AMD fixes, the lighting came out dark and a basket glowed white.

Image

Status

Part Status
Foliage, floor grass, ivy, vines, the door of Link's house Working on the test laptop
Lighting and bloom on AMD Working on the test laptop
GPUs with mesh shaders (the #640 device loss) Fixed in the code, not tested on such a GPU
Texture noise after alt-tab on low-VRAM cards Working, but further testing needed
New: slow resume after pause or alt-tab, black hearts, broken minimap cone (4 GB cards) Working, but further testing needed
New: glowing stems and flowers in Kokiri Forest Fixed by upstream PR #759 (a translator bug on every GPU)
New: black patch on the Deku Tree's floor, terrain upload stutters on AMD Fixed by upstream v2.0.6 commits

Tested on: a low-VRAM laptop

All of this was tested on one laptop: an AMD Radeon RX 5500M (RDNA1) with 4 GB of VRAM, on Windows 11 with AMD's driver. Vulkan lets Hookshot use only about 3.1 to 3.4 GB of that, which is well below the 8 GB the readme recommends. Several of the bugs below only showed up because of the low VRAM: the lock-up, the white distant foliage, the texture noise, the slow resume and the black hearts. The Minimum (4GB) preset is the one to use on a card like this.

These frame rates come from the logs. They are not a controlled benchmark.

State Frame rate GPU time per frame
Docked settings, before the grass fix (Link's house) ~15 fps ~54 ms
Minimum (4GB) preset (Link's house) ~19 fps ~41 ms
Minimum preset, after the fallback's later fixes (Link's house) ~34-38 fps ~17 ms
New: after the v2.0.6 texture fixes (inside the Deku Tree, so not the same spot) ~44 fps not measured

How it behaves

  • It is picked at every start. If the GPU has mesh shaders, the game draws as before and nothing of the fallback is built or cached. If it has none, the fallback turns on by itself, provided the driver has buffer device addresses (Vulkan 1.2).
  • There is a toggle: Insert menu → Graphics → Mesh shader fallback.
    • Auto (the default) uses the fallback only on GPUs without mesh shaders.
    • Off never uses it.
    • On uses it on every GPU and does not turn on the GPU's own mesh shaders.
    • A change applies after a restart. The setting is not part of the Quality preset, so changing it keeps the preset's VRAM caps. RX_VK_NO_MESH_FALLBACK=1 still works and has the same effect as Off.
  • On is for AMD cards that have mesh shaders but hang with them. Lichii's #814 found that native task and mesh draws hang AMD RX 6000 and newer cards. On Windows that shows as a driver timeout (TDR), and on Linux as GPU resets with RADV. Lichii therefore uses its fallback by default on AMD. This PR keeps Auto as the default, and players with those cards can switch to On. The hang is Lichii's finding and has not been confirmed in Hookshot yet.
  • Every fallback program is compiled at startup, so nothing pops in white during play. The third "Preparing shaders" count, about 11,000 on the test laptop, now includes about 1,350 fallback programs. AMD can take 10 to 12 seconds to compile one, so the first start after a translator change or a driver update takes a few minutes. Later starts load them from the driver's cache in seconds.
  • The launch warning now only appears when it's true: on a GPU with no mesh shaders where the fallback can't run either. It explains why: either the fallback was turned off, and it says how to turn it back on, or the driver lacks buffer device addresses. It no longer says Hookshot needs an RTX 20 / RX 6000 / Arc card.

What was broken, and the fixes

#640's shader translation

These were found with a new offline validator, scripts/validate_spirv.py. At first, almost every one of the ~1,350 fallback programs failed spirv-val. Now all 1,347 pass.

  • Payload address: the payload's address was read from the push constants using its byte offset (64) as a member index, so every task-driven program read and wrote its payload at a garbage address.
  • Missing Aligned: device-address loads and stores had no Aligned, which Vulkan requires.
  • Partial payload copy: each task thread copied only 1 word in 32 of its payload, because the copy loop stepped by thread.
  • Wrong mesh workgroup index: a mesh workgroup's index within its task was its global slot (up to 13,855) instead of 0–31, so it picked a meshlet past the end of its task's list.
  • Invalid SPIR-V: a loop counter was declared as a Function variable outside the first block, and a bool was stored into a u32.

#640's draw side

  • The freeze: a compute pipeline the driver refused was rebuilt on every draw, and the stages built before it leaked. AMD also took seconds to compile each pipeline inside the draw. Pipelines now compile on a worker thread, and the draw is skipped until they are ready.
  • The lock-up, and the white distant foliage: every program had its own buffers on every recording thread, about 2.2 GB for 95 programs. When VRAM ran out, new buffers could not be made, and those programs stopped drawing for good. Foliage that lost its colour pass but kept its depth pass came out white. All of a thread's fallback draws now share one 72 MB pair of buffers, with a barrier between draws.
  • Missing floor grass, part 1: a draw had only 1,024 slots, each sized for its program's maximum, so about 93% of the grass was dropped. There are now 16,384 slots. A mesh workgroup claims records only for the triangles it outputs, and the draw is indirect, sized from what was claimed. That also removed the GPU's work on empty vertices.
  • Missing floor grass, part 2: the grass is drawn with DrawMeshTasksIndirect, where the task count lives in GPU memory. The mesh stage was launched with the CPU's copy of that count, which is 0, so no mesh workgroup ever ran for indirect task draws. Task threads now count the mesh work on the GPU, and the mesh stage is launched from that count.
  • Why the door and ivy appeared before the grass: they started drawing once the payload and index fixes above were in, but the grass is an indirect draw, so it needed the fix just above as well.
  • The door was white at spawn while its colour pass waited for slow compiles. The fix is the startup precompile described above.
  • New, hardening: the generated vertex stage now starts every output at zero, as the console's stage buffer does, instead of leaving unwritten components undefined. Fragment-stage butterfly shuffles now take the same AMD workaround as the other shuffles. Neither was the cause of the glowing foliage below.

GPUs with mesh shaders: the #640 regression

  • #640 changed build_graphics (which builds every graphics pipeline) to drop all TASK and MESH stages for every program. Fallback stages are compute stages, so in practice this only ever hit native mesh programs: on a GPU with mesh shaders, every native mesh pipeline was built with only a fragment stage. That is very likely the RTX 4060 Ti device loss, which happened with the fallback off.
  • build_graphics now skips only a fallback program's compute stages.
  • The push-constant range is back to main's size when the fallback is off.
  • Vertex input is read from stage 0 again.
  • Native task and mesh shader code is byte for byte main's.

AMD lighting and bloom (not specific to the fallback)

  • 64-lane graphics warps: AMD only allows compute shaders to be pinned to 32 lanes, so fragment shaders run 64 lanes wide while the translator assumed 32. Graphics stages on such GPUs now use the translator's 64-lane path, and compute keeps the old one. Giving compute the 64-lane path too made shader preparation take hours. The program-cache stamp now includes these settings.
  • Lichii #273: fragment-shader shuffles on AMD's proprietary driver now go through the broadcast workaround (a basket glowed white). Depth-compare textures are never treated as integer images.

Texture noise after alt-tab (working, but further testing needed)

This one is a bug in Hookshot's own texture code, unrelated to the fallback:

  1. With the game in the background, about a thousand textures got evicted.
  2. Re-making them requires the BC3 re-encode scratch, whose output buffer could only live in VRAM, and VRAM was full.
  3. Every re-make failed, but each texture was marked as uploaded anyway, so the empty BC3 images were drawn as noise and never retried.

Now the buffer falls back to system memory, and a failed BC3 upload is retried 30 frames later. Each failure is logged as vk: BC3 texture ... its upload failed (...). This was only tested briefly.

New: slow resume after pause or alt-tab, black hearts, broken minimap cone (working, further testing needed)

Pausing and resuming, including after the game loses focus, is now fast. Before, unpausing could hang for a long time on a 4 GB card.

The cause was a texture budget that ratcheted down:

  1. Every 10 frames, when the driver said Hookshot was above 97% of its allowance, the texture budget was set below what was held, and textures unused for 60 frames (about a second) were evicted.
  2. Evicting frees space inside the memory blocks the allocator already holds, but the driver still counts those blocks. So its count stayed at the limit (3,337 of 3,337 MB, with 1,224 MB free inside the blocks), and 10 frames later the budget shrank again.
  3. While paused, nothing in the world is drawn, so its textures drained from 2,851 MB to 1,137 MB in a few seconds. Unpausing had to make about 1,000 of them again: ASTC decode, BC3 re-encode and upload.

The same squeeze broke the HUD. The hearts and the minimap's view cone are small textures the GPU draws once and the game samples again about 60 frames later. Their pixels exist only in the GPU image. Once evicted, they came back over memory that never got those pixels, as black glyphs, until an area change made the game redraw them.

Fix: the 97% squeeze no longer runs while the allocator's blocks hold at least an eighth of the allowance free. New textures go into that space without the driver's count growing. The other rule, which judges pressure after taking that free space off, still evicts when memory is really short. The VRAM log line says not squeezed when this applies, and RX_VK_SQUEEZE_ALWAYS=1 restores the old behaviour. The log also now names each small GPU-drawn texture evicted with its pixels lost.

New: glowing "radioactive" stems and flowers (fixed by upstream PR #759)

In Kokiri Forest, stems glowed cyan, petals bloomed white and bushes had coloured streaks. Upstream had the same report on GPUs with mesh shaders (#757 on an RTX 5080, #874, #986), so this was not the fallback.

  • How it was found: the new readback tool showed the geometry was right. The depth prepass and the colour pass had bit-identical positions and triangles in the same frame, with no NaNs. The stems' colour pass computes light-probe (spherical-harmonics) lighting per vertex, and about a third of the stems came out at hundreds to thousands, some negative.
  • The cause: the translator renames high registers through temporary markers, and a marker could take the number of an operand the same FFMA still read, so the math used the wrong value.
  • The fix: upstream PR #759 (deaf4ff), cherry-picked with its author kept. Its regression test fails on this branch's base without the fix and passes with it. In game, the stems and flowers look right.

New: black patch on the Deku Tree's floor, and terrain stutters on AMD (fixed by upstream v2.0.6 commits)

Inside the Deku Tree, a large stepped region of the floor sometimes turned black. It stayed with every fallback draw hidden, so the fallback doesn't draw it. This branch still had two AMD texture paths that upstream fixed in v2.0.6 (tracking issue #1006):

  • The uneven-BC3 resize (#705) re-encoded the terrain's 1365×1365×90 material arrays at 1536×1536. Upstream measured 7× the GPU time on an RX 6700 XT with it. Here each of those 57 MB arrays took 80-93 ms to upload, a stutter every time.
  • The BC3 re-encode of ASTC arrays, several layers per run, left dark blocks on terrain on AMD.

Both upstream commits are cherry-picked: dc2f4dc (the resize is opt-in again, RX_VK_UNEVEN_BC=0 turns it on) and ec89f92 (#956, one layer per run). In game, the black patch did not come back, walking is smoother, the frame rate rose (table above), and VRAM is no longer pinned at the limit. The thin grid lines at terrain seams that upstream names as the trade-off were not seen.

New: debugging tools

A live switch file, hookshot_fallback_debug.txt in the game folder, is read again within half a second of a change, with no restart and no shader preparation:

  • off hides every fallback draw, and a key's hex prefix hides that program.
  • list logs each fallback program drawn.
  • read <prefix> copies a program's fallback records back to the CPU, all in one frame, then logs each value's range, NaNs and infinities, and saves them to a .bin file.
  • info <prefix> logs a draw's render targets and formats, write masks, blending and depth test.

RX_VK_WARP64=0|1 and RX_VK_FRAG_SHUFFLE_FALLBACK=0|1 force the AMD warp settings for A/B tests.

Commits

  1. 5918e80 Re-apply PR #640 (undoes 806c566 and 449382e)
  2. dc61d57 tools: offline SPIR-V validator (scripts/validate_spirv.py)
  3. b17e373 AMD: 64-lane graphics warps, Lichii #273
  4. 6b5e629 vk: GPUs with mesh shaders draw as main again (#640's broken native mesh pipelines)
  5. cb0d3ef menu: Mesh shader fallback (Auto / Off / On), and the launch warning text
  6. 4e36b5d mesh fallback: shader translation, draw side, memory, floor grass, startup precompile
  7. 353bd4a textures: retry failed BC3 uploads, with a system-memory fallback for the scratch buffer (working, further testing needed)
  8. 214096b docs: docs/mesh_fallback_rdna1.md and readme updates

New in this update:

  1. 5c2004f mesh fallback: zeroed vertex-stage outputs, butterfly shuffles through the AMD workaround, the live debug switch
  2. 55b7ee5 docs: glowing stems and flowers, as a known issue at the time
  3. 9218be7 mesh fallback debug: read (records readback) and info (draw state) in the live switch
  4. 46e8c3a cherry-pick of upstream PR #759 (deaf4ff): register markers never alias an original operand (glowing stems)
  5. f8c00a4 docs: the cause of the glowing stems and flowers (item 11)
  6. 5ff934f textures: no driver's-limit squeeze while the allocator's blocks hold room (fast resume, HUD fixed); log lost GPU-drawn textures
  7. 2709765 docs: the squeeze fix (item 12)
  8. a35cbf7 cherry-pick of upstream ec89f92 (#956): BC3 re-encode of ASTC arrays one layer per run
  9. a34f6f6 cherry-pick of upstream dc2f4dc: the AMD uneven-BC3 resize (#705) opt-in again
  10. 16229ef docs: the Deku Tree black patch and terrain stutters (item 13)

Testing

  • Played on the RX 5500M:
    • the foliage, floor grass, ivy, vines and door draw, the lighting and bloom are correct, there's no lock-up, and alt-tabbing no longer leaves noise
    • new: pausing and resuming, including after alt-tab, is quick; the hearts and the minimap cone stay correct
    • new: the Kokiri stems and flowers no longer glow; the Deku Tree's floor patch did not come back; walking stutters less
  • spirv-val: 1,347 fallback programs, 0 failures (run with --uniform-buffer-standard-layout --scalar-block-layout).
  • Tests:
    • new or extended: tests/test_mesh_fallback.py, tests/test_amd_fragment_warp.py, tests/test_bc_upload_retry.py (now also checks the squeeze rule), tests/test_validate_spirv.py
    • new, from upstream #759: tests/test_lower_register_markers.py, 89 cases, 0 failures. It fails on this branch's base without the fix.
    • also passing: test_mesh_point_topology, test_mesh_stage_buffer_zero, test_program_cache_reader, test_bc_levels, nx2/test_mesh_dynamic_state and the menu tests
    • the stale #640 check in test_mesh_fallback.py was removed; it was passing only because the same text appears on an unrelated line
  • Per commit: each commit's tests were run against that commit's own state. Only the final states were compiled and played. test_fps_overlay_order fails the same way on main: its fixture is missing <initializer_list>.

Needs testing / review notes

  • A GPU with mesh shaders, ideally an RTX card: please confirm the device loss is gone and that nothing else changed.
  • Other GPUs without mesh shaders (RX 5000, RX 400/500, GTX 10 series) are untested.
  • AMD RX 6000 or newer: try Mesh shader fallback set to On if they hang.
  • The texture fixes need more play time, and testing on other low-VRAM cards. With the squeeze fix the driver may page VRAM to system memory instead of Hookshot evicting; that wasn't seen here, but it's the risk.
  • #759 on AMD: upstream has an open report of green diagonal stripes on an RX 7600 that may come from #759 (its A/B switch, RX_VK_MARKERS_V205=1, is on main, not on this branch). None were seen on the RX 5500M.
  • Rebase: the branch is based on 232eecb, and main has moved on since, so it will need a rebase. Commits 12, 16 and 17 are cherry-picks of commits already on main, so a rebase should drop them as duplicates.

Known limits

  • Draw order: semi-transparent mesh geometry can occasionally draw out of order, because records land in the order the GPU claims them. Opaque foliage isn't affected. Lichii's fallback has the same limit.
  • Only one GPU tested so far.
  • One big hitch remains on 4 GB cards: the 12288×12288 shadow map (288 MB) was once made again and uploaded whole, taking about 2 seconds, probably on an area change. Other first-use texture uploads can still take up to ~100 ms.

Credits

#640 (shayder) for the fallback this builds on. Lichii (zortos293 and contributors) for the fallback design (#569), the AMD fixes (#273), the AMD hang finding (#814), and their notes on how the game draws its foliage. Upstream PR #759 for the register-marker fix, YunaraDebugs for #956, and NotRexed for making the AMD resize opt-in. The reporters of #757, #874, #986 and #1006, whose reports pointed to the upstream fixes.

# Mesh shader fallback for GPUs without mesh shaders (RDNA1 and older), and fixes for AMD This brings back #640 and makes it work. On an AMD RX 5500M (RDNA1, no `VK_EXT_mesh_shader`), the foliage, the floor grass, the ivy and the door of Link's house now draw. It also fixes what is likely the cause of the RTX 4060 Ti device loss that got #640 reverted (`806c566`), two AMD rendering bugs, and several texture problems on low-VRAM cards. The full write-up, covering what was broken, how it was fixed and how it behaves, is in [`docs/mesh_fallback_rdna1.md`](https://github.com/rakefire-impala/Hookshot-w-Mesh-Shader-Fallback/blob/working_fallback_mesh_shaders_RDNA1/docs/mesh_fallback_rdna1.md). > [!IMPORTANT] > **Updated since the first version (10 new commits, 9-18 below):** > - **Pausing and resuming after alt-tab is fast again.** While the game was paused or in the background, a texture budget that kept shrinking evicted about 1,000 textures (2,851 MB down to 1,137 MB) in a few seconds, and unpausing had to make them all again. That no longer happens. > - **No more black glyphs instead of the hearts, and no broken minimap view cone** on 4 GB cards. Same cause as above. > - **The "radioactive" glowing stems and flowers are fixed.** This was a shader translator bug on every GPU, not the fallback. Upstream fixed it in PR #759, which is cherry-picked here. > - **The black stepped patch on the Deku Tree's floor is gone, and walking stutters less**, from two upstream v2.0.6 AMD texture fixes, cherry-picked. > - **New debugging tools:** a live switch file to hide, list, read back or inspect fallback draws, with no restart. ## Screenshots Both are from an AMD RX 5500M laptop (4 GB of VRAM, no mesh shaders), with the fallback on Auto, 50% resolution, the handheld profile and frame generation 2x. **Kokiri Forest:** the foliage, ground cover, trees and textures all draw, and the minimap is correct. <img width="1917" height="1077" alt="Image" src="https://github.com/user-attachments/assets/b0428db4-ca40-4aeb-be5a-8da310eaafee" /> **House:** the floor grass, the ivy, the door and the lighting are correct. Without the fallback, the floor here is flat moss with no grass and the door does not render. On this card, before the AMD fixes, the lighting came out dark and a basket glowed white. <img width="1917" height="1071" alt="Image" src="https://github.com/user-attachments/assets/91170287-1ee9-40e8-8522-28416380b196" /> ## Status | Part | Status | |---|---| | Foliage, floor grass, ivy, vines, the door of Link's house | Working on the test laptop | | Lighting and bloom on AMD | Working on the test laptop | | GPUs **with** mesh shaders (the #640 device loss) | Fixed in the code, **not tested** on such a GPU | | Texture noise after alt-tab on low-VRAM cards | **Working, but further testing needed** | | **New:** slow resume after pause or alt-tab, black hearts, broken minimap cone (4 GB cards) | **Working, but further testing needed** | | **New:** glowing stems and flowers in Kokiri Forest | **Fixed** by upstream PR #759 (a translator bug on every GPU) | | **New:** black patch on the Deku Tree's floor, terrain upload stutters on AMD | **Fixed** by upstream v2.0.6 commits | ## Tested on: a low-VRAM laptop All of this was tested on one laptop: an **AMD Radeon RX 5500M (RDNA1) with 4 GB of VRAM**, on Windows 11 with AMD's driver. Vulkan lets Hookshot use only about 3.1 to 3.4 GB of that, which is well below the 8 GB the readme recommends. Several of the bugs below only showed up because of the low VRAM: the lock-up, the white distant foliage, the texture noise, the slow resume and the black hearts. The **Minimum (4GB)** preset is the one to use on a card like this. These frame rates come from the logs. They are not a controlled benchmark. | State | Frame rate | GPU time per frame | |---|---|---| | Docked settings, before the grass fix (Link's house) | ~15 fps | ~54 ms | | Minimum (4GB) preset (Link's house) | ~19 fps | ~41 ms | | Minimum preset, after the fallback's later fixes (Link's house) | ~34-38 fps | ~17 ms | | **New:** after the v2.0.6 texture fixes (inside the Deku Tree, so not the same spot) | **~44 fps** | not measured | ## How it behaves - **It is picked at every start.** If the GPU has mesh shaders, the game draws as before and nothing of the fallback is built or cached. If it has none, the fallback turns on by itself, provided the driver has buffer device addresses (Vulkan 1.2). - **There is a toggle:** Insert menu → Graphics → **Mesh shader fallback**. - **Auto** (the default) uses the fallback only on GPUs without mesh shaders. - **Off** never uses it. - **On** uses it on every GPU and does not turn on the GPU's own mesh shaders. - A change applies after a restart. The setting is not part of the Quality preset, so changing it keeps the preset's VRAM caps. `RX_VK_NO_MESH_FALLBACK=1` still works and has the same effect as Off. - **On is for AMD cards that have mesh shaders but hang with them.** Lichii's [#814](https://github.com/zortos293/Lichii/pull/814) found that native task and mesh draws hang AMD RX 6000 and newer cards. On Windows that shows as a driver timeout (TDR), and on Linux as GPU resets with RADV. Lichii therefore uses its fallback by default on AMD. This PR keeps Auto as the default, and players with those cards can switch to On. The hang is Lichii's finding and has not been confirmed in Hookshot yet. - **Every fallback program is compiled at startup, so nothing pops in white during play.** The third "Preparing shaders" count, about 11,000 on the test laptop, now includes about 1,350 fallback programs. AMD can take 10 to 12 seconds to compile one, so the first start after a translator change or a driver update takes a few minutes. Later starts load them from the driver's cache in seconds. - **The launch warning now only appears when it's true:** on a GPU with no mesh shaders where the fallback can't run either. It explains why: either the fallback was turned off, and it says how to turn it back on, or the driver lacks buffer device addresses. It no longer says Hookshot needs an RTX 20 / RX 6000 / Arc card. ## What was broken, and the fixes ### #640's shader translation These were found with a new offline validator, `scripts/validate_spirv.py`. At first, almost every one of the ~1,350 fallback programs failed `spirv-val`. Now **all 1,347 pass**. - **Payload address:** the payload's address was read from the push constants using its byte offset (64) as a member index, so every task-driven program read and wrote its payload at a garbage address. - **Missing `Aligned`:** device-address loads and stores had no `Aligned`, which Vulkan requires. - **Partial payload copy:** each task thread copied only 1 word in 32 of its payload, because the copy loop stepped by thread. - **Wrong mesh workgroup index:** a mesh workgroup's index within its task was its global slot (up to 13,855) instead of 0–31, so it picked a meshlet past the end of its task's list. - **Invalid SPIR-V:** a loop counter was declared as a Function variable outside the first block, and a `bool` was stored into a `u32`. ### #640's draw side - **The freeze:** a compute pipeline the driver refused was rebuilt on every draw, and the stages built before it leaked. AMD also took seconds to compile each pipeline inside the draw. Pipelines now compile on a worker thread, and the draw is skipped until they are ready. - **The lock-up, and the white distant foliage:** every program had its own buffers on every recording thread, about **2.2 GB for 95 programs**. When VRAM ran out, new buffers could not be made, and those programs stopped drawing for good. Foliage that lost its colour pass but kept its depth pass came out white. All of a thread's fallback draws now share one 72 MB pair of buffers, with a barrier between draws. - **Missing floor grass, part 1:** a draw had only 1,024 slots, each sized for its program's maximum, so about 93% of the grass was dropped. There are now 16,384 slots. A mesh workgroup claims records only for the triangles it outputs, and the draw is indirect, sized from what was claimed. That also removed the GPU's work on empty vertices. - **Missing floor grass, part 2:** the grass is drawn with `DrawMeshTasksIndirect`, where the task count lives in GPU memory. The mesh stage was launched with the CPU's copy of that count, which is 0, so **no mesh workgroup ever ran for indirect task draws**. Task threads now count the mesh work on the GPU, and the mesh stage is launched from that count. - **Why the door and ivy appeared before the grass:** they started drawing once the payload and index fixes above were in, but the grass is an indirect draw, so it needed the fix just above as well. - **The door was white at spawn** while its colour pass waited for slow compiles. The fix is the startup precompile described above. - **New, hardening:** the generated vertex stage now starts every output at zero, as the console's stage buffer does, instead of leaving unwritten components undefined. Fragment-stage butterfly shuffles now take the same AMD workaround as the other shuffles. Neither was the cause of the glowing foliage below. ### GPUs with mesh shaders: the #640 regression - #640 changed `build_graphics` (which builds every graphics pipeline) to drop all **TASK and MESH** stages for every program. Fallback stages are compute stages, so in practice this only ever hit **native** mesh programs: on a GPU with mesh shaders, every native mesh pipeline was built with only a fragment stage. That is very likely the RTX 4060 Ti device loss, which happened with the fallback off. - `build_graphics` now skips only a fallback program's compute stages. - The push-constant range is back to `main`'s size when the fallback is off. - Vertex input is read from stage 0 again. - Native task and mesh shader code is byte for byte `main`'s. ### AMD lighting and bloom (not specific to the fallback) - **64-lane graphics warps:** AMD only allows compute shaders to be pinned to 32 lanes, so fragment shaders run 64 lanes wide while the translator assumed 32. Graphics stages on such GPUs now use the translator's 64-lane path, and compute keeps the old one. Giving compute the 64-lane path too made shader preparation take hours. The program-cache stamp now includes these settings. - **[Lichii #273](https://github.com/zortos293/Lichii/pull/273):** fragment-shader shuffles on AMD's proprietary driver now go through the broadcast workaround (a basket glowed white). Depth-compare textures are never treated as integer images. ### Texture noise after alt-tab (working, but further testing needed) This one is a bug in Hookshot's own texture code, unrelated to the fallback: 1. With the game in the background, about a thousand textures got evicted. 2. Re-making them requires the BC3 re-encode scratch, whose output buffer could only live in VRAM, and VRAM was full. 3. Every re-make failed, but each texture was marked as uploaded anyway, so the empty BC3 images were drawn as noise and never retried. Now the buffer falls back to system memory, and a failed BC3 upload is retried 30 frames later. Each failure is logged as `vk: BC3 texture ... its upload failed (...)`. This was only tested briefly. ### New: slow resume after pause or alt-tab, black hearts, broken minimap cone (working, further testing needed) **Pausing and resuming, including after the game loses focus, is now fast.** Before, unpausing could hang for a long time on a 4 GB card. **The cause was a texture budget that ratcheted down:** 1. Every 10 frames, when the driver said Hookshot was above 97% of its allowance, the texture budget was set below what was held, and textures unused for 60 frames (about a second) were evicted. 2. Evicting frees space *inside* the memory blocks the allocator already holds, but the driver still counts those blocks. So its count stayed at the limit (3,337 of 3,337 MB, with 1,224 MB free inside the blocks), and 10 frames later the budget shrank again. 3. While paused, nothing in the world is drawn, so its textures drained from 2,851 MB to 1,137 MB in a few seconds. Unpausing had to make about 1,000 of them again: ASTC decode, BC3 re-encode and upload. **The same squeeze broke the HUD.** The hearts and the minimap's view cone are small textures the GPU draws once and the game samples again about 60 frames later. Their pixels exist only in the GPU image. Once evicted, they came back over memory that never got those pixels, as black glyphs, until an area change made the game redraw them. **Fix:** the 97% squeeze no longer runs while the allocator's blocks hold at least an eighth of the allowance free. New textures go into that space without the driver's count growing. The other rule, which judges pressure after taking that free space off, still evicts when memory is really short. The VRAM log line says `not squeezed` when this applies, and `RX_VK_SQUEEZE_ALWAYS=1` restores the old behaviour. The log also now names each small GPU-drawn texture evicted with its pixels lost. ### New: glowing "radioactive" stems and flowers (fixed by upstream PR #759) In Kokiri Forest, stems glowed cyan, petals bloomed white and bushes had coloured streaks. Upstream had the same report on GPUs with mesh shaders (#757 on an RTX 5080, #874, #986), so this was not the fallback. - **How it was found:** the new readback tool showed the geometry was right. The depth prepass and the colour pass had bit-identical positions and triangles in the same frame, with no NaNs. The stems' colour pass computes light-probe (spherical-harmonics) lighting per vertex, and about a third of the stems came out at hundreds to thousands, some negative. - **The cause:** the translator renames high registers through temporary markers, and a marker could take the number of an operand the same FFMA still read, so the math used the wrong value. - **The fix:** upstream PR #759 (`deaf4ff`), cherry-picked with its author kept. Its regression test fails on this branch's base without the fix and passes with it. In game, the stems and flowers look right. ### New: black patch on the Deku Tree's floor, and terrain stutters on AMD (fixed by upstream v2.0.6 commits) Inside the Deku Tree, a large stepped region of the floor sometimes turned black. It stayed with every fallback draw hidden, so the fallback doesn't draw it. This branch still had two AMD texture paths that upstream fixed in v2.0.6 (tracking issue #1006): - **The uneven-BC3 resize (#705)** re-encoded the terrain's 1365×1365×90 material arrays at 1536×1536. Upstream measured 7× the GPU time on an RX 6700 XT with it. Here each of those 57 MB arrays took 80-93 ms to upload, a stutter every time. - **The BC3 re-encode of ASTC arrays, several layers per run,** left dark blocks on terrain on AMD. Both upstream commits are cherry-picked: `dc2f4dc` (the resize is opt-in again, `RX_VK_UNEVEN_BC=0` turns it on) and `ec89f92` (#956, one layer per run). In game, the black patch did not come back, walking is smoother, the frame rate rose (table above), and VRAM is no longer pinned at the limit. The thin grid lines at terrain seams that upstream names as the trade-off were not seen. ### New: debugging tools A live switch file, `hookshot_fallback_debug.txt` in the game folder, is read again within half a second of a change, with no restart and no shader preparation: - `off` hides every fallback draw, and a key's hex prefix hides that program. - `list` logs each fallback program drawn. - `read <prefix>` copies a program's fallback records back to the CPU, all in one frame, then logs each value's range, NaNs and infinities, and saves them to a `.bin` file. - `info <prefix>` logs a draw's render targets and formats, write masks, blending and depth test. `RX_VK_WARP64=0|1` and `RX_VK_FRAG_SHUFFLE_FALLBACK=0|1` force the AMD warp settings for A/B tests. ## Commits 1. `5918e80` Re-apply PR #640 (undoes `806c566` and `449382e`) 2. `dc61d57` tools: offline SPIR-V validator (`scripts/validate_spirv.py`) 3. `b17e373` AMD: 64-lane graphics warps, Lichii #273 4. `6b5e629` vk: GPUs with mesh shaders draw as `main` again (#640's broken native mesh pipelines) 5. `cb0d3ef` menu: Mesh shader fallback (Auto / Off / On), and the launch warning text 6. `4e36b5d` mesh fallback: shader translation, draw side, memory, floor grass, startup precompile 7. `353bd4a` textures: retry failed BC3 uploads, with a system-memory fallback for the scratch buffer (working, further testing needed) 8. `214096b` docs: `docs/mesh_fallback_rdna1.md` and readme updates **New in this update:** 9. `5c2004f` mesh fallback: zeroed vertex-stage outputs, butterfly shuffles through the AMD workaround, the live debug switch 10. `55b7ee5` docs: glowing stems and flowers, as a known issue at the time 11. `9218be7` mesh fallback debug: `read` (records readback) and `info` (draw state) in the live switch 12. `46e8c3a` cherry-pick of upstream PR #759 (`deaf4ff`): register markers never alias an original operand (glowing stems) 13. `f8c00a4` docs: the cause of the glowing stems and flowers (item 11) 14. `5ff934f` **textures: no driver's-limit squeeze while the allocator's blocks hold room (fast resume, HUD fixed); log lost GPU-drawn textures** 15. `2709765` docs: the squeeze fix (item 12) 16. `a35cbf7` cherry-pick of upstream `ec89f92` (#956): BC3 re-encode of ASTC arrays one layer per run 17. `a34f6f6` cherry-pick of upstream `dc2f4dc`: the AMD uneven-BC3 resize (#705) opt-in again 18. `16229ef` docs: the Deku Tree black patch and terrain stutters (item 13) ## Testing - **Played on the RX 5500M:** - the foliage, floor grass, ivy, vines and door draw, the lighting and bloom are correct, there's no lock-up, and alt-tabbing no longer leaves noise - **new:** pausing and resuming, including after alt-tab, is quick; the hearts and the minimap cone stay correct - **new:** the Kokiri stems and flowers no longer glow; the Deku Tree's floor patch did not come back; walking stutters less - **`spirv-val`:** 1,347 fallback programs, 0 failures (run with `--uniform-buffer-standard-layout --scalar-block-layout`). - **Tests:** - new or extended: `tests/test_mesh_fallback.py`, `tests/test_amd_fragment_warp.py`, `tests/test_bc_upload_retry.py` (now also checks the squeeze rule), `tests/test_validate_spirv.py` - **new, from upstream #759:** `tests/test_lower_register_markers.py`, 89 cases, 0 failures. It fails on this branch's base without the fix. - also passing: `test_mesh_point_topology`, `test_mesh_stage_buffer_zero`, `test_program_cache_reader`, `test_bc_levels`, `nx2/test_mesh_dynamic_state` and the menu tests - the stale #640 check in `test_mesh_fallback.py` was removed; it was passing only because the same text appears on an unrelated line - **Per commit:** each commit's tests were run against that commit's own state. Only the final states were compiled and played. `test_fps_overlay_order` fails the same way on `main`: its fixture is missing `<initializer_list>`. ## Needs testing / review notes - **A GPU with mesh shaders, ideally an RTX card:** please confirm the device loss is gone and that nothing else changed. - **Other GPUs without mesh shaders** (RX 5000, RX 400/500, GTX 10 series) are untested. - **AMD RX 6000 or newer:** try Mesh shader fallback set to On if they hang. - **The texture fixes** need more play time, and testing on other low-VRAM cards. With the squeeze fix the driver may page VRAM to system memory instead of Hookshot evicting; that wasn't seen here, but it's the risk. - **#759 on AMD:** upstream has an open report of green diagonal stripes on an RX 7600 that may come from #759 (its A/B switch, `RX_VK_MARKERS_V205=1`, is on `main`, not on this branch). None were seen on the RX 5500M. - **Rebase:** the branch is based on `232eecb`, and `main` has moved on since, so it will need a rebase. Commits 12, 16 and 17 are cherry-picks of commits already on `main`, so a rebase should drop them as duplicates. ## Known limits - **Draw order:** semi-transparent mesh geometry can occasionally draw out of order, because records land in the order the GPU claims them. Opaque foliage isn't affected. Lichii's fallback has the same limit. - **Only one GPU tested so far.** - **One big hitch remains on 4 GB cards:** the 12288×12288 shadow map (288 MB) was once made again and uploaded whole, taking about 2 seconds, probably on an area change. Other first-use texture uploads can still take up to ~100 ms. ## Credits #640 (shayder) for the fallback this builds on. [Lichii](https://github.com/zortos293/Lichii) (zortos293 and contributors) for the fallback design (#569), the AMD fixes (#273), the AMD hang finding (#814), and their notes on how the game draws its foliage. Upstream PR #759 for the register-marker fix, YunaraDebugs for #956, and NotRexed for making the AMD resize opt-in. The reporters of #757, #874, #986 and #1006, whose reports pointed to the upstream fixes.
👍 1
hckrman101 commented 2026-11-09 18:39:58 +00:00 (Migrated from github.com)

Hi! This PR has merge conflicts with main now (a lot has landed today). Could you merge main in or rebase and resolve them? Happy to give it a proper review once it's up to date.

Hi! This PR has merge conflicts with `main` now (a lot has landed today). Could you merge `main` in or rebase and resolve them? Happy to give it a proper review once it's up to date.
hckrman101 (Migrated from github.com) reviewed 2026-11-09 23:29:02 +00:00
hckrman101 (Migrated from github.com) left a comment

Thanks for this, and for the detailed write-up. The indirect-task fix and the shared, allocated records are real wins over what's on main. But main has moved a long way since 232eecb, and a good part of this branch is now already there. So here's what's duplicated, what's still new, and how I'd suggest landing the rest.

Already on main (drop these when you rebase)

  • Commit 1 (re-apply #640): landed as #779 (07773ab + aec016f). emit_spirv.cpp is byte-identical to main.
  • Commit 4 (build_graphics keeps native task/mesh stages): aec016f on main (p.fallback && (COMPUTE|TASK|MESH)), which does the same thing. The vertex_info early return is a no-op next to it.
  • From commit 6, the four fixes that came over in #1314 (from #792): the payload member index (it was the byte offset 64), Aligned 4 on PhysicalStorageBuffer accesses, the started word as a u32 / spirv-val-clean Strided (main unrolls the loop, you used a Private var; either is fine, keep main's), and the mesh workgroup's index within its task.
  • Commits 12, 16, 17 (#759, ec89f92, dc2f4dc): already on main, as you said. git cherry marks them as such. The depth-compare-is-never-integer half of commit 3 is on main too (46c5be8, #834).
  • Also on main and missing here: #1314's descriptor fix (dw().prog != &p in fallback_record and in same_desc). It has to survive the rebase: with the shared buffers, a dispatch between two fallback draws still clobbers the pushed compute descriptors.

Still new, and the part worth landing (one PR: "mesh fallback: indirect task draws, shared allocated records, async compile")

  • GPU-side mesh dispatch: task threads append their slot to a map and bump a VkDispatchIndirectCommand, and the mesh pass runs vkCmdDispatchIndirect from it. Main still dispatches groups * threads from the CPU, which is 0 for DrawMeshTasksIndirect, so the floor grass never draws on main. That's the most valuable piece here.
  • Records taken by an atomic allocator, drawn with vkCmdDrawIndirect, and capacity 1024 → 16384.
  • One shared records/payload pair per recording thread instead of one per program, plus the per-task payload copy (no longer one per slot).
  • Fallback compute pipelines built on a worker via compute_state, plus the precompile pass. Note that #1317 (still open, not merged) solves "a failed pipeline is retried every draw" another way. Your state-2-with-empty-fallback_compute covers that case too, so only one of them should land. Worth coordinating with CiLiBi18.
  • Zeroed vertex-stage outputs and the butterfly shuffle routing are small and fine to keep with this.

I read through the new SPIR-V and the dispatch carefully, and the core logic holds up:

  • The map indices stay below CAPACITY, since each started slot is unique and below it.
  • The allocator's "fits" check keeps the drawn records a prefix of the buffer, so vertexCount only covers written records, and taken is a multiple of corners.
  • The header fills don't overlap: draw command in words 0-3, allocator in word 4, dispatch in words 5-7.
  • The records buffer has INDIRECT_BUFFER_BIT.
  • full_barrier is ALL_COMMANDS / MEMORY_READ|WRITE, so it covers the indirect reads. vkCmdFillBuffer and vkCmdDispatchIndirect both bump work_epoch, so the dedupe can't drop the barrier after the fills.
  • The two extra workgroup barriers in MeshFallbackEpilogue sit at top level and correctly broadcast base through mesh_primitive_count.

Things to watch in the rebase are inline: main's EmitIsbeRead and runtime_info changes collide with yours.

Unrelated, so please split out (separate PRs, separate testing)

  1. AMD 64-lane graphics warps + the proprietary-driver fragment shuffle workaround (commit 3, plus the RX_VK_WARP64 / RX_VK_FRAG_SHUFFLE_FALLBACK switches). This changes the SPIR-V of every graphics program on every AMD GPU whose graphics stages can't be pinned, including RDNA2/3 and RADV/Steam Deck players who use native mesh shaders. It also changes the cache stamp, so they all re-translate on first start. It may well be right, but it needs its own PR and testing on RDNA2+ (both drivers), not just the 5500M.
  2. Texture changes (commits 7 and 14): the BC3 scratch output in system memory, the 30-frame BC3 retry, the "no squeeze while VMA blocks hold 1/8 free" rule, and the evicted-texels log. These are a VRAM-policy change with its own risk (the driver paging instead of us evicting, as you noted). That area also moved on main (3e66cf9 BC3 scratch kept, the v2.0.11 iGPU caps, #840 batched eviction writebacks merged today), so it will conflict anyway.
  3. The Mesh shader fallback menu row (Auto/Off/On) and the launch warning text. They're small and could go with the core PR, but please rebase them over main's current vk_main.cpp.
  4. scripts/validate_spirv.py + its test. It's handy, but it's tooling, and main now has RX_VK_DUMP_MESH_SPV for the same job. It can stand on its own.
  5. The live debug switch (hookshot_fallback_debug.txt, read/info). See the inline comment: I'd gate it behind an env var or leave it out of the core PR.
  6. Docs: 400 lines in docs/mesh_fallback_rdna1.md is a lot to keep in sync. Most of it is PR history. I'd keep a short "how the fallback works / switches" section and leave the narrative in the PR description.

Suggested path
Start a fresh branch from current main. Cherry-pick only the fallback parts of commits 4e36b5d/5c2004f (the dispatch map, allocator, indirect draw, shared buffers, async build + precompile, zeroed outputs), resolving against #1314's code as noted inline. Drop commits 1, 4, 12, 16, 17 and the parts of commit 6 that #1314 already covers. Then open the warp/shuffle change and the texture changes as two more PRs. That also makes the "needs a mesh-shader GPU to confirm" question easy: the core PR doesn't touch native mesh paths at all once commit 4 is dropped.

CI hasn't run here (the branch conflicts), so nothing to report there yet.

Thanks for this, and for the detailed write-up. The indirect-task fix and the shared, allocated records are real wins over what's on main. But main has moved a long way since `232eecb`, and a good part of this branch is now already there. So here's what's duplicated, what's still new, and how I'd suggest landing the rest. **Already on main (drop these when you rebase)** - Commit 1 (re-apply #640): landed as #779 (`07773ab` + `aec016f`). `emit_spirv.cpp` is byte-identical to main. - Commit 4 (build_graphics keeps native task/mesh stages): `aec016f` on main (`p.fallback && (COMPUTE|TASK|MESH)`), which does the same thing. The `vertex_info` early return is a no-op next to it. - From commit 6, the four fixes that came over in #1314 (from #792): the payload member index (it was the byte offset 64), `Aligned 4` on PhysicalStorageBuffer accesses, the started word as a u32 / spirv-val-clean `Strided` (main unrolls the loop, you used a Private var; either is fine, keep main's), and the mesh workgroup's index within its task. - Commits 12, 16, 17 (#759, `ec89f92`, `dc2f4dc`): already on main, as you said. `git cherry` marks them as such. The depth-compare-is-never-integer half of commit 3 is on main too (`46c5be8`, #834). - Also on main and missing here: #1314's descriptor fix (`dw().prog != &p` in `fallback_record` and in `same_desc`). It has to survive the rebase: with the shared buffers, a dispatch between two fallback draws still clobbers the pushed compute descriptors. **Still new, and the part worth landing (one PR: "mesh fallback: indirect task draws, shared allocated records, async compile")** - GPU-side mesh dispatch: task threads append their slot to a map and bump a `VkDispatchIndirectCommand`, and the mesh pass runs `vkCmdDispatchIndirect` from it. Main still dispatches `groups * threads` from the CPU, which is 0 for `DrawMeshTasksIndirect`, so the floor grass never draws on main. That's the most valuable piece here. - Records taken by an atomic allocator, drawn with `vkCmdDrawIndirect`, and capacity 1024 → 16384. - One shared records/payload pair per recording thread instead of one per program, plus the per-task payload copy (no longer one per slot). - Fallback compute pipelines built on a worker via `compute_state`, plus the precompile pass. Note that #1317 (still open, not merged) solves "a failed pipeline is retried every draw" another way. Your state-2-with-empty-`fallback_compute` covers that case too, so only one of them should land. Worth coordinating with CiLiBi18. - Zeroed vertex-stage outputs and the butterfly shuffle routing are small and fine to keep with this. I read through the new SPIR-V and the dispatch carefully, and the core logic holds up: - The map indices stay below CAPACITY, since each started slot is unique and below it. - The allocator's "fits" check keeps the drawn records a prefix of the buffer, so `vertexCount` only covers written records, and `taken` is a multiple of `corners`. - The header fills don't overlap: draw command in words 0-3, allocator in word 4, dispatch in words 5-7. - The records buffer has `INDIRECT_BUFFER_BIT`. - `full_barrier` is ALL_COMMANDS / MEMORY_READ|WRITE, so it covers the indirect reads. `vkCmdFillBuffer` and `vkCmdDispatchIndirect` both bump `work_epoch`, so the dedupe can't drop the barrier after the fills. - The two extra workgroup barriers in `MeshFallbackEpilogue` sit at top level and correctly broadcast `base` through `mesh_primitive_count`. Things to watch in the rebase are inline: main's `EmitIsbeRead` and runtime_info changes collide with yours. **Unrelated, so please split out (separate PRs, separate testing)** 1. AMD 64-lane graphics warps + the proprietary-driver fragment shuffle workaround (commit 3, plus the `RX_VK_WARP64` / `RX_VK_FRAG_SHUFFLE_FALLBACK` switches). This changes the SPIR-V of every graphics program on every AMD GPU whose graphics stages can't be pinned, including RDNA2/3 and RADV/Steam Deck players who use native mesh shaders. It also changes the cache stamp, so they all re-translate on first start. It may well be right, but it needs its own PR and testing on RDNA2+ (both drivers), not just the 5500M. 2. Texture changes (commits 7 and 14): the BC3 scratch output in system memory, the 30-frame BC3 retry, the "no squeeze while VMA blocks hold 1/8 free" rule, and the evicted-texels log. These are a VRAM-policy change with its own risk (the driver paging instead of us evicting, as you noted). That area also moved on main (`3e66cf9` BC3 scratch kept, the v2.0.11 iGPU caps, #840 batched eviction writebacks merged today), so it will conflict anyway. 3. The Mesh shader fallback menu row (Auto/Off/On) and the launch warning text. They're small and could go with the core PR, but please rebase them over main's current `vk_main.cpp`. 4. `scripts/validate_spirv.py` + its test. It's handy, but it's tooling, and main now has `RX_VK_DUMP_MESH_SPV` for the same job. It can stand on its own. 5. The live debug switch (`hookshot_fallback_debug.txt`, read/info). See the inline comment: I'd gate it behind an env var or leave it out of the core PR. 6. Docs: 400 lines in `docs/mesh_fallback_rdna1.md` is a lot to keep in sync. Most of it is PR history. I'd keep a short "how the fallback works / switches" section and leave the narrative in the PR description. **Suggested path** Start a fresh branch from current main. Cherry-pick only the fallback parts of commits 4e36b5d/5c2004f (the dispatch map, allocator, indirect draw, shared buffers, async build + precompile, zeroed outputs), resolving against #1314's code as noted inline. Drop commits 1, 4, 12, 16, 17 and the parts of commit 6 that #1314 already covers. Then open the warp/shuffle change and the texture changes as two more PRs. That also makes the "needs a mesh-shader GPU to confirm" question easy: the core PR doesn't touch native mesh paths at all once commit 4 is dropped. CI hasn't run here (the branch conflicts), so nothing to report there yet.
gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_mesh.cpp
hckrman101 (Migrated from github.com) commented 2026-11-09 23:29:02 +00:00

Rebase hazard: main (#1314, 14e0865) already rewrites workgroup here as gl_WorkGroupID.x % task_threads. With your map, gl_WorkGroupID.x is the dispatch index, not the slot, so main's version would be wrong on top of your TaskEpilogue. Keep yours (FallbackSlot % threads) and drop main's, or the meshlet index will be garbage again. The same goes for PayloadWord: it has to use the slot from the map (FallbackSlot / threads), not main's WorkgroupId.

Rebase hazard: main (#1314, `14e0865`) already rewrites `workgroup` here as `gl_WorkGroupID.x % task_threads`. With your map, `gl_WorkGroupID.x` is the dispatch index, not the slot, so main's version would be wrong on top of your TaskEpilogue. Keep yours (`FallbackSlot % threads`) and drop main's, or the meshlet index will be garbage again. The same goes for `PayloadWord`: it has to use the slot from the map (`FallbackSlot / threads`), not main's `WorkgroupId`.
gpu/vk/eden/shader_recompiler/runtime_info.h
@ -108,2 +108,4 @@
/// (hookshot) the fallback's record slots: mesh workgroups past this one write nothing
u32 mesh_fallback_capacity{};
/// (hookshot) the task program's threads: a fallback mesh workgroup is the slot task workgroup * threads + its index in the task
u32 mesh_fallback_task_threads{1};
hckrman101 (Migrated from github.com) commented 2026-11-09 23:29:02 +00:00

Main added this field too (#1314), defaulting to {} (0) and set to task_t ? workgroup_size[0] : 0. Main guards its OpUMod with && mesh_fallback_task_threads; yours always emits OpUMod/OpUDiv by Const(threads). When you merge, keep the default of 1 (and the ? : 1 in vk_shader.cpp), or a 0 will turn into a mod/div by zero in the generated SPIR-V.

Main added this field too (#1314), defaulting to `{}` (0) and set to `task_t ? workgroup_size[0] : 0`. Main guards its `OpUMod` with `&& mesh_fallback_task_threads`; yours always emits `OpUMod`/`OpUDiv` by `Const(threads)`. When you merge, keep the default of 1 (and the `? : 1` in `vk_shader.cpp`), or a 0 will turn into a mod/div by zero in the generated SPIR-V.
gpu/vk/vk_draw.cpp
@ -4527,0 +4647,4 @@
/// range and mean, and how many were NaN, infinite, zero, or denormal (an integer seen as a float: word 0, the written flag, is)
void fallback_read_report(const FallbackDebug::Slot& slot, uint32_t i) {
invalidate_mapped(fdbg.mem, kFallbackReadBytes * kFallbackReadSlots);
const uint32_t* w = fdbg.mapped + i * (kFallbackReadBytes / 4);
hckrman101 (Migrated from github.com) commented 2026-11-09 23:29:02 +00:00

Looks right to me: buffers per thread, full_barrier at the top of fallback_record ordering against the previous fallback draw's vertex reads, and pipeline barriers covering earlier command buffers on the same queue. One thing to double-check: how many recording threads you get at most. At ~72 MB each, a machine with many parallel-replay threads can still add up on a 4 GB card. A log line with the thread count next to the existing buffer pairs, N MB line would make that visible.

Looks right to me: buffers per thread, `full_barrier` at the top of `fallback_record` ordering against the previous fallback draw's vertex reads, and pipeline barriers covering earlier command buffers on the same queue. One thing to double-check: how many recording threads you get at most. At ~72 MB each, a machine with many parallel-replay threads can still add up on a 4 GB card. A log line with the thread count next to the existing `buffer pairs, N MB` line would make that visible.
gpu/vk/vk_draw.cpp
@ -4527,0 +4714,4 @@
else if (line.rfind("info ", 0) == 0) fdbg.infos.push_back(line.substr(line.find_first_not_of(" \t", 5)));
else fdbg.prefixes.push_back(line);
}
logf("vk: fallback debug: %s%s, %zu program prefixes left out, %zu read",
hckrman101 (Migrated from github.com) commented 2026-11-09 23:29:02 +00:00

This overlaps #1317 (open), which remembers a failed (key, variant) in a set under the old mutex. Your state 2 with an empty fallback_compute already stops the retry storm, so if this lands, #1317 isn't needed (or the other way round). Also, on rebase, keep main's dw().prog != &p check above the call in fallback_record (#1314, 1b41593).

This overlaps #1317 (open), which remembers a failed (key, variant) in a set under the old mutex. Your state 2 with an empty `fallback_compute` already stops the retry storm, so if this lands, #1317 isn't needed (or the other way round). Also, on rebase, keep main's `dw().prog != &p` check above the call in `fallback_record` (#1314, `1b41593`).
gpu/vk/vk_draw.cpp
hckrman101 (Migrated from github.com) commented 2026-11-09 23:29:02 +00:00

This runs on every fallback draw on every recording thread. It takes a global mutex, and every 500 ms one of those draws fopens a file in the game folder (on Windows with AV, that can be milliseconds on the recording thread). It's a nice tool, but I'd rather it cost nothing by default: gate the whole thing behind e.g. RX_VK_FALLBACK_DEBUG=1, read once into a static, and keep it out of the core fallback PR.

This runs on every fallback draw on every recording thread. It takes a global mutex, and every 500 ms one of those draws `fopen`s a file in the game folder (on Windows with AV, that can be milliseconds on the recording thread). It's a nice tool, but I'd rather it cost nothing by default: gate the whole thing behind e.g. `RX_VK_FALLBACK_DEBUG=1`, read once into a static, and keep it out of the core fallback PR.
gpu/vk/vk_shader.cpp
@ -502,0 +501,4 @@
// (vk_main.cpp). The translator then keeps each 32-lane guest warp to its half of the host's, as Eden does on AMD; the
// long path is right at 32 lanes too. The ballots and shuffles of the game's fragment shaders ran across both halves
p.warp_size_potentially_larger_than_guest =
dev.subgroup_size > 32 && (dev.sg32_stages & VK_SHADER_STAGE_ALL_GRAPHICS) != VK_SHADER_STAGE_ALL_GRAPHICS;
hckrman101 (Migrated from github.com) commented 2026-11-09 23:29:03 +00:00

This flips the warp path for every graphics program on any AMD GPU that can't pin graphics stages to 32 lanes, including RDNA2/3 on both AMD's driver and RADV, where native mesh shaders are used and this PR's fallback isn't involved. With the stamp change, those players all re-translate on their first start. Please move it (and the fragment-shuffle workaround below) into its own PR with testing on an RDNA2+ card, so a regression there doesn't hold up the fallback work.

This flips the warp path for every graphics program on any AMD GPU that can't pin graphics stages to 32 lanes, including RDNA2/3 on both AMD's driver and RADV, where native mesh shaders are used and this PR's fallback isn't involved. With the stamp change, those players all re-translate on their first start. Please move it (and the fragment-shuffle workaround below) into its own PR with testing on an RDNA2+ card, so a regression there doesn't hold up the fallback work.
gpu/vk/vk_texture.cpp
hckrman101 (Migrated from github.com) commented 2026-11-09 23:29:03 +00:00

This is a VRAM policy change for every GPU, not just the fallback: at the driver's limit with 1/8 free in the blocks, we now stop evicting and rely on the driver not paging. It may be the right trade-off, but it belongs in a separate texture PR with testing on a couple of low-VRAM cards. This area also changed on main today (#840 batched eviction writebacks), so it will conflict.

This is a VRAM policy change for every GPU, not just the fallback: at the driver's limit with 1/8 free in the blocks, we now stop evicting and rely on the driver not paging. It may be the right trade-off, but it belongs in a separate texture PR with testing on a couple of low-VRAM cards. This area also changed on main today (#840 batched eviction writebacks), so it will conflict.
NotRexed commented 2026-11-10 02:08:13 +00:00 (Migrated from github.com)

Thanks for all the work on the RDNA1 fallback. The fallback itself is on main, the v2.0.12 branch carries its crash fixes (#1314, #1317), and #1383 is next. This branch conflicts with main in 10 files and mixes many changes. Closing; any fix not covered by those would be very welcome as a small PR of its own.

Thanks for all the work on the RDNA1 fallback. The fallback itself is on main, the v2.0.12 branch carries its crash fixes (#1314, #1317), and #1383 is next. This branch conflicts with main in 10 files and mixes many changes. Closing; any fix not covered by those would be very welcome as a small PR of its own.
rakefire-impala commented 2026-11-10 05:44:35 +00:00 (Migrated from github.com)

Thanks, that makes sense. I've split what main doesn't already cover into small PRs, each fixing one thing, all on current main:

  • #1505 rainbow block-noise textures on low-VRAM cards: a failed BC3 upload is tried again, not drawn as noise (#993)
  • #1506 slow resume and black HUD hearts on 4 GB cards: the texture squeeze waits while VMA's blocks have room
  • #1507 dark lighting and glowing bloom on AMD: 64-lane graphics warps and Lichii #273's fragment shuffle fallback
  • #1508 scripts/validate_spirv.py: an offline spirv-val pass over the program cache
  • #1509 mesh fallback: wait for the last draw before clearing the shared buffers (the hazard raised in #1383's review)

They merge cleanly with main and with each other. #1505–#1507 were tested on an RX 5500M (RDNA1, 4 GB, no mesh shaders). #1507 changes the graphics translation on AMD cards that can't pin graphics subgroups, so an RDNA2+ tester would help before merging it.

Thanks, that makes sense. I've split what main doesn't already cover into small PRs, each fixing one thing, all on current main: - #1505 rainbow block-noise textures on low-VRAM cards: a failed BC3 upload is tried again, not drawn as noise (#993) - #1506 slow resume and black HUD hearts on 4 GB cards: the texture squeeze waits while VMA's blocks have room - #1507 dark lighting and glowing bloom on AMD: 64-lane graphics warps and Lichii #273's fragment shuffle fallback - #1508 `scripts/validate_spirv.py`: an offline `spirv-val` pass over the program cache - #1509 mesh fallback: wait for the last draw before clearing the shared buffers (the hazard raised in #1383's review) They merge cleanly with main and with each other. #1505–#1507 were tested on an RX 5500M (RDNA1, 4 GB, no mesh shaders). #1507 changes the graphics translation on AMD cards that can't pin graphics subgroups, so an RDNA2+ tester would help before merging it.
CiLiBi18 commented 2026-11-10 06:42:43 +00:00 (Migrated from github.com)

Tested on a GTX 1070 (Pascal, no mesh shaders) and it brings the missing geometry back.

Setup: this PR's head (16229ef) built as-is on Linux, GTX 1070 8 GB / NVIDIA 580.178.04, i7-7700HQ, output 1920x1080 at 75%. The comparison is v2.0.11 + #1314 + #1317, played through the same part of the game the night before.

  • Geometry: on v2.0.11 + #1314/#1317, some mesh-shader geometry ran through the compute fallback but wasn't drawn, while its collision stayed. Inside the Deku Tree, the walls and the block you push to climb up and go on were invisible, which made the dungeon very hard to follow. With this PR they are all drawn. I haven't checked Link's house door yet.
  • Stability: 0 driver faulted, 0 failed fallback pipelines. The only failed programs are the 10 movie-upscaler networks, which are left out on purpose on this card.
  • Frame rate, gameplay only: mostly 30–40 fps, with GPU time 15–18 ms and submit 12–20 ms a frame, at 500–700 draws and ~90 fallback dispatches. It isn't a like-for-like comparison with my v2.0.11 runs (median 25 fps), because the scenes weren't the same.
  • Stalls: the first two minutes in a new area had seven submits of 1.7–9.8 s, almost all of it in draw. One fallback compute pipeline took 101 s to compile on a worker: fallback: the compute pipeline of 55d1da3fea1603d4 stage 1: SLOW compile, 101648 ms (on a worker). The 9.8 s stall came right after it. This was a fresh game folder with an empty pipeline cache (precompile: ... 0 graphics), so part of this is first-time compiling. After those two minutes, I saw no more stalls.

One more data point: the storage buffer NOT bound (reads zeros) lines for the far-foliage programs (68642bb810ee3e7c, 51183fe750d5d0c8, d905f52b7517c681…, at c5+0x60 / c5+0x1d0) still show up here and reach the log's cap of 40. So they weren't what hid the geometry on this card.

Happy to run anything else on Pascal if it helps get this merged.

**Tested on a GTX 1070 (Pascal, no mesh shaders) and it brings the missing geometry back.** Setup: this PR's head (`16229ef`) built as-is on Linux, GTX 1070 8 GB / NVIDIA 580.178.04, i7-7700HQ, output 1920x1080 at 75%. The comparison is v2.0.11 + #1314 + #1317, played through the same part of the game the night before. - **Geometry:** on v2.0.11 + #1314/#1317, some mesh-shader geometry ran through the compute fallback but wasn't drawn, while its collision stayed. Inside the Deku Tree, the walls and the block you push to climb up and go on were invisible, which made the dungeon very hard to follow. With this PR they are all drawn. I haven't checked Link's house door yet. - **Stability:** 0 `driver faulted`, 0 failed fallback pipelines. The only failed programs are the 10 movie-upscaler networks, which are left out on purpose on this card. - **Frame rate, gameplay only:** mostly 30–40 fps, with GPU time 15–18 ms and submit 12–20 ms a frame, at 500–700 draws and ~90 fallback dispatches. It isn't a like-for-like comparison with my v2.0.11 runs (median 25 fps), because the scenes weren't the same. - **Stalls:** the first two minutes in a new area had seven submits of 1.7–9.8 s, almost all of it in `draw`. One fallback compute pipeline took 101 s to compile on a worker: `fallback: the compute pipeline of 55d1da3fea1603d4 stage 1: SLOW compile, 101648 ms (on a worker)`. The 9.8 s stall came right after it. This was a fresh game folder with an empty pipeline cache (`precompile: ... 0 graphics`), so part of this is first-time compiling. After those two minutes, I saw no more stalls. One more data point: the `storage buffer NOT bound (reads zeros)` lines for the far-foliage programs (`68642bb810ee3e7c`, `51183fe750d5d0c8`, `d905f52b7517c681`…, at `c5+0x60` / `c5+0x1d0`) still show up here and reach the log's cap of 40. So they weren't what hid the geometry on this card. Happy to run anything else on Pascal if it helps get this merged.
rakefire-impala commented 2026-11-10 11:09:40 +00:00 (Migrated from github.com)

Thanks for testing this on Pascal. It's the first report from a second GPU vendor, and the Deku Tree result is exactly the kind of thing this needed.

#1023 was closed so its fixes could go in as small PRs. The parts you tested are now:

  • #1510 and #1511: the missing geometry. Indirect task draws ran no mesh workgroups, and a draw had only 1024 slots (now 65535). These are most likely what brought your walls and the push block back. #1511 is stacked on #1510, so #1510 goes first.
  • #1512: the fallback's compute pipelines are built on a worker, and in the precompile at start for the programs the cache already has, instead of inside a draw. Whether that covers your 101 s compile of 55d1da3fea1603d4 and the 1.7-9.8 s stalls in a new area's first minutes is worth checking on a second start.
  • #1509 and #1513: fallback flicker and coloured specks. #1581: shadow flicker on cards where the TAAU network is left out.

All of them merge cleanly with main and with each other. If you could run main + #1509-#1513 (and #1581 if you like) on the GTX 1070 and post what you see on #1510 or #1511, that would be the strongest case for merging them.

Thanks for testing this on Pascal. It's the first report from a second GPU vendor, and the Deku Tree result is exactly the kind of thing this needed. #1023 was closed so its fixes could go in as small PRs. The parts you tested are now: - **#1510 and #1511: the missing geometry.** Indirect task draws ran no mesh workgroups, and a draw had only 1024 slots (now 65535). These are most likely what brought your walls and the push block back. #1511 is stacked on #1510, so #1510 goes first. - **#1512: the fallback's compute pipelines** are built on a worker, and in the precompile at start for the programs the cache already has, instead of inside a draw. Whether that covers your 101 s compile of `55d1da3fea1603d4` and the 1.7-9.8 s stalls in a new area's first minutes is worth checking on a second start. - **#1509 and #1513:** fallback flicker and coloured specks. **#1581:** shadow flicker on cards where the TAAU network is left out. All of them merge cleanly with main and with each other. If you could run main + #1509-#1513 (and #1581 if you like) on the GTX 1070 and post what you see on #1510 or #1511, that would be the strongest case for merging them.
CiLiBi18 commented 2026-11-10 14:32:05 +00:00 (Migrated from github.com)

Thanks for splitting it up. I'll build main (88b554e) + #1509–#1513 + #1581 as they are on the GTX 1070 (they merge cleanly here too) and report on #1510. I'll check:

  • Geometry: the Deku Tree walls and the push block, and Link's house (the missing door and the black faceted shape), which I hadn't checked on #1023.
  • Stalls, for #1512: the first start with an empty pipeline cache, then a second start in the same folder over the same route. I'll look for 55d1da3fea1603d4 and the submit times in the first minutes of a new area.
  • Flicker and specks (#1509, #1513), and tree shadows (#1581), on NVIDIA.
Thanks for splitting it up. I'll build main (`88b554e`) + #1509–#1513 + #1581 as they are on the GTX 1070 (they merge cleanly here too) and report on #1510. I'll check: - **Geometry:** the Deku Tree walls and the push block, and Link's house (the missing door and the black faceted shape), which I hadn't checked on #1023. - **Stalls, for #1512:** the first start with an empty pipeline cache, then a second start in the same folder over the same route. I'll look for `55d1da3fea1603d4` and the submit times in the first minutes of a new area. - **Flicker and specks (#1509, #1513), and tree shadows (#1581),** on NVIDIA.
Commenting is not possible because the repository is archived.
Reviewers
No reviewers
Labels
Clear labels   
accessibility
Barrier affecting people with disabilities

  
autorelease: pending

   
autorelease: tagged

   
bug
Something isn't working

  
dependencies
Pull requests that update a dependency file

  
documentation
Improvements or additions to documentation

  
duplicate
This issue or pull request already exists

  
enhancement
New feature or request

  
github_actions
Pull requests that update GitHub Actions code

  
good first issue
Good for newcomers

  
help wanted
Extra attention is needed

  
invalid
This doesn't seem right

  
question
Further information is requested

  
testing wanted

   
wontfix
This will not be worked on

No labels
accessibility
autorelease: pending
autorelease: tagged
bug
dependencies
documentation
duplicate
enhancement
github_actions
good first issue
help wanted
invalid
question
testing wanted
wontfix
Milestone
Clear milestone
No items
No milestone
Assignees
Clear assignees
No assignees
1 participant
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set

Reference
MagicBeanDev/Hookshot!1023
No description provided.