Compare commits

...
Sign in to create a new pull request.
Author SHA1 Message Date
Takoroka
50fa86bbbc nx2: fix the Linux (GCC) build of the vi2 present stats
std::chrono::microseconds::rep is long on Linux and long long on Windows,
so std::max over the long long stats and a converted duration did not
compile with GCC. Return long long from the conversion.
2026-10-29 19:23:21 -04:00
Takoroka
c7fb540505 Update README.md 2026-10-29 23:37:31 +01:00
Takoroka
ab12442b56 Update NX2_NOTES.md
gave context to tested games, removed the third game as it was only tested once
2026-10-29 23:19:06 +01:00
Takoroka
b7ff03036c nx2: support mapping modify so NVN2 places compressible textures
NVN2 gives a texture in a compressible memory pool its GPU address only
after NvRmGpuMappingModify switched the pool's mapping to the texture's
kind. NvRm fails that call unless the GPU characteristics report
NVGPU_GPU_FLAGS_SUPPORT_MAPPING_MODIFY, so such textures kept address 0
and draws into them went nowhere: title 1's cup-opening movie stayed on
its first frame (also its 640x360 and 512x512 passes).

Report the flag and answer the request nn::gpu::sdk then sends, gpu
sub-session 212 (NormalModifyMemoryKind4K), with the wanted kind and a
signalled job slot. The movie now plays on screen and the race follows.
2026-10-29 18:01:44 -04:00
Takoroka
278a2a0e05 core: use the 12 GiB layout and auto-stub IPC for Switch 2 titles
Switch 2 titles size their heap from the total memory and stop at the
first unimplemented command, so they only booted with the 12 GiB layout
and EDEN_NX2_AUTO_STUB set by hand. Detect them from main.npdm before the
kernel starts and apply the layout as a per-game value (the configured
layout comes back once the game is closed), and auto-stub unknown
commands from Switch 2 processes.
2026-10-29 15:56:56 -04:00
Takoroka
a482274421 build: fixes for clang-cl
- (std::min)/(std::max) where windows.h's macros are visible
- LAMBDA_FORCEINLINE: always_inline under clang (clang-cl defines
  _MSC_VER but rejects [[msvc::forceinline]] on lambdas)
- ImageBase::FindView: find_if, as MSVC's STL takes its vectorized find
  for ImageViewInfo under clang-cl and fails a size assertion
- parentheses for -Wlogical-op-parentheses in the Ampere lowering
2026-10-29 15:56:45 -04:00
Takoroka
df081e2f9d video_core: stop resuming pending channels once the system is powered off
At shutdown is_powered_on is cleared before the GPU thread is stopped.
A channel suspended on a semaphore acquire that can resume then went
through DispatchCalls, which returns at once when powered off, so the
acquire stayed pending and ResumePending counted it as progress forever:
the GPU thread never got back to its stop request and closing the game
hung in ThreadManager::NotifyShutdown.
2026-10-29 15:56:45 -04:00
Takoroka
9fa008fda3 nx2: decode Switch 2 video (nn::codec::video) with FFmpeg
The decoder service used to complete every submitted frame at once
without decoding it. Each decoder service (gsv 902/906/1002 -> TIPC 1)
now decodes on a thread of its own with FFmpeg (VP9, VP8, H.264;
video_core/host1x/nx2_video_decoder), in software or on the host GPU per
the NVDEC emulation setting.

- 101 takes the codec name ("video/vp9") from its A buffer.
- 109 queues the A buffer's frame; the ready event fires while the queue
  has room.
- 311/314 reply a NativeSyncHandle {~syncpoint, 0, value}: the service
  allocates a syncpoint and increments it per finished frame, frames
  filling output images in register/release order like the client's
  ring. 314 waits for the client's release fence before reuse.
- 213 returns the frame's size, pts (TimeSpan, ns) and end of stream.
- gsv 911 creates a decoder image pool (count images in one memory) and
  gsv 961 returns the 128-byte CpuUserImageDescriptor that
  nn::codec::ReadDecoderOutputImage* copy frames out with; images are
  pitch-linear NV12, RGBA or BGRA.
- Each new video command is logged once in full; EDEN_NX2_VDEC_DUMP=<dir>
  saves every 120th decoded frame's luma.

Title 1's Mushroom Cup opening (2560x1440 VP9, 1423 frames) now decodes
in real time and the game goes on to the race. The picture on screen
still stays on the first frame: the game memcpys each frame into GPU
memory and the texture cache keeps its first upload (notes 148).
2026-10-29 15:55:22 -04:00
Takoroka
06087bfc57 nx2: syncpt cancel always signals; a new slot request starts unsignalled
nn::syncpt::sdk's CancelSyncPointEvent sends 202 and then always waits on
the slot's event, so 202 now signals it also when no request was pending.
Title 1's movie player cancels each output-buffer slot before its first
request and hung there forever (the cup-opening screen stayed black and
the decoder was never started).

Slots are reused by different client objects: the movie player got slot
0 after NVN's fence waits had left its event signalled, read that as a
decoded frame, found none and aborted with 2111-3604. 201 now clears the
slot's event before registering the new request.
2026-10-29 15:55:00 -04:00
Takoroka
92f0d80513 nx2: gsv 906 opens a video decoder session
Newer SDKs (nn::codec::video::detail::InitializeServerContext, title 4)
send 906(u8, u64 decoder library version) where older ones send 902.
Auto-stubbed, it returned no handle and title 4 aborted with
InvalidHandle (0xE401) when its first movie started, about 30 minutes
into a new game. 906 now hands out the same session as 902.
2026-10-29 14:59:22 -04:00
Takoroka
de004b6e9f vulkan: leave the alpha test reference out of pipelines that don't test alpha
FixedPipelineState stored alpha_test_ref whatever the alpha test
function. Title 1 animates the register with the test off (text fading
in and out), so every value made a new pipeline whose shaders were
lowered and translated on the GPU thread: bursts of 600-1,900 pipelines
in 10 s (up to 5 s of GPU thread) whenever text was on screen. The
reference is now kept only while the test compares with it (not Always
or Never). Driving title 1's menus: 400+ pipelines for already-known
shaders before, 0 after; one shader lowered 79 times before, 2 after.

EDEN_NX2_PIPELINE_DIFF=1 logs which fixed-state fields differ when a new
pipeline reuses known shaders.
2026-10-29 14:28:34 -04:00
Takoroka
ac7ccd35f9 video_core: GPU thread split on by default for Switch 2 titles; cheaper front paths
Clean A/B in title 1 (3 rounds each): 1.46x the draws a second of one
thread, frames 17% shorter, no crashes, front register copies equal to
the back's in 57,089 checks. EDEN_NX2_SPLIT=0 runs everything on one
thread, =1 keeps macros on the back.

The front's per-write paths no longer check guarded statics
(GpuSplit::Get, Nx2Profile::Enabled, the skip-same switches) and the
ring's method writes are inline. Notes row 145.
2026-10-29 11:55:07 -04:00
Takoroka
b884c25d6e nx2: notes 142-144 (GPU thread split, scheduler race, front register copies) 2026-10-29 11:18:19 -04:00
Takoroka
7ef288ce7e video_core: split the GPU thread in two (EDEN_NX2_SPLIT, off by default)
The GPU thread (front) takes the CPU's commands, parses command lists
and runs the puller; engine method calls and everything behind them
(registers, draws, dispatches, the rasterizer and its caches) run in
order on a second thread (back), fed through a single-producer
single-consumer ring (gpu_split.h). Work the front did on back-owned
state (end of command list flushes, channel binds, semaphore releases,
syncpoint increments, flush/invalidate requests, presentation) is
deferred into the same ring; the front waits for the back only where it
needs a result from that state: semaphore acquires and MME DMA reads.

EDEN_NX2_SPLIT=1 runs macros on the back. EDEN_NX2_SPLIT=2 runs the 3D
and compute macros on the front against copies of the registers they
read (Maxwell3D/KeplerCompute Front*), drops state writes that keep a
register's value there, and sends the rest. EDEN_NX2_SPLIT_VERIFY=1
compares the copies with the back's registers every 4096 macro calls
(0 differences in 40,193 checks in title 1); EDEN_NX2_SPLIT_CHECK=1
logs the stack of any other thread using the Vulkan scheduler while the
back runs. Switch 2 titles, Vulkan, non-accurate GPU only.

Title 1, 3D attract loop: ~44,000 draws a second on one thread, ~47,400
with EDEN_NX2_SPLIT=1, 50,000-66,500 with EDEN_NX2_SPLIT=2.
2026-10-29 11:16:43 -04:00
Takoroka
60308ba40c vulkan: never flush the scheduler from the async present thread
Swapchain::AcquireNextImage waits for the main scheduler's tick of the
image's last use; when that tick was not submitted yet, Scheduler::Wait
flushed, recording into the current chunk on the present thread while
the GPU thread recorded into it too (corrupt command buffers, crashes
in vkCmdEndRenderPass and in CommandChunk::Record). With async
presentation the present thread now waits up to 20 ms for the
recording thread to submit the tick, and goes on without the wait
otherwise (the copy to the swapchain has its own fences).
2026-10-29 11:15:48 -04:00
Takoroka
9bc673df95 nx2: profile engine method calls and semaphore acquires; notes 140-141
EDEN_NX2_GPU_PROFILE now times engine method calls (the work a split GPU thread's back half
would run) and counts semaphore acquires and those that wait. Title 1: method calls ~343 ms
a second (6.9 million), front-half work (command parsing, macros) ~333 ms, acquires ~2,700 a
second of which ~220 wait. Notes 140 (macro compiler) and 141 (draw setup).
2026-10-29 09:45:30 -04:00
Takoroka
3755efe1c9 video_core: compile Turing MME macro code to x86-64
Title 1 calls macros ~1.3 million times a second; most were replayed from the macro memo,
whose lookups missed the CPU caches inside the GPU thread (~8%), and the rest interpreted.
macro_turing_jit.h compiles all of the uploaded macro memory with Xbyak: each instruction a
block falling through to the next, branches and loops through a jump table, with the
interpreter's semantics (lanes reading operands before either writes, predication as two
variants, the END delay slot, loops, carry). Loop checks only where a loop can end, END
checks only after an END, the instruction limit counted on indirect jumps. Emits and register
reads call back into C++. It is recompiled when the macro code changes and used for 3D and
compute macros on x86-64; EDEN_NX2_MME_JIT=0 goes back to the memo and the interpreter.

Checked against the interpreter on 200,000 recorded calls (registers, data RAM, every
method write, FIFO use): identical. In the game: macro calls 148 ns instead of 174 (two
profiled runs each), no failures, 0 crashes in 3 runs. NX2_GPU_THREAD_SPLIT.md describes
the next step, splitting the GPU thread.
2026-10-29 09:39:25 -04:00
Takoroka
a8ee576823 nx2: notes 138-139 (async presentation retested, buffer collector on its own size) 2026-10-29 08:51:18 -04:00
Takoroka
3434220055 buffer_cache: collect on the buffers' own size for Switch 2 titles
The buffer collector compared the whole device's memory use with thresholds from the
device size, so with title 1's textures filling an 8 GB GPU (~6.8 GB used, buffers ~200
MiB of it) it ran aggressively every frame: up to 64 buffers evicted and recreated later,
the GPU-written ones downloaded first with a wait for the GPU (~5% of the GPU thread).
For Switch 2 titles it now runs when the buffers themselves pass 1 GiB (aggressive past
2 GiB); below the aggressive mode it also keeps buffers holding GPU-written data, which
it would have to download. Other titles keep the device-wide thresholds. Buffer
downloads by the collector: 5.4% -> 0%; the cache settles around 540 MiB.
2026-10-29 08:38:09 -04:00
Takoroka
24d0773c53 maxwell_3d: no method dispatch for non-executable register writes; notes 133-137
ProcessMethodCall and the draw manager behind it act only on executable methods, so
CallMethod stops at the dirty flags for the rest (macro replays write ~2.4 million of
those a second). No change beyond run-to-run noise in two profiled runs; it only removes
work. Notes: row 133 corrected (those courses are bound by the GPU thread, not the host
GPU: 20-40% busy), rows 135-137 (same-value writes, native compute registers, 2D blits).
2026-10-29 08:12:52 -04:00
Takoroka
4d8b512cff kepler_compute: store plain native register writes directly
Most native compute register writes (QMD fields its macros replay for thousands of
dispatches a second) only store the value; they went through ProcessNativeMethod's
macro, QMD, upload and semaphore checks. CallMethod now stores them through a table of
plain methods. ProcessNativeMethod's own time 2.2% -> 1.6% of the GPU thread.
2026-10-29 08:04:42 -04:00
Takoroka
72d9d1b64d maxwell_3d: drop state register writes that keep the register's value
Title 1 makes ~6 million 3D register writes a second and ~60% of them write the value
the register already holds (counted with EDEN_NX2_GPU_PROFILE), most from macro memo
replays. For registers without side effects (not in execution_mask: draws, clears,
constant buffer data and binds, semaphores, ...) such a write changes nothing, but it
set dirty flags that made the next draw re-check that state. CallMethod and the sunk
writes now skip them after shadow RAM processing. Two profiled runs each: macro calls
169 ns instead of 185, draw setup 5.80 us instead of 5.95. EDEN_NX2_SKIP_SAME=0 turns it
off.
2026-10-29 07:49:09 -04:00
Takoroka
719cd0b315 buffer_cache: one GPU wait for all buffers the collector evicts
The buffer collector downloaded each evicted buffer's GPU-written data with its own
wait for the GPU. Without per-fence downloads (d530bc4) most of title 1's evicted
buffers are GPU-written, and on an 8 GB GPU near its budget it evicts some every frame.
Their ranges now go to one staging buffer, copied with one wait.
2026-10-29 07:26:48 -04:00
Takoroka
5eb7c1944b nx2: notes 131-134 (buffer downloads, inline upload lookups, VRAM-bound courses, deferred inline invalidation) 2026-10-29 07:09:52 -04:00
Takoroka
3e0b6e8728 vulkan: defer the caches' invalidation for inline writes outside GPU-written data
Title 1 makes thousands of small inline-to-memory writes between draws, and each one
invalidated the buffer, texture, shader and query caches on its own (~7.5% of the GPU
thread after 0fe581f). Guest memory is still written at once, query results pending
there are still dropped at once, and writes into GPU-written buffer data are still
uploaded in place; the buffer, texture and shader caches now get the merged ranges at
the next FlushCaching, which runs before every draw, dispatch, clear, copy and 2D blit.
Inline writes: 4.4% of the GPU thread (deferred invalidation included in FlushCaching,
3.9% -> 3.1%). EDEN_NX2_INLINE_DEFER=0 invalidates at once.
2026-10-29 07:09:36 -04:00
Takoroka
39dd849e24 texture_cache: count evictions that download; EDEN_NX2_TC_MIN_AGE knob
The collector's periodic log now also counts evictions that downloaded the image first
(each a full GPU wait). EDEN_NX2_TC_MIN_AGE=<frames> keeps images used more recently
than that (a VRAM thrash test on 8 GB: 120 frames made no difference over three runs).
2026-10-29 07:05:03 -04:00
Takoroka
ead04e6aec common: pass-through hash for integer and pointer keys; RangeSet::Intersects
Common::unordered_map/set defaulted to std::hash, a byte-wise FNV-1a on MSVC for
integers and pointers, which boost's flat containers then mix again: the GPU thread's
page maps (shader and query cache invalidation, per inline upload) spent ~1% in it.
Integer and pointer keys now go through as they are and boost mixes them.
BufferCache::IsRegionGpuModified visited every GPU-written range inside the query just
to answer yes or no; without per-fence downloads those ranges stay fragmented and it was
~4% of title 1's GPU thread (inline uploads). RangeSet::Intersects answers with one
lookup. Inline uploads: 10.5% -> 7.5% of the GPU thread.
2026-10-29 07:05:03 -04:00
Takoroka
148d0522f3 buffer_cache: no buffer downloads at every fence for Switch 2 titles
Eden copied every GPU-written buffer range back to guest memory at the next fence.
Title 1 signals a fence per compute dispatch, so that was 250-330 MB/s in 13,000-20,000
ranges a second, with the fencing thread holding the buffer cache while it wrote them.
Measured: the native MME's DMA reads never touched those pages (they read query words
written on the CPU side), and with reactive flushing the guest CPU read GPU-written
memory at most 64 times in 200 s. Without them, for Switch 2 titles (native tiling):
fences ~1.1 us instead of ~5, draw setup 6.1 us instead of 6.4-6.7, macro calls 160 ns
instead of 167-170, GPU-thread lock spinning 1.8% (was 5.5%). Regions stay marked GPU
written and are still flushed on demand. EDEN_NX2_BUFFER_DOWNLOADS=0/1 forces either
way. Reactive flushing reads of GPU-written memory are now counted in the log.
2026-10-29 06:22:54 -04:00
Takoroka
a6b3589eec nx2: notes 123-130 (fences, lock handoffs, translation cache crash, macro parameters, counter polling, worker wake-ups, memo compaction, measurement pitfalls) 2026-10-29 05:48:05 -04:00
Takoroka
963e843721 video_core: compact the macro memo arena; download profiling counters
Calls stored again with another shape leave their old record in the memo's arena;
title 1 filled its 16 MiB every ~2 s and clearing it made the next ~20,000 calls
interpret again. The live records are now copied to a new arena (a clear only when
they alone fill three quarters of it). The memo statistics count hits on the last
call's slot and compactions. With EDEN_NX2_GPU_PROFILE the GPU thread also reports
how much downloaded buffer data lands on pages the native MME has DMA-read.
2026-10-29 05:46:59 -04:00
Takoroka
556144e266 vulkan: dispatch to the worker every 32 draws, wake it only when it waits
Every 8 draws the GPU thread handed a command chunk to the Vulkan worker and woke it
with a system call (~2% of the GPU thread at ~55,000 draws a second). Chunks now go
every 32 draws (a full chunk still goes on its own), and the worker is woken only
when it is waiting; a busy worker pops the queue before waiting again.
2026-10-29 05:46:59 -04:00
Takoroka
e2054a2d96 core: pause in guest loops that poll the system counter
Title 1's main thread reads CNTPCT ~17 million times a second while it waits for the
GPU (~23% of its samples are the counter read). On an SMT core that spin takes
resources from the sibling thread, often the GPU thread. Two reads within ~1.7 us of
each other now execute two PAUSE instructions first; the guest sees the same values.
EDEN_NX2_CNTPCT_PAUSE=0 turns it off.
2026-10-29 05:46:59 -04:00
Takoroka
b0f8b14cce video_core: copy whole macro calls to the stack
Nearly every macro call arrives whole in one method run (~1,000,000 a second in title
1); its parameters were appended to a std::vector first (~2% of the GPU thread). They
are now copied to a stack array. Still a copy: the command list is often guest memory
read in place, which the guest may reuse while the macro runs, and the macro memo
stores the parameters after the call. Accurate GPU mode keeps the vector (it rewrites
parameters in place). EDEN_NX2_MACRO_COPY=1 always uses the vector.
2026-10-29 05:46:58 -04:00
Takoroka
af1b2fd5cd video_core: release already-signaled fences together
The fencing thread now takes every fence the host GPU has already passed (up to 64)
and pops their texture and buffer downloads under one lock acquisition per cache,
instead of one per fence (title 1 signals 5,000-7,000 a second; the GPU thread spun,
slept and woke the fencing thread for each). Queries and operations keep their
per-fence order. In the sampler, the GPU thread's spinning on the cache locks went
from 7.0% to 5.5%; frame times are within the attract loop's run-to-run noise.
EDEN_NX2_FENCE_BATCH=0 releases one fence at a time.
2026-10-29 05:46:58 -04:00
Takoroka
d7d289636d core: per-thread device memory translation cache
DeviceMemoryManager kept one four-entry page translation cache for every thread. Map
(on a guest service thread, e.g. nx2:gpu 211 job mappings) cleared it while the GPU
thread was between its null check of an entry and the copy from it: the GPU thread
then read from address 0 plus the page offset. Title 1 crashed this way in about half
of its 150 s test runs (BindHostGraphicsUniformBuffer -> ReadBlockUnsafe, fault at
0x600). Each thread now keeps its own cache, which Map and Unmap invalidate through a
generation counter. 0 crashes in 6 runs after.
2026-10-29 05:46:38 -04:00
Takoroka
6cfd9bf041 video_core: skip empty fences; one lock per cache per fence; single-pass overlap subtract
Switch 2 titles signal a fence for each compute dispatch's release (5,000-7,000 a
second in title 1). With fast GPU fences the callback runs at once, so a fence with
no downloads, queries or queued operations does nothing: those (~20%) are no longer
made (one in 64 still is, for the operations fences carry). Each cache's lock is now
taken once per fence (the flush check took both together first). The overlap range
set's Subtract restarted its search after every erase (quadratic); the fencing thread
runs it for every downloaded range while holding the buffer cache. The GPU profiler
counts fences, skipped fences and buffer download bytes.
2026-10-29 04:40:26 -04:00
Takoroka
e090f4e946 nx2: notes 120-122 (cache lock contention, shader translation, final measurements) 2026-10-29 02:25:45 -04:00
Takoroka
9472293e0e nx2: macro memo tries each macro's last call first; FIFO filled only when interpreting
46% of title 1's macro calls repeat the previous call of the same
macro exactly: a per-start slot holding that record's offset skips the
hash and the map probe for them (the record is still in the CPU cache).
The interpreter's parameter FIFO was filled before the memo lookup; it
is now filled only for calls that are interpreted.

Checked against the previous interpreter on 200,000 recorded 3D-scene
calls: identical writes, registers and data RAM.
2026-10-29 02:15:53 -04:00
Takoroka
6967b0f199 shader_recompiler: linear-time breadth-first search; cache storage buffer tracking
Translating title 1's 28,102-instruction compute shader took 3.4 s on
the GPU thread (a freeze the first time it runs), three quarters of it
in GlobalMemoryToStorageBufferPass:
- BreadthFirstSearch checked "visited" with a linear scan per argument,
  quadratic once a search walks thousands of instructions; past 64 it
  now uses a hash set (same visit order, same result).
- The pass searched again for every global memory access; accesses
  through the same address now share the result (the IR does not
  change while storage buffers are collected).
That shader now takes 0.85 s; other new pipelines also cost less.
2026-10-29 02:06:56 -04:00
Takoroka
db61b2bf32 vulkan: log runtime pipeline builds on the GPU thread over 2 ms and their total every 10 s
Tells how much of the stutter on new content is shader translation
(the driver compile already runs on workers).
2026-10-29 01:51:42 -04:00
Takoroka
fa6b3e1a31 video_core: less waiting on the GPU caches' locks
- The texture and buffer caches' mutexes spin briefly before sleeping
  (Common::SpinRecursiveMutex): the GPU thread, the fencing thread and
  CPU write traps take them tens of thousands of times a second for
  microseconds, and a sleeping waiter's wake-up took longer than that.
  Title 1's GPU thread slept 7.7% of its time there, now 3.7%.
- Fences commit and pop texture and buffer downloads separately, and
  only when there are any (one fence per compute dispatch in Switch 2
  titles): the fencing thread no longer holds both caches while it
  copies buffer downloads.
- Inline uploads and pre-draw invalidation take the caches one at a
  time again: holding one while std::lock waited for the other waited
  longer.
2026-10-29 01:51:42 -04:00
Takoroka
ee69e5684c nx2: notes 117-119 (macro memo, async presentation crashes, tester log findings) 2026-10-29 01:26:52 -04:00
Takoroka
daee990471 video_core: cheaper small GPU reads, buffer usage tracking and compute uniforms
- MemoryManager::ReadBlockUnsafe reads within one 4 KiB page through a
  single translation (as Read<T> does) instead of the page walk:
  texture and sampler descriptors, several per draw (2.6% -> 0.1% of
  title 1's GPU thread).
- UsageTracker skips ranges it tracked just before (until Reset): draws
  bind the same multi-megabyte storage buffers again and again, and
  each Track wrote a word per 4 KiB (2.0% -> 0.7%).
- Small compute uniform buffers the GPU has not written are streamed,
  as graphics ones are; synchronizing them re-protected pages an inline
  upload had just unprotected, two VirtualProtect calls per dispatch
  (1.9% -> 0.7%, page protection 3.2% -> 2.3%).
2026-10-29 01:26:52 -04:00
Takoroka
1b852c1e3e video_core: guard the fence manager's uncommitted operations
SyncOperation appended to uncommitted_operations without a lock while
SignalFence moved that list into pending_operations under `guard` on
another thread; a list corrupted this way crashed the fencing thread.
They now share a lock of their own, held for no callbacks (SignalFence
runs callbacks under `guard`, so taking `guard` here could deadlock).
2026-10-29 01:26:52 -04:00
Takoroka
30b75cd85b texture_cache: Switch 2 collector thresholds follow the device memory budget
Native titles' thresholds came from Eden's device memory figure, which
is capped near 7 GB: on a tester's RTX 4070 SUPER (11.2 GB budget) title
1 kept evicting 800-1,600 textures per 600 frames with 3 GB unused. They
now follow the budget (critical at budget - 768 MiB, expected at budget
- 1536 MiB, re-read every 60 frames); the EDEN_NX2_TC_* overrides still
win, and the thresholds are logged only when they move by 256 MiB.
2026-10-29 00:54:03 -04:00
Takoroka
f5f885d3fc nx2:vi2: quiet stubs for 104 and 105
Title 1 calls both ~75 times a second with the display name "default".
The auto stub's reply (success, no data) works; it also logged two lines
per call, which made a tester's 18 minute log 50 MB. Same reply, no log.
2026-10-29 00:54:02 -04:00
Takoroka
a50d4a80ca vulkan: one lock for both caches in InnerInvalidation
It runs before most draws (constant buffer updates go through the
invalidation accumulator) and took the texture and buffer cache mutexes
one after the other.
2026-10-29 00:54:02 -04:00
Takoroka
a4674d242e nx2: memo for compute macros; one lock per inline upload; cheaper shader invalidation
- Compute engine macros replay repeated calls like 3D ones (title 1's
  compute macro path fell from ~27% to ~14% of the GPU thread).
- AccelerateInlineToMemory takes the buffer and texture cache mutexes
  together instead of one after the other.
- ShaderCache::InvalidatePagesInRegion looks at the rejected programs
  only when the write can reach one (bounds kept as they are added);
  it ran a locked scan for every inline upload and pre-draw
  invalidation.

Inline uploads went from ~10% to ~4% of the GPU thread.
2026-10-29 00:45:20 -04:00
Takoroka
3434eababe nx2: macro memo hashes its keys as they are, rewrites re-stored calls in place
The flat map re-hashed the already mixed u64 keys with MSVC's byte-wise
FNV-1a (~1.4% of the GPU thread), and calls stored again after a data
RAM guard failed appended new records, so the arena filled with stale
copies and was cleared every few seconds. EDEN_NX2_MME_MEMO_STATS logs
the counters: in title 1, 97.3% of calls replay, 2.1% fail a guard.
2026-10-29 00:37:06 -04:00
Takoroka
ee9e787fc8 nx2: poll pending channels briefly before sleeping; profile zones on the TSC
- A channel waiting on a semaphore the CPU releases slept 1 ms per poll
  (~25 waits a second in title 1); poll for 0.5 ms first.
- The GPU profile zones read the wall clock (the TSC) instead of
  QueryPerformanceCounter, which cost ~5% of the GPU thread when on.
2026-10-29 00:27:35 -04:00
Takoroka
8d83cef520 nx2: replay repeated native macro calls instead of interpreting them
Title 1 makes ~1,000,000 native macro calls a second in 3D scenes, 7
instructions each on average, and 96% repeat an earlier call exactly.
A call whose effects depend only on its start, its parameters and the
data RAM words it read (no MME_READ, no memory reads, no re-read of
its own data RAM writes, no code upload while it ran) is remembered
with those reads as guards; a later call with the same start and
parameters whose guards still hold replays its data RAM and method
writes. Uploading macro code clears the memo. EDEN_NX2_MME_MEMO=0
interprets every call.

Checked offline against the previous interpreter on 200,000 recorded
3D-scene calls (method writes, registers and data RAM all identical,
94% replayed, 1.9x faster); in the game the macro zone fell from ~0.26
to ~0.20 us per call.
2026-10-29 00:27:35 -04:00
Takoroka
ab80749e94 nx2: present stats report the worst frame and frames over 100 ms; notes 110-116
The per-600-present log only had the average frame time, which hides
shader-compile hitches; it now also gives the longest frame and how
many took over 100 ms.

Notes: today's profile of title 1 and what came of it (TSC on AMD,
rejected-shader rethrows, debug hooks, texture churn from VMA slack,
reactive flushing, macro interpreter).
2026-10-28 23:58:55 -04:00
Takoroka
d5d4b61b5b nx2: EDEN_NX2_MME_RECORD records macro calls for offline interpreter checks
Writes the first 200,000 native macro calls (start, parameters) and the
macro code, so an interpreter change can be replayed against the
previous interpreter (all 200,000 matched for the operand-slot rewrite,
which was then dropped as no faster).
2026-10-28 23:46:54 -04:00
Takoroka
ed9262e642 vulkan: 64 MiB VMA blocks; device memory usage without VMA block slack
Title 1's texture cache deleted 1,400-2,600 textures per 600 frames
and created them again: device memory sat at the collector's critical
line, but ~38% of it was empty space in 256 MiB VMA blocks. Freeing an
image returns that space only once its whole block is empty, so the
caches kept collecting without the usage dropping (one mass purge
deleted 4,300 textures).

- 64 MiB blocks on discrete GPUs too (as on integrated ones).
- GetDeviceMemoryUsage subtracts the unused bytes in VMA blocks, so
  each eviction counts at once.

Now 0-224 deletions per 600 frames, ~4,400 textures resident, and
driver usage 6.24 GB of a 6.97 GB budget.
2026-10-28 23:46:53 -04:00
Takoroka
a89892695e kernel: spin briefly before sleeping on a contended KSpinLock
The scheduler lock is taken by every emulated core thousands of times a
frame (title 1's job system: WaitForAddress / SignalToAddress / yield)
and held for microseconds. As a std::mutex, a contended Lock put the
host thread to sleep at once. Now a three-state futex lock: one atomic
op uncontended, ~128 pause spins when contended, then
std::atomic::wait. Time blocked in Lock fell ~20-25% on title 1's cores.
2026-10-28 23:46:53 -04:00
Takoroka
2671edfe36 nx2: measure the TSC on AMD, no rethrows for rejected shaders, debug hooks off the hot path
- AMD CPUs have no CPUID leaf 0x15, so the TSC frequency stayed 0 and
  every guest CNTPCT read went through QueryPerformanceCounter (~10% of
  the game's main core in title 1). Measure it like Intel CPUs without
  a crystal frequency.
- A rejected native shader was rethrown (two C++ throws) every time it
  was bound; return false instead (~1-2% of the GPU thread).
- The semaphore watch, pushbuffer entry hashing, first-write trace and
  per-macro debug logs ran on every command word or macro call; they
  now cost a flag check unless their environment variable is set
  (~4% of the GPU thread). Also starts the macro engine's profile zone.
2026-10-28 23:07:49 -04:00
Takoroka
51791d0fa4 nx2: GPU-thread zone timers and an in-process sampling profiler
EDEN_NX2_GPU_PROFILE=1 times the GPU thread's stages (command parsing,
macros, draw setup by step, compute, texture uploads) and logs each
one's share of every second.

EDEN_NX2_SAMPLE=<thread names> (Windows) pauses the named threads about
600 times a second for EDEN_NX2_SAMPLE_SECONDS, walks their stacks and
writes per-function counts plus folded stacks (EDEN_NX2_SAMPLE_OUT).
2026-10-28 23:07:49 -04:00
Takoroka
8c0579e555 nx2:gpu: free GPU-only region memory when the region is destroyed
Sub 215 dropped the region's GPU mapping but kept the kernel pages 214
had allocated for it, and their SMMU mapping (cached per CPU range).
Title 1 creates and destroys these regions as scenes load, so the pool
ran dry within minutes and later regions got no memory at all (37
allocation failures in 85 s, 165 in a longer run): GPU writes to them
were lost.

Destroying a region now drops its GPU mapping, then (through a new
registry release step, in GPU-thread order) each address space's SMMU
mapping and cache entry, and frees the pages once the GPU thread is
past both. A 4-minute run destroyed 116 regions with no later failures
and renders cleanly; only the 5.8 GB virtual texture regions and three
170-340 MB regions at startup still do not fit the pool.
2026-10-28 22:24:39 -04:00
Takoroka
dcf6212ddd nx2: video decoder service protocol (no decoding yet)
Read from the SDK's nn::codec VP9 decoder: the decoder service session
(vdec TIPC 1) takes 200 initialize (work buffer transfer memory), 101
configure (X: codec configuration, width, height), 3 its submission
event (copy handle at word 3), 6/207/8 start/stop/flush, 115 finalize,
311/314/212 register/release/unregister an output image (reply a
16-byte NativeSyncHandle; all zero reads as signalled), 109 submit a
frame (A: bitstream, u64 pts, u8 end of stream) and 213 the next
finished frame (image id, size, colour space, pts, end of stream).
Output images come from gsv 911, the same request as 429.

The service completes each submitted frame at once into the next free
image (no picture yet) and keeps its event signalled. Title 1 now sets
the decoder up without aborting; in the runs so far it has not submitted
frames yet. Auto-stubbed requests on nx2:vdec* sessions log 48 words.
2026-10-28 22:05:08 -04:00
Takoroka
8965d72141 nx2: present no longer waits for its own frame; cancelled slot events signal
- vi2 229 now waits for the previous frame's fence by default, so the
  game prepares frame N+1 while the GPU thread runs frame N. Waiting for
  the current frame serialized them: 8-56 ms of each 34-105 ms frame was
  spent blocked. Title 1 now blocks under 1 ms per frame on average and
  its GPU thread runs ~90% busy, which is the next limit.
  EDEN_NX2_FRAMES_AHEAD=0 restores the old wait.
- syncpt 202 (cancel a slot request) signals the slot's event when the
  request was still pending. NvOsEventPool waits on it again after its
  ~100 ms wait times out; title 1 hung there forever whenever a mid-frame
  wait outlasted 100 ms with the GPU thread a frame behind (2 of 4
  runs). Two 185 s runs with 2-3 cancels each completed.
- EDEN_NX2_VSYNC_HZ=<rate> sets the vi2 vsync event rate (experiment;
  240 Hz did not shorten title 1's frames).
2026-10-28 21:22:29 -04:00
Takoroka
3a3cdd4e39 Notes: rows 105-109 (video sessions, interpreter, GPU ASTC transcode, present serialization, job-system SVCs) 2026-10-28 20:43:46 -04:00
Takoroka
65a52cffa5 nx2: video decoder sessions; present waits on an event; present timeline
- gsv 901/1001 succeed and 902/1002 return sessions; their TIPC 1
  (CreateVideoDecoderService) returns another, auto-stubbed. Title 1's
  movie player no longer aborts with 2001-0114 (no frames are decoded).
- The present's wait for its fence is woken by the GPU thread through a
  syncpoint action instead of 200 us sleeps (~1 ms each on Windows).
- Every 600 presents, log the average frame time and how much of it the
  present thread spent blocked on the GPU fence, compositing and pacing.
2026-10-28 20:43:01 -04:00
Takoroka
1522ce5e43 ASTC: decode and re-compress to BC1/BC3 on the GPU; texture cache budget
- accelerate_astc=Gpu with astc_recompression=Bc1/Bc3 now decodes ASTC
  with the GPU decoder into a temporary RGBA8 image (dedicated memory,
  freed once the GPU is past it) and compresses it with a new compute
  shader (bc_encoder.comp: bounding-box endpoints along the colours'
  covariance, inset 1/16, nearest palette entry; BC3 alpha ramp; BC1
  punch-through), then copies the blocks into the image. No CPU decode,
  so streamed textures are ready the frame they are first used instead
  of showing placeholders for seconds; the decode pass no longer waits
  for the GPU on this path.
- The CPU path (CpuAsynchronous) fills a new image with zero blocks
  before its first decode instead of showing stale VRAM, and the
  transcode pool is a quarter of the hardware threads.
- Switch 2 titles: collector thresholds near the device budget, re-read
  every 60 frames from VK_EXT_memory_budget (another game shrank it from
  7.0 to 5.4 GB).
- EDEN_NX2_HIDE_WINDOW=0 now leaves the window visible.
2026-10-28 20:43:01 -04:00
Takoroka
9686a8f94b Native MME: decode each instruction once; reuse the parameter FIFO
The Turing macro interpreter re-extracted every field of each 96-bit
instruction on every execution and copied the register file per
instruction; title 1's GPU thread spent about half its busy time there.
Instructions are now decoded once per PC (re-decoded when their words
change) and both lanes' operands are read up front instead of copying
the registers. The parameter FIFO is a member instead of a vector
allocated per macro call. Rendering checked unchanged on the attract
loop; the GPU thread went from ~68% to ~45-50% busy.
2026-10-28 20:43:00 -04:00
Takoroka
69e7033d82 Texture cache: stop the churn, the per-frame image leak and garbage textures
- Unmapped GPU addresses get a stand-in CPU address fixed by the GPU
  address. A fresh one per lookup meant FindImage never found an image
  there again: title 1's bloom chain at 0xf0000-0x15b600 was re-created
  on every use (~5,400 images per 600 frames), filling VRAM.
- ASTC whose blocks carry no alpha (exact endpoint-mode scan reusing the
  decoder's block parsing) is re-compressed to BC1 instead of BC3, half
  the size; about 70% of title 1's ASTC qualifies. Copies between BC1
  and BC3 hosts are skipped (their guest data is uploaded instead).
  EDEN_NX2_ASTC_BC1=0 turns it off.
- Deleting an image cancels every decode still queued for it. An image
  re-decoded mid-decode has two; the first to finish cleared IsDecoding,
  the second was not cancelled and was uploaded into the next image to
  take the slot: garbage textures (the title logo, road patches) and a
  staging overrun when that image was smaller.
- Diagnostics: images created and deleted per 600 frames by shape,
  collector deletions, async decodes, BC1/BC3 split; collector
  thresholds overridable (EDEN_NX2_TC_EXPECTED_MIB/CRITICAL_MIB); GPU
  ASTC decode wait timing (EDEN_NX2_ASTC_FLUSH=1 submits without it).

Title 1 attract loop: device memory 4.8-5.1 GB instead of 5.9-6.1 GB, no
decoded textures evicted, ~17 -> 22-23 fps.
2026-10-28 19:32:48 -04:00
Takoroka
eb9f32a073 eden-cli: crash handler prints the faulting stack first; log close requests
MiniDumpWithIndirectlyReferencedMemory left an empty file on Eden's
address space, so print the exception, address and stack as
module+offset before writing a smaller dump, and report a failed write.
A window close or quit request is now logged, to tell it from a crash.
2026-10-28 19:32:48 -04:00
Takoroka
525929363f nx2:gsv 991: size decoder image pools (title 1 movie abort)
nn::codec::CalculateDecoderImagePoolBuffer sends 991(preference, count)
and reads a size and an alignment. Auto-stubbed, the reply kept the
request words, the frame size read as 0x167027f'00000001 and title 1's
movie player aborted on its budget check when its first movie started.
Answer a block-linear NV12 size (4 bytes per pixel for other layouts).
2026-10-28 19:32:48 -04:00
Takoroka
669dcfacee Texture cache: log VMA blocks per memory type with the periodic memory log
Device memory use is ~3 GB above what the images account for; this shows
where. The per-decode log drops to debug level.
2026-10-28 18:36:20 -04:00
Takoroka
cfb9245911 eden-cli: EDEN_NX2_CRASH_DUMP writes a minidump on unhandled exceptions
A debugger changes the timing enough to hide some crashes; this catches
them at full speed. dbghelp.dll is loaded only when a dump is written.
2026-10-28 18:36:20 -04:00
Takoroka
9a86143287 Notes: row 102 resolved (async decode use-after-free) 2026-10-28 17:35:05 -04:00
Takoroka
5505e1ddae Texture cache: fix use-after-free in async texture decode
TickAsyncDecode locked a finished decode's mutex, erased the context (freeing
the mutex) and only then let the lock go out of scope, unlocking freed
memory. With ASTC re-compressed on the CPU (accelerate_astc=CpuAsynchronous)
title 1 hit this a few minutes in: cdb reported "Free Heap block modified
after it was freed" at the mutex offset, and a run without the debugger
crashed. The lock is now released before the erase.

Also, deleting an image while it is being decoded (deletions other than the
collector, which already skips them) cancels the decode instead of
uploading into a freed slot.

Verified under cdb past 330 s (failed at 243 s twice before).
2026-10-28 17:34:48 -04:00
Takoroka
9677090712 nx2 services: callbacks no longer outlive their sessions
Under cdb title 1 reported "Free Heap block modified after it was freed"
when it recreated its display session (~4 min in); a run without the
debugger crashed at 113 s. The vi2 vsync timer callback (HostTiming thread)
captured the session and signalled its events; the destructor unscheduled
the timer (which does not wait for a running callback) and closed the
events. The syncpt slot actions (GPU thread) had the same race with
CloseEvent, as did the frame-capture and dump callbacks.

vsync events now live in state shared with the callback and are emptied
under its lock before closing; syncpt actions check a shared "alive" flag
under a mutex the destructor takes; capture/dump callbacks hold no session
pointers.
2026-10-28 17:22:54 -04:00
Takoroka
eb3edd5a84 Notes: rows 99-102 (VRAM/BC3, transcode pool, Medal layer, host crash) 2026-10-28 17:16:06 -04:00
Takoroka
d30b8dcef3 ASTC transcode pool: an eighth of the hardware threads (at least 2)
With half (8 on a Ryzen 7 8845HS), re-compressing streamed ASTC textures to
BC3 took ~3.5 cores away from the six emulated cores and the GPU thread.
Title 1 attract loop, 65-150 s: 8 threads 16.0 fps, 4 threads 12.6, 2
threads 18.9 (scene noise is large). EDEN_NX2_TRANSCODE_THREADS overrides.
2026-10-28 17:15:59 -04:00
Takoroka
81e7bc4112 ASTC: re-compress to BC3 except huge arrays, which stay RGBA8 and trimmed
Title 1 has ~700 ASTC 8x8 textures (~280 MB as stored); decoded to RGBA8
they took ~4.5 GB, pushing Eden past the 8 GB of an RTX 4070 Laptop
(8.4-9.8 GB in use) and the texture cache into permanent aggressive
collection. With astc_recompression=Bc3 (CPU, asynchronous) they fit
(4.6-5.6 GB). The two 8192x8192x68 terrain arrays would still not fit as
BC3, so ASTC 2D arrays over 1 GiB as RGBA8 always use the GPU decoder with
trimmed levels (RGBA8 host image, UNORM/sRGB views), and the decoder pass
now exists whenever the GPU lacks native ASTC.

Also logs every 600 frames: device memory against the collector thresholds,
the largest images, totals by format, and buffer cache totals.

Title screen 65-150 s: 10.1 -> 16.0 fps.
2026-10-28 17:08:43 -04:00
Takoroka
b661f30210 Notes: title 1 rows 92-98 (colour cast, terrain, sin/cos, movie abort, pop-in, frame rate) 2026-10-28 16:44:43 -04:00
Takoroka
e69090383e Performance diagnostics and small wins for title 1
- eden-cli: SDL called SDL_AppIterate back to back, keeping a host core at
  100% polling input; run the main callbacks at 240 Hz.
- EDEN_NX2_IDLE_SPIN_US: optional spin before an idle core blocks (measured:
  no gain on title 1, so off by default).
- EDEN_NX2_SVC_TRACE=<first>:<last>: trace SVCs (with thread ids) for a range
  of presents only.
- MME DMA reads flush GPU-written memory first (never triggered on title 1,
  kept for correctness); render-enable override histogram instead of the
  first-400 log.

Measured on title 1's title screen (RTX 4070 Laptop): GPU ~20% busy, the six
emulated cores ~0.9 host cores in total, Eden's GPU thread 45-70% busy and
largely waiting: on host fences for semaphore acquires (GPU accuracy High
delays every release to host completion), on ASTC upload Finish(), and in a
1 ms poll sleep. gpu_accuracy=Low: ~12.5 -> ~16 fps.
2026-10-28 16:44:24 -04:00
Takoroka
64c39aa0b7 Switch 2 shaders: MUFU.SIN/COS take turns, Eden's Maxwell MUFU radians
SM70+ removed RRO: compilers multiply the angle by 1/2pi (FMUL.RZ by
0.15915494) and MUFU.SIN/COS compute sin/cos(2pi x). Eden's Maxwell MUFU
computes sin/cos of radians (its RRO is a plain move), so every lowered sine
and cosine used the wrong angle (wind animation, rotations, the ocean FFT
twiddles). Scale the operand by 2pi first.
2026-10-28 15:44:48 -04:00
Takoroka
b8a07d6d5b Keep huge ASTC arrays from a lower mip on the host (title 1 terrain)
Title 1's terrain virtual textures are 8192x8192x68 ASTC 8x8 sRGB arrays with
12 levels: 1.5 GB in guest memory (fully backed) but 23 GB each as the RGBA8
Eden decodes ASTC to, so they were created unbacked and the near terrain
sampled black. A GPU-decoded ASTC 2D array over 1 GiB decoded now keeps only
the levels from the first one that fits in 512 MiB (here 3: 1024x1024, 380
MiB); host mip 0 is that guest level. Uploads read each layer's tail of the
mip chain into a compact buffer and decode only those levels, views shift
their level range, and these images are never downloaded. Normalized
coordinates and derivative LOD address the smaller image correctly; the
nearest terrain is lower resolution than on the console.

Also logs how unbacked images are used (views, render targets, blits, DMA).
2026-10-28 15:37:34 -04:00
Takoroka
6d7acfb990 Fix the P6 scratch save in R2UR/MATCH/HFMA2.RELU; GPU memory dump tool
The colour cast on title 1 (magenta/cyan, then red/green, changing per
course) came from the sky/atmosphere volume shader (a pixel shader drawing
a 32^3 RGBA16F table slice by slice), which starts with an unguarded R2UR.
R2UR, MATCH and HFMA2.RELU use P6 as scratch and saved/restored it with the
uniform-guard mask, 1 << (6 - 7) for an unguarded instruction: undefined, in
practice bit 31, which shifted into the opcode bits and turned the save into
another instruction. They now always save and restore P6.

Debugging aids used to find it:
- EDEN_NX2_DUMP=<present>:<addr>:<size>,...: flush GPU memory ranges at a
  present (on the GPU thread) and save them next to the frame captures;
  nx2work/vol3d.py untiles Switch 2 3D textures from such dumps.
- GPU::RunOnGpuThread.
- Logs of Switch 2 3D images and descriptors.
2026-10-28 15:19:44 -04:00
Takoroka
2d48fb6539 Jump-table programs end at their self-branch; aligned wide zero stores
- With a BRX the jump targets live in a constant buffer, so capture and
  lowering keep going past EXITs to the compiler's closing branch-to-self
  (a title 1 compute program jumped to 0xae0, just past the first EXIT).
- 128-bit stores of RZ use R248-R251 (Eden wants 4-aligned data registers).
- The rejected-program cache covers environment construction, where native
  programs are lowered (it never applied before), and logs the stage slot.

Title 1 reaches its title screen with the ground drawn.
2026-10-28 14:50:32 -04:00
Takoroka
410ef99109 IMAD.SHL as a shift; EDEN_NX2_REJECT; log where rejected programs are bound
- IMAD by a power of two lowers to SHL imm, the form Eden's jump-table
  tracking expects (a title 1 compute program failed "track indirect branch").
- EDEN_NX2_REJECT=<hex addresses>: treat those native programs as rejected,
  to bisect a behaviour change without rebuilding.
- A rejected graphics program logs its stage slot, header and slot registers
  once (for the task/mesh pipeline setup).
- Notes: rows 87-92, title 1 status.
2026-10-28 14:44:18 -04:00
Takoroka
033012cb0a Switch 2 shaders: LEA with a negated addend (ISCADD bit 48) 2026-10-28 14:39:20 -04:00
Takoroka
a9f07b419b Switch 2 shaders: compute forms (LDS uniform base, ATOMG, RED, STG.U16) and more
- LDS [Ra.Xn + URb]; ATOMG.ADD at a register address; RED at [URa+offset];
  8/16-bit STG; REDUX.MAX; R2UR and REDUX into any uniform register.
- UIMAD immediate; F2I from a uniform register; FSET.BF combined with a
  predicate; TXQ single component to an odd register; SHF.L.W.HI of an
  immediate by a register.

Title 1's dumped programs: 1654 of 1667 lower (the rest: 12 task/mesh
programs using ISBERD/ISBEWR, and one SURED).
2026-10-28 14:36:54 -04:00
Takoroka
27301da553 Cache rejected native shaders and log each rejection once
A rejected program was lowered again (and logged) on every draw or
dispatch: 1.7M log lines in 13 minutes of title 1. The shader cache keeps
the first exception per address until guest memory there is written, and
the error logs are deduplicated.
2026-10-28 14:31:00 -04:00
Takoroka
fcfcb5049f Switch 2 shaders: calls, uniform predicates, wide multiply and more forms
- RET: accept any copy of the call link into the return register in the
  function (MOV or IMAD.MOV); library functions may EXIT; @Pg EXIT Pc.
- PLOP3 with any table (SEL to 0/~0, LOP3 sets the predicate), including a
  uniform third source; UPLOP3 (with a uniform guard); UP2UR; ULOP3 with a
  uniform predicate output and guard.
- UIMAD.WIDE.U32, uniform USHF, SGXT.U32, FMNMX with a uniform source,
  I2F.U8 from a uniform register, I2F.U16 (.H1), F2I.FTZ floor/nearest,
  TXQ width/height masks, FSWZADD any swizzle, PRMT any byte selection.
- LEA.HI with an immediate addend; 0x923 is ISBERD, not PIXLD (removed).

Title 1's dumped programs: 1379 -> 1446 of 1473 lower.
2026-10-28 14:31:00 -04:00
Takoroka
a4b40f7ba3 Switch 2 shaders: IADD3 negated third source, LEA.HI with addend, uniform SHF, UR16-UR27
- IADD3 with a third source negated (bit 75).
- LEA.HI Rd, Ra, Rb, RZ: Rb plus the high word; bit 73 sign-extends.
- SHF with a uniform shift amount (0xc19) and SHF.R.U32 of the low word.
- Uniform registers UR0-UR21 in R216-R237 and UR22-UR27 in R240-R245 (pairs
  stay adjacent); ULDC and uniform operands use the mapping.
Title 1: 1359 of 1434 dumped programs lower.
2026-10-28 14:00:31 -04:00
Takoroka
2f657ac0a1 Vulkan: unbacked sparse images for textures too large to back
Title 1 samples an 8192x8192, 12-level, 68-layer array (a sparse virtual
texture, ~23 GB): vkCreateImage failed and Eden aborted, or VMA put it in
system memory and staged a full upload (34 GB private). Images over 2^28
texels are now created sparse-resident with no memory bound; the texture
cache neither uploads nor downloads them. Images merely over the VRAM
budget are allocated without the budget limit instead of failing.
2026-10-28 13:55:29 -04:00
Takoroka
bb0eb5ec48 Switch 2 shaders: centroid IPA, 8/16-bit LDG, UP0-UP6, PIXLD.COVMASK, USHF.R, 2D_ARRAY surfaces
Title 1 (768 dumped programs): 386 lowered cleanly, now 701.
- IPA: bits 76-77 are the sample mode (most varyings are centroid), mapped to
  Maxwell bits 52-53.
- LDG: U8/S8/U16/S16 sizes as Maxwell's.
- Uniform predicates UP0-UP6 (UISETP outputs, USEL/VOTEU, uniform execution
  guards) in R239 bit 6-n, swapped into P(6-n).
- PIXLD.COVMASK reads all-ones (no sample mask input; a mask write ANDs it).
- USHF.R.U32 of the low or high register.
- SULD/SUST 2D_ARRAY (dimension 4, as Maxwell's).
2026-10-28 13:55:29 -04:00
Takoroka
dae499692d nx2: notes on title 1 reaching its title screen 2026-10-28 13:23:14 -04:00
Takoroka
9e3f5cea40 Switch 2: nx2:gpu 215 DestroySparseRegion; log 214 arguments
nn::gpu::sdk names 214 CreateSparseRegion and 215 DestroySparseRegion.
215 (range in words 3-6, job slot in word 8) drops the GPU mapping and
completes its job behind the GPU; it was auto-stubbed and never signalled.
2026-10-28 13:22:59 -04:00
Takoroka
4fc2ffa89b video_core: 64 GiB device address space
Switch 2 titles map most of their memory twice plus GPU-only regions;
title 1 ran out of the 16 GiB SMMU space. The buffer cache page table and
memory tracker cover the wider space.
2026-10-28 13:22:59 -04:00
Takoroka
e1cc47ac43 Switch 2: per-channel MME data RAM; one QMD table for all channels
Macro code stays shared by the channels of a process, but each channel
now has its own MME data RAM. Title 1's channels keep their own QMD ring
base in data word 0x27; with one shared RAM every channel used the last
channel's ring, QMDs collided, releases came early or never and the GPU
waited on semaphores (4 s timeouts per frame).

The QMD scheduler state is one table for all channels: a QMD's
dependency counter is state of the QMD, scheduled and decremented from
different channels.

Diagnostics: EDEN_NX2_QMD_TRACE (scheduler loads, counts, launches) and
EDEN_NX2_MME_DUMP (compute macro code, calls, data word 0x27 changes).
2026-10-28 13:22:59 -04:00
Takoroka
862880b400 nx2: notes on title 1's vsync hang and what it hits next 2026-10-28 12:31:54 -04:00
Takoroka
ed97c9b9dc debugger: monitor handles lists each handle's type and service event name
gdbrsp.py gains raw and monitor commands.
2026-10-28 12:31:45 -04:00
Takoroka
508f3400bd Switch 2: signal nx2:vi2 31/32 events on every vblank
Title 1's main thread waits on the vi2 32 event until it has fired twice
before it submits GPU work; Eden's stub never signalled it and the title
hung before its first frame. Signal the 31 and 32 events at 60 Hz.
2026-10-28 12:31:44 -04:00
Takoroka
da234f31e6 nx2: notes on the GPU-only memory fix and title 1's hang 2026-10-28 12:12:26 -04:00
Takoroka
8e7edfaa98 Switch 2: back GPU-only regions outside the application pool
nx2:gpu 214 regions are memory the graphics server keeps itself, but Eden
took their pages from the application pool: title 4's six 300 MB regions
left it 1.8 GB less than the 9.55 GiB the kernel reports, all of which
its heap and VAMM plan to use. Add 8 GiB of DRAM to the 12 GB layout for
the non-secure system pool and back the regions from there; the device
memory tracker now covers 32 GiB of physical memory.
2026-10-28 12:03:11 -04:00
Takoroka
98871aae52 kernel: EDEN_NX2_MEM_TRACE logs heap, physical memory and memory info requests 2026-10-28 12:03:10 -04:00
Takoroka
f93966bc22 nx2: notes on the 100 MiB ColPaintBuild failure behind the recurring crash 2026-10-28 11:46:21 -04:00
Takoroka
47bd739011 nx2: notes on the six-core debugger run 2026-10-28 11:35:31 -04:00
Takoroka
08c5ffd642 core: emulate six CPU cores
Switch 2 applications are given cores 0..5 (title 4's npdm) and start a
job worker per core. With four emulated cores, 4 and 5 were folded onto
0 and 1, so two cores ran two workers each. Switch 1 titles keep using
cores 0..3; core 3 stays the shared system core.
2026-10-28 11:29:20 -04:00
Takoroka
cea331e4c7 nx2: notes on which heap the scene's GPU pools come from 2026-10-28 11:14:08 -04:00
Takoroka
d9ee15f545 nx2: notes on the scene heap running out at the recurring crash 2026-10-28 11:05:44 -04:00
Takoroka
0d1fab27e8 nx2: notes on the recurring worker-thread crash; gdbrsp waits without a timeout 2026-10-28 10:49:26 -04:00
Takoroka
c489bae2d0 nx2: read debug switches once; key scripts from a file
Several debug switches were read with std::getenv on every macro call, QMD,
semaphore write or draw. The UCRT getenv locks and scans the environment,
thousands of times a frame; a long EDEN_NX2_KEY_SCRIPT made title 4 crawl
at 2 fps. Each is now read once. EDEN_NX2_KEY_SCRIPT=@<file> reads the
script from a file. gdbrsp.py gains 'hits' (log registers, frame chain and
object words at breakpoints, following the stopped thread).
2026-10-28 07:47:45 -04:00
Takoroka
fcd02f8778 nx2: unmap GPU memory in submission order
NormalUnmapGpuMemory4K and destroying GPU memory are jobs the console's
graphics server completes behind the work already submitted. Run the GMMU
unmaps on the GPU thread after the command lists queued before them
(GPU::RunInOrder) and signal the job's event only then, instead of unmapping
and signalling at once while Eden's GPU thread may still be behind.

Also EDEN_NX2_FRAMES_AHEAD (experimental): wait for the previous frame's
fence at present instead of this one's.
2026-10-28 07:47:45 -04:00
Takoroka
36b6b30452 nx2: ATOMG and local spills with an immediate offset
ATOMG.E.ADD through a uniform pointer and STL/LDL.32 at RZ only took a zero
offset; title 4's loading screen uses [UR + 4] and [RZ + 4]. Maxwell ATOM
has a signed 20-bit offset at bits 28-47, STL/LDL a 24-bit one at 20-43.
2026-10-28 07:47:45 -04:00
Takoroka
a89b3087a9 nx2: method data continues across GPFIFO entries
Title 4 calls a 3D macro with a 6-word IncreaseOnce header, ends the main
GPFIFO entry after the first word and supplies the other five in a 5-word
subroutine entry. The split-method parking (an entry elsewhere was treated
as a call with a clean method state) fed those parameters to the parser as
headers: the macro ran with garbage, later draws read megabytes of unmapped
memory and the host driver timed out at the loading screen after the dock
scene. As on the hardware, the method state now simply continues; the old
guess stays behind EDEN_NX2_PARK_SPLIT_METHODS.

Debug aids: splits are logged at debug level with the words around them,
and a 3D write that interrupts a macro's parameters dumps the last 48 writes.
2026-10-28 07:47:45 -04:00
Takoroka
5046a37c95 nx2: notes for the semaphore, pacing and open-file fixes 2026-10-28 05:55:09 -04:00
Takoroka
c81a146fd5 vfs: keep the open-file cache below the CRT limit
The real VFS kept up to 8192 host files open, which is also the most the
Windows C runtime allows (_setmaxstdio(8192)). A game read from a loose RomFS
folder fills that within minutes (title 4: ~7300 handles after 160 s); from
then on every other open failed with "too many files open" - its next scene
pack, its save file, the shader cache - and the game waited forever on the
load. Cap the cache 512 below _getmaxstdio() on Windows.
2026-10-28 05:55:09 -04:00
Takoroka
42989ad8a9 nx2: pace Switch 2 presents to a 60 Hz vblank
The console flips one frame per vblank and holds the rest in the swapchain,
which caps a game at 60 fps. vi2 229 reports every buffer slot free at once,
so nothing paced title 4: it presented at 90-160 fps and its frame-stepped
timing, audio included, ran ahead. Hold each present until the next vblank,
scaled by the speed limit, without banking time lost to slow frames.
2026-10-28 05:55:09 -04:00
Takoroka
5d27ccd0a0 nx2: semaphore releases latch the acquires they satisfy
Switch 2 channels that wait for an exact semaphore value poll on the console
and see every release. Eden resumes a waiting channel only between command
lists and writes releases when fences signal, often several back to back, so
an acquire-equal could miss its value and wait forever. Title 4's 3D queue
ticks a counter that two other queues wait on value by value; under the
loading after the post-customization save one of them missed a tick and the
game froze (channel 4 waiting for 0x5293 with the counter at 0x5310).

Pending acquires now register with Nx2SemaphoreWaiters before fences run,
and every query-cache payload write and QMD add/max release latches the
waiters it meets at the moment it is written.
2026-10-28 05:41:45 -04:00
Takoroka
fd768d8530 nx2: notes for the freed-memory crash and the ASTC UI fix 2026-10-28 05:16:04 -04:00
Takoroka
034b9cab96 nx2: decode ASTC textures with the Switch 2 GOB layout
The GPU ASTC decoder unswizzles block-linear data itself and never applied
the native (Switch 2) GOB swizzle that the other unswizzle paths use, so
ASTC textures came out with 16-byte blocks swapped wherever the image
varies. Title 4's UI is ASTC (10x5, 4x4, 8x8): portraits, buttons and the
squid cursor had hatched edges and clean interiors. Pass the native tiling
flag to the decoder and swap offset bits 5 and 6 as
block_linear_unswizzle_2d.comp does.
2026-10-28 05:14:02 -04:00
Takoroka
910e56d372 nx2: unmap GPU memory the game frees
nx2:gpu 213 (NormalUnmapGpuMemory4K) and 102 (destroy GPU-accessible
memory) were no-ops, so Eden kept freed memory mapped for the GPU. Title 4
frees its customization-room memory right after saving and reuses it as
ordinary heap; pending GPU->CPU write-backs for the old buffers then landed on
the new objects, and two threads crashed on garbage pointers ~5 s later
(the post-customization freeze).

Unmap the range from every address space (the GMMU unmap drops cached
buffers/images and their pending downloads), keep the registry's mapping
list in step, re-map 211 requests even at the memory's own address, and
drop freed regions from the semaphore mirror table, which every 8 KiB
allocation used to fill for good after 32 entries.
2026-10-28 04:56:16 -04:00
Takoroka
d4dc5af3be nx2: frame-locked input replay and quieter logs
EDEN_NX2_KEY_RECORD records keyboard transitions as "f<frame>:code" against a
global count of guest presents; EDEN_NX2_KEY_SCRIPT replays them on the same
frames (or by seconds, as before), so a played session reproduces exactly.

Debug aids: the PC/LR/SP of the first unmapped reads, NSO load addresses,
EDEN_NX2_IPC_SUMMARY (request counts every 5 s), EDEN_NX2_FORCE_LOD0.
Unmapped-access logging is rate limited and the per-draw/semaphore spam is
demoted, so long sessions stay under the log size cap.
2026-10-28 04:56:14 -04:00
Takoroka
6f38217ed3 nx2: lower what title 4 needs after its intro
- 64-bit adds: IADD3/UIADD3 with a carry-out predicate (unsigned compare of
  the sum) and IADD3.X/UIADD3.X with a carry-in predicate (SEL + add); uniform
  predicates go through R239 like UISETP. Used by a compute pass the game runs
  thousands of times while loading the next area.
- TMML (0xb69) LOD query, 2D: 8.8 fixed-point results as on Maxwell.
- Vertex stores to a[0x5c] (an Ampere per-primitive hint, likely shading rate)
  are dropped; Eden had no attribute for it and refused the pipelines.

All 864 shaders from a 5-minute scripted session past the intro now lower,
except the opt-in tensor-core upscaler (bilinear stand-in).
2026-10-28 03:21:26 -04:00
Takoroka
7ca1e4f86b nx2: notes for the intro cutscene milestone and the next blockers 2026-10-28 03:11:51 -04:00
Takoroka
fe5cb5efdf nx2: title 4's intro cutscene renders correctly
Native shader capture counted CALL.REL.NOINC as a function call, so a program
using it waited for a RET after its EXIT and read on into the following
shaders until a garbage branch target made it fail. NOINC calls push no
return address (the compiler uses them as jumps; the lowering already treats
them so): count them as forward branches.

That shader builds title 4's 3D colour-grading LUT once per scene; with it
rejected the LUT stayed empty and the tonemap turned every frame black. The
intro cutscene now renders lit and graded, with subtitles, through the
bilinear stand-in for the tensor-core upscaler.

Also: rejected captures dump to EDEN_NX2_SHADER_DUMP as rejected-<addr>.bin
and report the offending instruction.
2026-10-28 02:57:31 -04:00
Takoroka
465eeb8598 nx2: title 4's cutscene scene renders; blit constants for fragment shaders
- Render area falls back to colour target 0 when the surface clip is unset
  (Switch 2 titles may never program it); it feeds the synthesized blit block.
- BlitConstantsPass also covers fragment shaders: UV clamp rectangle 0..1
  (0x30-0x3c) and zero dither bias.
- Debug: EDEN_NX2_SKIP_VS=<addr,...>, EDEN_NX2_TRACE_VS=<addr>,
  EDEN_NX2_UPSCALE_SOURCE[_FORMAT], EDEN_NX2_TRACE_AFTER_UPSCALE,
  EDEN_NX2_SAMPLE_AT, skipped-draw reasons, clear tracing.

Showing the raw HDR scene colour through the upscale fallback gives the
correct lit cockpit scene; the game's tonemap pass still outputs black.
2026-10-28 02:45:00 -04:00
Takoroka
7e5b415e08 nx2: stand in for title 4's tensor-core upscaler and its blit constants
- When a compute program is rejected for tensor-core math, blit its input
  image (handle at cbuf0+0x90) to its output (cbuf0+0xa0) with bilinear
  filtering. Title 4 renders at 2560x1440 and upscales to 3840x2160 with a
  DLSS-style network; this keeps the 4K output fed. EDEN_NX2_TENSOR=stub keeps
  the program with HMMA as D = C (for experiments).
- BlitConstantsPass: vertex shaders reading bank 24 (lowered to 17) get the
  console-provided blit block synthesized from the render area (texel size,
  size, UV scale 1, rectangle 0..W-1,H-1).
- Debug: EDEN_NX2_CB24_TRACE, upscaler resource description and render target
  sampling at the fallback.

The scene behind the intro dialogue is still black: the tonemap output that
feeds the upscaler has no colour although earlier targets do.
2026-10-28 01:55:10 -04:00
Takoroka
26abba68f5 nx2: lower every shader title 4 uses in game
- Tensor-core ops for its neural compute pass: HMMA.16816.F16 (SHFL + HFMA2),
  MOVM.16.MT88 and LDSM .MT88. Opt-in via EDEN_NX2_TENSOR: the expansion
  (372 HMMAs) takes the host driver minutes to compile.
- HSETP2 with an immediate half pair; skip @!PT instructions.
- Constant bank 24, read only by three NVN2 fullscreen vertex shaders and never
  bound by the game, reads from bank 17 (also never bound: zeros, as on hw).
- Debug dumps of the vertex stage's bound constant buffers on rejection.

All 214 shaders captured after the menu lower (213 without EDEN_NX2_TENSOR).
2026-10-28 01:18:23 -04:00
Takoroka
1fefe5a339 nx2: title 4 runs past its loading screen without GPU stalls
- Native QMD: a launch whose shader Eden can't translate still completes and
  fires its releases (the console never fails a launch; queues wait on them).
- QMD releases with flag bit 19 are max releases: NVN2 tracks QMD completion
  in a word other releases also advance, so a plain write moved it backwards.
- Ampere lowering: LDSM.16.M88.{1,2,4} (SHFL + LDS), IMAD with a negated
  uniform addend, strong LDG (bit 77), 2D TLD4 gather, RED MIN/MAX/..., AST of
  RZ. 210 of the 214 shaders title 4 uses after its menu now lower.
- Maxwell3D: ignore constant buffer binds to slots >= 18 instead of writing
  past the state array.
- Debug: EDEN_NX2_SEM_WATCH covers QMD and compute-macro releases and logs
  address spaces; EDEN_NX2_3D_METHODS logs the first write to each 3D method;
  constant buffer binds are logged once per stage/slot.
2026-10-28 00:50:13 -04:00
Takoroka
f6cf565bbb nx2: get title 4 past the difficulty confirm
- MME: implement MME_DMA_READ / MME_DMA_READ_FIFOED (Turing macros read query
  results into their load FIFO; the conditional-render macro at 0x142 needs it).
- vi2 present waits up to 4 s for the frame fence (was 250 ms) so shader
  compile stalls don't let the game lap its equal-value gate semaphores.
- Native QMD: fire the third release slot at word 52, which NVN2 patches for
  queue-to-queue syncs (compute queue releasing a gate the 3D queue waits on).
- Compute class SET_REPORT_SEMAPHORE (0x6c0-0x6c3) releases.
- Ampere lowering: AOFFI, any LOP3 truth table, STG RZ, LDG 0x981, IABS,
  I2FP rounding, IMAD negated addend, SHFL 0x989 (all 77 title-4 shaders lower).
- Debug: EDEN_NX2_SEM_WATCH=<addr> traces acquires, releases, method
  arguments and compute macro emits for one semaphore block;
  EDEN_NX2_KEY_SCRIPT scripted key presses; per-submit logs moved to TRACE.
2026-10-28 00:24:14 -04:00
Takoroka
732778b124 Switch 2: present unflipped; fix frame-capture orientation
The vertical flip at present was a misreading: EDEN_NX2_FRAME_CAPTURE wrote the
TGA origin bit backwards (invert_y means the data is bottom-up), so captures
were upside down while the screen was right. Present unflipped again.
2026-10-27 23:10:17 -04:00
Takoroka
b8a278125e Switch 2: title 4 menu text renders upright
- Present flips the image vertically: title 4 renders with a lower-left origin
  (negative viewport scale), so its frames are stored bottom-up.
- Lowering: TEX/TEXS FP16 results (bit 79, packed half pairs) - the glyph
  shaders unpack them with HADD2.F32; LDG.E with a register address;
  SHF.R.U32 low word with a zero high word.
- Debug aids: EDEN_NX2_TRACE_PS / EDEN_NX2_WINDOW_TRACE_SKIP for the window
  trace, EDEN_NX2_TEXTURE_DUMP, EDEN_NX2_TRACE_DST for copies; RT trace keyed
  by shader pair. NX2_NOTES: blockers 53-57, next steps.
2026-10-27 23:04:22 -04:00
Takoroka
d754fff776 Switch 2: back GPU-only regions, pace presents, interleaved split methods
- gpu 214 creates GPU-only memory the graphics server backs itself (title 4
  renders everything into six 300 MB regions at 0x81_0000_0000). Back each
  with application-pool pages mapped into the alias code region and map that
  for the GPU. This is what made title 4's frames appear.
- vi2 229 waits (bounded) for the presented frame's GPU fence: the reply
  reports every buffer free, so without this the game laps Eden's GPU thread
  and rewrites command memory it has not read yet.
- Split methods: keep several parked methods keyed by resume address (NVN
  interleaves command buffers) and drop them when the submission ends.
- hid 214 (per-frame vibration send) answered quietly.
- Debug aids: EDEN_NX2_WINDOW_TRACE / EDEN_NX2_RT_TRACE (render targets,
  clears, textures), EDEN_NX2_AUDIO_LEVEL (output peak), EDEN_NX2_PB_DUMP
  (command list for offline replay), parsed-entry change check.
2026-10-27 22:28:10 -04:00
Takoroka
9bee875c48 Switch 2: Turing macros, native compute, Ampere shader lowering, native tiling
From changes.patch: the TU104/GA10x MME interpreter (shared macro context for
the process's channels, also on the compute class), native compute with QMD
3.0 dependency scheduling, the Ampere-to-Maxwell shader lowering, native
vertex stream sizes, and the console's GOB swizzle for textures, DMA copies
and the host unswizzle shaders. Switch 2 channels turn all of it on at init.

The lowering is widened for title 4 (all 32 of its programs lower): general
TEX (any dimension/LOD/mask, operand split 1+1/2+1/2+2/4+n, any but consistent
descriptor buffer), cube/array, SULD channel masks, 128-bit STG, indexed IPA,
VOTE, SHFL with any immediate or predicate output, R2P/P2R byte select,
uniform ULOP3/ULEA/UBREV/UFLO.SH/UISETP signed/UMOV URZ, IMNMX cbuf, IADD3
and FFMA uniform operands, FMUL /4 /8 x8, and BSSY/BSYNC/BREAK/WARPSYNC as
no-ops. The capture stops at unmapped memory. EDEN_NX2_SHADER_DUMP dumps every
native program; tools in C:\dev\nx2work test the lowering offline.
2026-10-27 22:26:16 -04:00
Takoroka
09630c34ed Switch 2: suspend channels on semaphore acquires; fix pushbuffer desyncs
NVN2 gates command buffers with semaphores the CPU releases after writing
them. Eden read past unmet SEMAPHORED acquires, so it executed stale commands
and desynchronised. Channels now suspend on an unmet acquire and resume from
the same word once it is met (scheduler keeps the waiting channels, the GPU
thread polls them), as in changes.patch.

Also: methods split across GPFIFO entries around a called command buffer are
set aside and resumed; subchannel 2 defaults to inline-to-memory; unbound
subchannels and out-of-range compute/DMA methods are dropped instead of
crashing; 4-byte semaphore releases no longer write 16 bytes; Ampere
SEM_PAYLOAD_HI/SEM_EXECUTE no longer spin as Kepler acquires; unmapped GPU
reads log instead of asserting. Debug aids: EDEN_NX2_PB_CHECK (pushbuffer
changed after submit), EDEN_NX2_QUERY_TRACE, a one-time pushbuffer dump.
2026-10-27 20:53:15 -04:00
Takoroka
c43191ee4d Switch 2: present vi2 native images and capture frames
vi2 229 composites the presented native image (block-linear RGBA8 in the
console's GOB layout, from changes.patch) at its SMMU address, which the GPU
memory registry now records per mapped range. EDEN_NX2_FRAME_CAPTURE=<dir>
saves every 120th presented frame as a TGA for hidden test runs. Also mirrors
releases for every two-page gsv semaphore region and logs gsv sub 101 ranges.
2026-10-27 20:53:14 -04:00
Takoroka
0fbf68fa15 Switch 2: vi2 present/cancel replies and syncpt QueryCounter
vi2 229/228 (PresentClientBuffer/CancelClientBuffer) reply with the mask of
slots the client may acquire next; syncpt 103 (QueryCounter) returns the host1x
value of ~counter. Names and reply shapes from changes.patch. Title 4 now runs
its render loop, presenting window images 1-3 in turn (730 GPU submissions in
60 s); nothing is shown yet.
2026-10-27 19:49:20 -04:00
Takoroka
f7c3638d6e Switch 2: native memory/images, GPU-accessible attribute, TSG subcontexts, syncpoint shim
- gsv 111/112/114/423/429/430/314/315/435 track native memory and images with
  real ids (meanings from changes.patch); 1510/435 use its block-linear layout.
- svc 0x80 sets a real GpuAccessible (0x20) memory attribute (from changes.patch).
- Characteristics report TSG subcontexts (flag 22, 64 per TSG at 0xb4); NvRm
  otherwise allows one subcontext and a second NVN queue fails its setup.
- Syncpoint shim: flag 26, address-space ioctl 0xB, and GPU reads of the shim
  page return guest syncpoint values (scheme from changes.patch).
- EDEN_NX2_IPC_TRACE also logs TIPC commands; gdbrsp.py watch/mem improvements.
2026-10-27 19:30:20 -04:00
Takoroka
c066219bc6 NX2_NOTES: keep the blocker table in order 2026-10-27 17:46:15 -04:00
Takoroka
ac3c8d25cd Switch 2 (title 1): syncpt handshake, gpu user/gsv partitions
Newer SDKs open the syncpt session with TIPC 0(u64) and need it echoed, and
query the GPU address space partitions with gpu 201/202 instead of 203. With
these and the larger memory budget, title 1 reaches its main loop and the same
empty-command-buffer stop as title 4.
2026-10-27 17:46:02 -04:00
Takoroka
5fada8e590 Switch 2 GPU memory handles and mappings at requested addresses
nx2:gpu 101 (CreateGpuAccessibleMemory) returns real GpuMemoryHandles; 211
(NormalMapGpuMemory4K) maps handle+offset at the requested GPU address,
including sparse regions, through Nx2GpuMemoryRegistry::MapAt (SMMU ranges are
reused). 102/213/214 complete through their job-slot events. gdbrsp.py waits
for a real thread before driving the stub and gains a write-watch op. Notes
record the current stop: an NVN command buffer with no command memory.
2026-10-27 17:39:54 -04:00
Takoroka
ff4b2bb33f NX2_NOTES: title 4 reaches its main loop; blockers 26-41, next steps 2026-10-27 17:22:25 -04:00
Takoroka
f756fe05f7 Switch 2: aud:u audio, REv revision tag, job-slot events, open-file limit
- aud:u (nn::audio::CreateUserClientSessionByHipc) command 2 returns Eden's
  IAudioRendererManager; the renderer accepts Switch 2's 'REv?' update tag by
  comparing revision numbers.
- nx2:gpu sub 211 signals the job slot in request word 17, not word 2; mapping
  jobs no longer stall once the SDK rotates slots.
- eden-cli raises the CRT open-file limit (_setmaxstdio(8192)) like the Qt
  frontend, so loose RomFS folders don't run out of handles.
- Per-mapping GPU logs moved to Debug.
2026-10-27 17:21:19 -04:00
Takoroka
50de2d4e7e Switch 2: window image sizes, queue semaphores, built-in fonts, memory budget
- gsv 1510 CalculateRequirements: block-linear size/alignment for window images.
- nvhost-gpu: Switch 2 channel event 7 is a real (unsignalled) event.
- Puller mirrors NVN2 queue semaphore releases into the page the CPU and GPU
  wait on (the first two-page nx2:gpu region); EDEN_NX2_METHOD_TRACE logs methods.
- pl:u 2000/2010 GetBuiltInFontSharedMemory/GetBuiltInFontData serve plain TTF
  copies of the shared fonts.
- Switch 2 processes count the alias region in their max memory, so the game's
  nninitStartup heap (total - 6.3 GiB) has room for its sound heap.
- gdbrsp.py: interrupt, threads, selthread, hits, continue_async, sleep.
2026-10-27 17:04:06 -04:00
Takoroka
563e727332 Switch 2 GPU: syncpt server, identity GPU mappings, Ampere classes, submit fences
- gsv 204: real syncpoint shared page (values mirrored from host1x, armed waits
  acknowledged by a 500us poller); first handle is the writable request event.
- syncpt 22/201/202: per-slot events signalled through host1x guest actions.
- nx2:gpu sub 101 memory is identity-mapped (GPU VA == CPU VA) into every
  nvhost-as-gpu GMMU via Nx2GpuMemoryRegistry; SMMU physical range widened to
  34 bits for the 10/12 GiB layouts.
- Puller runs Ampere 3D/compute/copy classes on the Maxwell engines; class 0 unbinds.
- Turing+ macros are skipped (Eden only has the Fermi macro ISA).
- Submits with increment_value get Eden-side syncpoint increments, since NVN2
  signals completion with semaphore releases only.
2026-10-27 16:19:35 -04:00
Takoroka
0d679f770e nvdrv: Switch 2 GetEngineInfo, TSG flag, Ampere object classes, channel ioctl passthrough
Once an nx2 NvRm identifies itself (0xD0-byte characteristics), nvhost-gpu and
nvhost-ctrl-gpu answer unknown ioctls with a logged input echo instead of
NotImplemented, ctrl-gpu reports GR + async-copy engines, and Ampere class numbers
(and class 0) map onto Eden's Maxwell object slots. gdbrsp.py gains trace/flow/
breakwhen/fp helpers.
2026-10-27 15:54:37 -04:00
Takoroka
fb7323f25c nvdrv: stub device for nvgpu nodes Switch 1 never had
Unknown /dev/nvhost-* nodes (first: /dev/nvhost-tsg-gpu, opened by the
Switch 2 NvRm during NVN device init) open as nx2_stub_device: every
ioctl succeeds with the input echoed and the rest zeroed, and is logged.
With this NVN device init completes and the title reaches
nn::wsi::ounce::native::NativeWindow::Acquire (first frame).
2026-10-27 15:27:35 -04:00
Takoroka
7cb5831340 Switch 2 GPU memory: svc 0x80, gpu sub-session jobs, VM range, SM 0x808
- svc 0x80 (address, size, mask 0x20, attr 0x20) marks memory GPU
  accessible for nn::gpu::sdk::CreateGpuAccessibleMemory; accept it.
- gpu 22's sub-session is served by the gpu object; sub 101 registers
  memory, sub 211 submits a job whose id event is signalled at once.
- gpu 203 reports [0x8000000, 0xC000000000) as the in-process
  VirtualMachine range that serves AllocSpace (NVN reserves fixed VA at
  0x8000000 up to 512 GiB).
- Report SM 0x808: NVN2's engine class table is only returned for it.
With this NVN device init allocates its GPU memory and opens
/dev/nvhost-tsg-gpu.
2026-10-27 15:17:39 -04:00
Takoroka
3bc0c71884 NX2_NOTES: graphics boot through vi2 and NVN device init 2026-10-27 14:58:15 -04:00
Takoroka
ff097d4c30 nvdrv: Switch 2 GPU address-space layout; gpu 33/203
- AllocAsEx from Switch 2 NvRm passes flags 5, no big page size, range
  [0x8000000, 0xC000000000) and split 0 (small pages only). Treat split 0
  as small pages up to the end, keep a 4 GiB big-page region after it, and
  give that allocator an absolute limit (Eden's length-as-limit breaks once
  the split is past the region length).
- gpu TIPC 33 returns a count, 203 a kind plus a VA range used by
  nn::gpu::sdk::CreateGpuAddressSpace.
- gdbrsp.py: unbreak, step, breaklr, unbreakpc, fp (frame walk), trace.
2026-10-27 14:57:51 -04:00
Takoroka
b2c084edcc nvdrv: serve Switch 2 nvsrv with Eden's nvdrv; Ampere characteristics
- gsv 1402's nvsrv session speaks the nvdrv INvDrvServices protocol;
  serve it with an NVDRV on the shared module, initialised for the
  calling process (Switch 2 doesn't call Initialize).
- GetCharacteristics3 callers asking for the 0xD0-byte Switch 2 layout
  get chip 0x17E (Ampere; the Switch 2 NvRm rejects Maxwell for NVN),
  Ampere classes, and s16 interface versions at 0x70..0x80.
- After that, unknown nvhost-ctrl-gpu Ioctl3 commands from such callers
  return a zeroed success (logged).
2026-10-27 14:28:53 -04:00
Takoroka
5a1f99fee6 tools/nx2: minimal GDB remote client for Eden's gdbstub 2026-10-27 14:13:03 -04:00
Takoroka
be9435e691 service: correct TIPC data replies; vi2 timestamp recorder count
ResponseBuilder::Push(Result) writes the CMIF two-word result, so TIPC
data pushed after it landed one word late and outside the copied size;
clients read zeros. TipcReplyWords writes the one-word TIPC result and
the data. vi2 TIPC 5 returns 7 (LayerHistory needs 7 recorders). With
this the title creates its layer and goes on to hid/acc/aoc and RomFS
asset loading.
2026-10-27 14:13:03 -04:00
Takoroka
ccd3885a62 service: Switch 2 vi2 placeholder session (gsv 305)
gsv 305 returns the session nn::vi2::sdk::Initialize keeps. Reply shapes
read from the SDK: TIPC 1 -> u32 < 0x400, 3 -> u32 < 8, 11/31/32 -> one
handle then result (events here). With this nn::vi::Initialize returns and
the title opens the default display and starts creating a layer.
2026-10-27 14:00:36 -04:00
Takoroka
f493daaac2 service: Switch 2 GPU placeholder session (gsv 497)
gsv 497 returns the session nn::gpu::sdk::BootWithRemoteMemory keeps.
On it: TIPC 11 returns an id range, 12(id) returns that id's event, 21(id)
starts an operation the SDK then waits for (signalled at once here), 22
returns another session. With this nn::gs::sdk::BootGraphicsSystem
returns and boot moves on to nn::vi.
2026-10-27 13:56:54 -04:00
Takoroka
a0364c2c46 NX2_NOTES: classification of all titles, graphics boot progress 2026-10-27 13:48:52 -04:00
Takoroka
b17bb70a42 eden-cli: EDEN_NX2_HIDE_WINDOW for unattended test runs
Creates the render window hidden and not focusable so test runs don't
cover or interrupt the desktop.
2026-10-27 13:48:52 -04:00
Takoroka
6e57c2f643 service: Switch 2 nvsrv and syncpt placeholder sessions behind gsv
- gsv 1402 returns an nvsrv session (handle at word 3, result at word 4);
  nn::nvsrv::BootWithRemoteMemory keeps it.
- gsv 202 returns a syncpt session; gsv 204 takes two transfer memories
  (0x1000 RW, 0x2000 RO) and must leave a header in the RO one (version 1,
  count, five region offsets) and return two event handles.
- syncpt TIPC 1 returns a u32 < 0x400 (32 here).
RequestParser doesn't skip TIPC handle words, so 204 reads raw data words.
2026-10-27 13:48:52 -04:00
Takoroka
d2b28c7620 am: GetPseudoDeviceId without control metadata
Loose ExeFS boots have no control.nacp, so GetPseudoDeviceId returned
ResultUnknown and the title aborted. Seed from the program id instead.
2026-10-27 13:36:58 -04:00
Takoroka
9f1d8357c0 am: placeholder Switch 2 graphics-system (gsv) session
ISelfController 2000 hands nn::gsv::BootWithRemoteMemory a raw session
handle it then speaks TIPC on (4, 11, 101, 1501, 1502, 1402 during boot).
Return a session to INx2GraphicsSystemSession, which answers every
command with an empty success and logs it. ReturnRawSession creates the
session outside the caller's domain, since gsv wants a real handle.
2026-10-27 13:32:41 -04:00
Takoroka
f7be327f97 kernel: stub Switch 2 svcGetInfo 31/32 (remote memory)
nn::os::GetRemoteMemoryUsage returns {id 31, id 32}; nn::gsv and
nn::nvsrv BootWithRemoteMemory use it while booting the graphics
system. Return 0 for both for now (logged).
2026-10-27 13:32:40 -04:00
Takoroka
b10c8887b5 tools/nx2: NSO symbol lookup; skip undecodable modules in coverage 2026-10-27 13:26:46 -04:00
Takoroka
4e5274820d service: EDEN_NX2_IPC_TRACE and EDEN_NX2_AUTO_STUB switches
IPC trace logs every dispatched command and non-zero result words.
EDEN_NX2_AUTO_STUB auto-stubs unknown commands without relying on the
use_auto_stub setting, which eden-cli's config rewrite resets.
2026-10-27 13:26:46 -04:00
Takoroka
3f5bdb81f9 am: Switch 2 graphics-scale interface and IApplicationFunctions 321
IApplicationProxy 60 returns the interface nn::oe::GetCurrentGraphicsScale
uses; its command 10 must return 10, 20 or 30, so return 10. Other
commands on it are auto-stubbed (nx2: prefix). IApplicationFunctions 321
is called right after 320; stub it.
2026-10-27 13:26:44 -04:00
Takoroka
1a54a1aea7 loader: give loose ExeFS boots a RomFS and an fs registration
The ExeFS directory loader never set its romfs member and, when used on
its own, never registered the process with the FileSystemController, so
fsp-srv SetCurrentProcess failed with ResultTargetNotFound (2002-1002).
Build the RomFS from a romfs/ directory next to exefs/ (or a .romfs file
in exefs/), and register the process when the loader isn't wrapped by
the NCA/NSP loaders.
2026-10-27 13:26:37 -04:00
Takoroka
975c0d66ca Add NX2_NOTES.md and Switch 2 analysis helpers
npdm_info.py dumps META/ACI0 and kernel caps; nso_dis.py disassembles
part of an NSO with clang + llvm-objdump.
2026-10-27 12:58:04 -04:00
Takoroka
640b42c818 am: stub IApplicationFunctions 320 (called by Switch 2 titles) 2026-10-27 12:57:38 -04:00
Takoroka
a9bd43b9d1 kernel: optional SVC trace (EDEN_NX2_SVC_TRACE=1) 2026-10-27 12:52:40 -04:00
Takoroka
fb516c2ce2 kernel: fold Switch 2 virtual cores 4/5 onto emulated cores
Switch 2 titles get core mask 0x3f and nn::oe creates a thread on core 5.
Eden emulates four cores, so fold virtual cores 4..7 onto 0..3 in
CreateThread and SetThreadCoreMask (after validating against the real
process mask).
2026-10-27 12:52:39 -04:00
Takoroka
6a58a33930 am: Switch 2 appletOE/IApplicationProxy entry points
- appletOE command 10 opens the application proxy (same arguments as 0).
- IApplicationProxy commands 40/50/.../90 are new getters; return an
  INx2StubInterface (nx2_stub_interface.h). Interfaces whose name starts
  with nx2: answer every call with an empty success and log it.
2026-10-27 12:52:39 -04:00
Takoroka
f8e3da63f5 kernel: svcGetInfo 33 is a page granule; give 29/30 a real region
The SDK checks module address/size against (id33 - 1) when initialising
DSOs, so id 33 is a page granule: return 0x1000. Report ids 29/30 as the
upper half of the alias region (a plausible VAMM address-space region,
backed through MapPhysicalMemory) instead of an empty region.
2026-10-27 12:52:39 -04:00
Takoroka
b71771c691 kernel: stub Switch 2 svcGetInfo ids 29, 30 and 33
nnosInitialize queries ids 29/30 between the alias (2/3) and ASLR (12/13)
regions and aborts if they fail; report an empty region after the alias
region. VammManager::Initialize (enabled when the system resource size is
non-zero) queries id 33 and passes it to MapPhysicalMemory as a size, so
return 2 MiB. Both are logged as stubs; the real meanings are unknown.
2026-10-27 12:20:28 -04:00
Takoroka
6a19f4189c eden-cli: fix SDL3 main-callback lifetime bugs
- SDL_AppInit never stored *appstate, so every later callback got null.
- SDL_AppInit/SDL_AppEvent returned SDL_APP_SUCCESS, which tells SDL3 to
  quit; they must return SDL_APP_CONTINUE.
- InputSubsystem was a local in SDL_AppInit but the window and its timers
  keep a pointer to it; keep it in SdlState, ahead of the window.
- Guard the mouse timer against firing before the mouse driver exists.
2026-10-27 12:00:22 -04:00
Takoroka
b2f11310f5 npdm: parse the Switch 2 ACID layout
Switch 2 ACID blocks have a 0x300-byte signature/key area (Switch 1:
0x200), so the ACID header starts 0x100 bytes later. Without this the
FAC offset was garbage and loading failed with ErrorBadFileAccessControl.
Also read FAC/FAH relative to their ACID/ACI block, as the format (and
the existing KAC read) does.
2026-10-27 12:00:22 -04:00
Takoroka
47a7f5c495 loader: load a loose ExeFS when given its main
'main' is also a valid NSO, so IdentifyFile picked the single-NSO loader,
which skips main.npdm, rtld and the sdk modules and then crashes writing
into an uninitialised process. Prefer the ExeFS directory loader when the
file is named main and sits in an ExeFS.
2026-10-27 12:00:22 -04:00
Takoroka
dadb3e86fd dynarmic: decode ARMv8.1-8.5 instructions used by Switch 2 code
Switch 2 titles are built for a Cortex-A78C class core. Scanning one title's
.text with dynarmic's own decoder table found ~115k RETAB, ~37k LDAPR,
~25k LSE atomics, ~6.6k CAS, ~2.5k SWP plus a few CASP/XPAC/PAC/STG, none of
which dynarmic decoded (only the declarations existed).

- LSE atomics, SWP, CAS, CASP: exclusive load + op + exclusive store; on
  store failure no register is written and the instruction is retried
  (CheckBit terminal), so it stays atomic across cores.
- LDAPR: ordered load, like LDAR.
- Pointer authentication: identity (nothing is ever signed). RETAA/RETAB,
  BRA*/BLRA*, LDRAA/LDRAB, PAC*/AUT*/XPAC*, PACGA.
- MTE: tags are always 0; STG/ST2G only write back, STZG/STZ2G zero the
  granules, STGP stores the pair, IRG/GMI/ADDG/SUBG ignore tags.

tools/nx2/a64_coverage.py reports which encodings in an NSO's .text the
decoder table does not match.
2026-10-27 11:46:04 -04:00
Takoroka
a5011891e1 Fix build with MSVC 14.51 (VS 2026)
- sparse_large_vector: ULONG_PTR -> u64 needs static_cast, not reinterpret_cast
- zbic_compression: silence C4505/C5245 for the amalgamated zstd.c, which
  is compiled as project code and hits /we4505
- launch_params: explicit int -> bool (C4800 is an error here)
2026-10-27 11:46:04 -04:00
208 changed files with 16248 additions and 410 deletions

63
NX2_GPU_THREAD_SPLIT.md Normal file
View file

@ -0,0 +1,63 @@
# Splitting the GPU thread
Title 1 (Mario Kart World) is bound by one host thread, the GPU thread: ~100% busy while the
host GPU is 20-40% busy and the game's own CPU thread mostly polls the counter. Per second of
3D scenes (profiled, 2026-10-29): command processing and register writes ~335 ms, macros
~200 ms (1.3 million calls), draw setup ~345 ms (55,000 draws), compute ~70 ms. Speeding up the
pieces (NX2_NOTES rows 123-139) took frames from ~22-38 ms to ~18-26 ms; a locked 60 fps on the
heavy courses needs work on more than one core.
## Split point: method calls
- **Front thread** (the current GPU thread): takes the CPU's command queue, parses GPFIFO
entries and command lists (DmaPusher), runs macros (compiled, macro_turing_jit.h), and keeps
what macros read: a mirror of each engine's registers, its macro code and data RAM.
- **Back thread**: executes the resulting stream of method writes in order: Maxwell3D /
KeplerCompute / Fermi2D / DMA engines' CallMethod and everything behind them (register
state, dirty flags, constant buffer and inline uploads, draws, dispatches, queries, fences,
the rasterizer and its caches). Only the back thread touches the rasterizer.
- The stream: an SPSC ring of {engine, method, argument} plus copied data for multi-word
methods (command lists are often guest memory read in place, which the guest reuses).
CPU-side commands (flush/invalidate requests, RunInOrder actions, presentation) are
forwarded into the same ring so they keep their order.
## What the front needs without waiting
- Register values macros read (MME_READ): the front mirrors every register write it sends.
Back-side changes to registers it must reproduce: constant buffer offset (+4 per
const_buffer.buffer word), shadow_scratch[0] = 1 (falcon[4]), shadow RAM control (Track /
Replay change the effective value), macro upload pointers (the front handles load_mme
itself). Compute: native_registers (0xad is changed temporarily by the inline QMD path).
- Macro code uploads and binds (load_mme.*, compute 0x45-0x48): handled by the front.
## Where the front waits for the back (sync points)
- Macro DMA reads (MME_DMA_READ / _FIFOED, ~450 a second): read guest memory as of their
position, so the back must have executed everything before them.
- Semaphore acquires that are not satisfied yet (~10,500 a second, cross-channel waits on QMD
releases): the release is executed by the back, so the front waits until the semaphore
holds (or parks the channel and runs another one, as the pending-channel code does now).
- Syncpoint waits, WaitForIdle-style methods, and anything else that reads guest memory or
host state on the front.
- Expected cost: each wait drains the back only up to the dependency, then the back idles
until the front sends the next method (a few microseconds), so even ~10,000 waits a second
leave most of the overlap. Ceiling: front ~35-40% of the work, back ~60-65% -> ~1.4-1.6x.
## Stages
1. Instrument: count sync points per second and the front/back share precisely (zones).
2. Ring + back thread executing method writes; the front still does everything else.
Correctness first: an EDEN_NX2_SPLIT=0 switch runs both on one thread.
3. Move macros to the front with the register mirrors; verify mirrors against the back's
registers in a debug mode (compare on every MME_READ).
4. Sync points; crash and screenshot tests across courses; then title 4.
## Status (2026-10-29)
Stages 1-3 are in (babd395, `gpu_split.h`), off by default: `EDEN_NX2_SPLIT=1` runs engine methods
on the back, `EDEN_NX2_SPLIT=2` also runs the 3D and compute macros on the front against register
copies. Measured in title 1 (NX2_NOTES rows 142-144): ~44,000 draws a second on one thread, ~47,400
with stage 2, 50,000-66,500 with stage 3. The front waits ~4,000-5,000 times a second (semaphore
acquires, MME DMA reads); the back still spends ~200 ms a second outside its zones (ring dispatch,
method sink). Clean A/B (NX2_NOTES row 145): 1.46x throughput, frames 17% shorter; on by
default since. The back is ~92% busy, ~53% of it draw setup: the next step is there, not the ring.

264
NX2_NOTES.md Normal file
View file

@ -0,0 +1,264 @@
# Eden on Switch 2 titles (branch `nx2`)
Local experiment: get Eden (Switch 1 emulator) to load a Switch 2 title's loose ExeFS and push
boot as far as it goes. Nothing here is upstreamed or pushed. Graphics (Ampere / NVN2) is out of
scope for this pass.
## Setup
- Source + build live in `C:\dev\eden` (a worktree of `eden-nx2/`, branch `nx2`). The original
working folder path is ~180 chars, and `cl.exe` can't open its PDB from a build folder there
(C1041), so the build moved to a short path.
- Toolchain: VS 2026 (MSVC 14.51) via `vcvars64.bat`, Ninja from VS, CMake 4.4.
- Configure:
`cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=RelWithDebInfo -DCMAKE_INTERMEDIATE_DIR_STRATEGY=SHORT -DENABLE_QT=OFF -DYUZU_CMD=ON -DENABLE_WEB_SERVICE=OFF -DYUZU_ROOM=OFF -DYUZU_TESTS=OFF -DENABLE_WERROR=OFF -DENABLE_UPDATE_CHECKER=OFF`
- Build: `cmake --build build --target yuzu-cmd` -> `build/bin/eden-cli.exe`.
- Run (portable mode, `build/bin/user/` holds config/logs; no keys or firmware installed):
`eden-cli.exe -g "<title>\exefs\main"`
- Debug helpers:
- `EDEN_NX2_SVC_TRACE=1` logs every SVC with arguments and results.
- `EDEN_NX2_IPC_TRACE=1` logs every dispatched CMIF command and non-zero result words.
- `EDEN_NX2_AUTO_STUB=1` answers every unknown command with an empty success (the
`use_auto_stub` setting gets reset by eden-cli's config rewrite). Used for all runs below.
- `eden-cli --filter "<log filter>"` sets the log filter for one run (e.g. `*:Info Loader:Debug`
to get module load addresses).
- `tools/nx2/nso_syms.py <nso> addr|got|find ...` maps module offsets to exported symbols, GOT
slots to imports, or searches symbol names. Eden's backtrace shows "unknown" modules for some
titles; with the `loaded module X @ addr` lines this turns frames into SDK names.
- `tools/nx2/a64_coverage.py <a64.inc> <nso>...` lists encodings in an NSO's .text that
dynarmic's decoder table doesn't match.
- `tools/nx2/npdm_info.py <main.npdm>` dumps the META/ACI0 header, services and kernel caps.
- `tools/nx2/nso_dis.py <llvm bin> <nso> <text_off> <count> <workdir>` disassembles part of an
NSO's .text (clang + llvm-objdump from VS 2022's LLVM).
- Under cdb, stop only on second-chance AVs (`sxd -c2 "..." av`): dynarmic fastmem uses
first-chance access violations.
## Titles
All titles share the same `main.npdm` layout: 64-bit, 39-bit address space, ThreadInfo cores
0..5 (core mask 0x3f), priorities 28..59, KernelVersion cap 9.1, SVCs 0x01-0x36 plus 0x80,
Switch 2 ACID layout.
All four are **Switch 2 builds**: none is a plain Switch 1 sanity check. Instruction counts that
stock dynarmic did not decode (`a64_coverage.py` against the unpatched `a64.inc`):
| Title (ExeFS) | Program id | RETAA/B | LDAPR | LSE atomics | CAS | SWP | other |
|---|---|---|---|---|---|---|---|
| Mario Kart World `exefs` | 0400C3F00006E000 | 114,767 | 37,560 | 24,718 | 6,640 | 2,470 | CASP 6, XPAC 7, PAC 2, STG 1 |
| Pokopia `exefs` and `updates/exefs` | 0400A3A01EE4C000 | 12 (rtld only) | - | - | - | - | main/sdk/subsdk0 don't decompress (broken dump) |
| Splatoon Raiders `exefs` (update `updates/65536` ~same) | 0400380020056000 | 134,911 | 100,520 | 32,110 | 8,667 | 3,733 | PACGA 58, PAC 7, CASP 6, XPAC 3, STG 1 |
With the patched decoder, games 1/3/4 only leave ~60-110 rare SIMD encodings in `sdk` undecoded
(e.g. `5f119f9c`, `4ea05c42`, `0f413044`, FP16/v8.2 forms).
## Blockers hit, in order
| # | Where it stopped | Cause | Fix (commit) |
|---|---|---|---|
| 1 | CMake try-compile | C1041, PDB path ~290 chars | build in `C:\dev\eden` |
| 2 | Compile | MSVC 14.51: `reinterpret_cast<u64>(ULONG_PTR)`, C4505 on amalgamated zstd.c, C4800 int->bool | Fix build with MSVC 14.51 |
| 3 | Decoder | dynarmic had no LSE atomics, CAS/CASP, LDAPR, RETAA/RETAB, BRA*/BLRA*, LDRA*, PAC*/AUT*/XPAC*, PACGA, MTE ops. Title 1 has ~115k RETAB, ~37k LDAPR, ~25k LSE, ~6.6k CAS, ~2.5k SWP | dynarmic: decode ARMv8.1-8.5 instructions (atomics are exclusive load/op/store with retry; PAC is identity; MTE tags are 0) |
| 4 | Loader | `main` identified as a bare NSO, so main.npdm/rtld/sdk were skipped, then a crash writing into an uninitialised process | loader: load a loose ExeFS when given its main |
| 5 | NPDM parse (ErrorBadFileAccessControl) | Switch 2 ACID has a 0x300-byte signature/key area (Switch 1: 0x200); FAC/FAH offsets were also read without the block base | npdm: parse the Switch 2 ACID layout |
| 6 | eden-cli crash on start/exit | `*appstate` never set, SDL3 callbacks returned SUCCESS (quit), InputSubsystem was a local of SDL_AppInit, mouse timer before mouse driver | eden-cli: fix SDL3 main-callback lifetime bugs |
| 7 | `nnosInitialize` abort | `svcGetInfo` 29/30 unknown (region between alias and ASLR) | kernel: stub ids 29/30 (empty region after alias) |
| 8 | `VammManager::InitializeIfEnabled` abort | `svcGetInfo` 33 unknown; it's a page granule (module ranges are checked against `value-1`; VAMM maps one granule first) | return 0x1000 |
| 9 | `nn::am::InitializeApplicationInterfaces` | `appletOE` cmd 10 (Switch 2 OpenApplicationProxy) | alias to cmd 0 |
| 10 | same | `IApplicationProxy` cmds 60 and 40: new getters | return a generic `nx2:` stub interface (`nx2_stub_interface.h`); calls on it are auto-stubbed and logged |
| 11 | `nn::oe::SetupGpuErrorHandler` -> CreateThread | thread on virtual core 5; Eden emulates 4 cores | fold virtual cores 4/5 onto 0/1 in CreateThread/SetThreadCoreMask |
| 12 | (title 3) `nn::oe` init | `IApplicationFunctions` 320 has no handler | empty-success stub |
| 13 | (titles 3/4) `nn::fs::QueryMountRomCacheSize` -> `GetFileSystemProxyServiceObject` abort 2002-1002 | fsp-srv `SetCurrentProcess` returned ResultTargetNotFound: the loose-ExeFS loader never registered the process with the fs controller, and never set its RomFS | loader: give loose ExeFS boots a RomFS (sibling `romfs/`) and an fs registration |
| 14 | (title 4) `IApplicationFunctions` 321 | unknown | empty-success stub |
| 15 | (title 4) `nn::oe::GetCurrentGraphicsScale` UnexpectedDefault | `IApplicationProxy` 60 is a graphics-scale interface; its cmd 10 must return 10/20/30 | `IGraphicsScaleController` returning 10 |
| 16 | (title 4) `nn::gs::sdk::BootGraphicsSystem` -> `nn::os::GetRemoteMemoryUsage` abort | `svcGetInfo` 31/32 (remote memory for gsv/nvsrv) | return 0 for both (logged) |
| 17 | (title 4) `nn::gsv::BootWithRemoteMemory` abort | it sends TIPC on a session handle from `ISelfController` 2000, which had no handler | `ISelfController` 2000 returns a raw session to `INx2GraphicsSystemSession` (`nx2_gsv.h`), which accepts every command |
| 18 | (title 3) `nn::oe::GetPseudoDeviceId` abort | needs control.nacp, which loose ExeFS lacks | seed from program id |
| 19 | (title 4) `nn::syncpt::sdk::InitializeEx` | gsv 202 must return the syncpt session; gsv 204 takes two transfer memories and must leave a header (version 1, count, offsets) in the read-only one and return two events; syncpt 1 returns a u32 < 0x400 | `INx2SyncptSession`, gsv 202/204 |
| 20 | (title 4) `nn::gpu::sdk::BootWithRemoteMemory` | gsv 497 returns the GPU session; gpu 11 (id range), 12(id) event, 21(id) start + wait on that event, 22 another session | `INx2GpuSession` |
| 21 | (title 4) `nn::vi2::sdk::Initialize` / `nn::vi::CreateLayer` | gsv 305 returns the vi2 session; vi2 1/3/5 return small u32s (5 = timestamp recorder count, 7 needed by LayerHistory), 11/31/32 return handles | `INx2Vi2Session` |
| 22 | all TIPC replies with data | Eden bug: `ResponseBuilder::Push(Result)` writes the 2-word CMIF result, so TIPC data lands one word late and past the copied size | `TipcReplyWords` |
| 23 | (title 4) NVN -> NvRm device open | nvsrv is the nvdrv INvDrvServices protocol; serve it with Eden's nvdrv, initialised without `Initialize` | nvdrv-backed `nx2:nvsrv` |
| 24 | NvRm chip check | Switch 2 NvRm asks for a 0xD0-byte characteristics block and accepts chips 0x17E/0x12B/0x12E, but NVN needs the Ampere one; s16 interface versions at 0x70..0x80 must be >= 6 etc. | report chip 0x17E, Ampere classes, versions 0x40; unknown ctrl-gpu ioctls zero-stubbed for such callers |
| 25 | NvRm address space | AllocAsEx(flags 5, no big pages, [0x8000000, 0xC000000000), split 0); NVN then reserves 4 KiB VA up to 512 GiB; Eden's split/limit handling made the reported regions unusable | small pages up to the end, absolute big-page limit |
| 26 | (title 4) NvRm channel setup | nvhost-gpu ioctls 0x1F/0x20, `/dev/nvhost-tsg-gpu`, GetEngineInfo, object class 0 | logged echo for unknown ioctls once an nx2 driver is seen; GR + async-copy engines; Ampere/class-0 objects on Maxwell slots (70f1651, 82dbd88) |
| 27 | GPU memory | svc 0x80 (make memory GPU-accessible), gpu sub-session 101 (register range) and 211 (job) | stubs (608f0d6) |
| 28 | `syncpt::sdk::InitializeEx` -> `InterProcessEvent::Signal` abort E401 | gsv 204's first handle is attached *writable* (client -> server request), the second readable; the shared page has u64 values (RO+0xc), armed bits (RW+0x18, RO+0x1c) and summary words | gsv 204 returns a writable event, lays out the page, a 500 us poller mirrors host1x values and acknowledges armed waits (88367ea) |
| 29 | `NvOsEventPoolWait` E401 | syncpt 22(slot) must return an event; 201(slot, ~syncpoint, value, mode) / 202(slot) request/cancel a signal | slot events via host1x guest actions (88367ea) |
| 30 | GPU never executed submissions | Switch 2 never maps through nvhost-as-gpu: memory registered with gpu sub 101 is used at GPU VA == CPU VA | `Nx2GpuMemoryRegistry` identity-maps every range into each GMMU; SMMU physical range 33 -> 34 bits for 10/12 GB layouts (88367ea) |
| 31 | Puller | binds Ampere classes C797/C7C0/C7B5 and class 0; NVN2 uploads Turing+ (MME64) macros, which Eden's Fermi macro JIT runs as garbage until it crashes | Ampere classes run on the Maxwell engines; macros skipped while `Tegra::g_nx2_ampere` (88367ea) |
| 32 | NVN device init waits on syncpoint 1 >= 1 | submits use `increment_value` (bit 8) and signal completion only with semaphore releases | Eden appends the syncpoint increments for such submits (88367ea) |
| 33 | `wsi::ounce::NativeWindow::Acquire` on a null window | gsv 1510 is `gsv::shim::CalculateRequirements`; size 0 made window init fail | block-linear size/alignment, 4 bytes/pixel (5fb37b5) |
| 34 | `NvOsEventPoolWait` E401 (second channel) | nvhost-gpu QueryEvent 7 | per-channel event (5fb37b5) |
| 35 | CPU spins in NVN for `[0x2106028010] >= 3` | queues release progress to page P+0x1000 and wait (GPU acquire + CPU poll) on the same offset in page P | puller mirrors releases into the first two-page gpu region (5fb37b5; guess, works) |
| 36 | font header check on a wild pointer | `pl:u` 2010 `GetBuiltInFontData` (bool at offset 0, then {flags, u32 offset, u32 size}) and 2000 (font shared memory + size) | plain TTF copies of Eden's shared fonts after the BFTTF data (5fb37b5) |
| 37 | null "SoundResource" heap (300 MB) | `nninitStartup` sizes the heap as total - used - 0x195000000; Eden capped total at the 8 GiB heap region, leaving 1.55 GiB | Switch 2 processes count the alias region in max memory (5fb37b5) |
| 38 | game waits for service `aud:u` | Switch 2 audio entry (`nn::audio::CreateUserClientSessionByHipc`); cmd 2 returns an IAudioRendererManager | `INx2AudioUserClient` backed by Eden's audren (2e6c39c) |
| 39 | every audio `RequestUpdate` fails | updates tagged `REv?` (lowercase v), renderer opened with `REV?` | compare revision numbers (2e6c39c) |
| 40 | main thread parked in `ExecuteVirtualMachineMapBufferEx` | gpu sub 211's job slot is request word 17, not word 2 | signal word 17's event (2e6c39c) |
| 41 | "too many files open" reading RomFS | eden-cli never raised the CRT stdio limit | `_setmaxstdio(8192)` (2e6c39c) |
| 42 | GPU memory model | gpu sub 101 = `CreateGpuAccessibleMemory` (reply word 3 = GpuMemoryHandle); 211 = `NormalMapGpuMemory4K` (words 3-4 first VA, 5-6 last, 7 handle, 8-9 offset, 17 slot), also into sparse regions at 0x81_xxxx_xxxx; 102/213 (`NormalUnmapGpuMemory4K`)/214 (`CreateSparseRegion`) complete through slot events (words 3/8/11) | real handles, `Nx2GpuMemoryRegistry::MapAt` with SMMU range reuse, slot signals |
| 43 | (title 1) `nn::syncpt::sdk::Initialize` abort | newer SDK opens syncpt with TIPC 0(u64 token) and aborts unless it is echoed | handshake echo |
| 44 | (title 1) NVN closes the GPU device after GetVARegions | newer SDK asks gpu 201/202(id) for the user/gsv partitions (two u64s each, stored as min/max) instead of using 203 | 201 = [0x8000000, 0xC000000000), 202 empty |
| 45 | NVN command buffer with no command memory (~12 s, titles 1 and 4) | `NvRmGpuTSGSubcontextCreate` failed (0x11): characteristics reported one VEID | flags bit 8/22 (TSG, TSG_SUBCONTEXTS) and 64 subcontexts at characteristics +0xb4 (12125bc) |
| 46 | NVN asserts in `GetSyncpointShimRoAddress` | Switch 2 reads syncpoints through a read-only GPU page (flag bit 26, nvhost-as ioctl 0xB) | shim at the top of the VA range; GPU reads of it return the syncpoint value (12125bc, from changes.patch) |
| 47 | window/image creation | gsv 111/112/114/423/429/430/435/1510 native memory and images; svc 0x80 is SetGpuMemoryAttribute | native object registry + layout calc (from changes.patch), real svc 0x80 (12125bc) |
| 48 | presents go nowhere | vi2 229 (word 5 image id, ~word 22/word 24 fence) | composite the block-linear image with native GOB tiling (f63ff90, from changes.patch) |
| 49 | crash in `DmaPusher::CallMethod` (null engine) | NVN uses subchannel 2 (inline-to-memory) without binding a class | bind KeplerMemory there by default; drop methods on unbound subchannels (9f30683) |
| 50 | GPU parses garbage, then stalls (~14-19 s) | NVN2 gates each command buffer with a SEMAPHORED acquire the CPU releases after writing it; Eden read past unmet acquires and ran stale commands (`EDEN_NX2_PB_CHECK` confirmed entries changed after submit) | suspend the channel on an unmet acquire and resume it from the same word (9f30683, approach from changes.patch) |
| 51 | desync after a called command buffer | NVN2 ends a GPFIFO entry inside a method's data, calls another buffer, and continues at the next address | set the unfinished method aside and resume it when the stream continues (9f30683) |
| 52 | semaphore words next to releases get clobbered | 4-byte releases (SEMAPHORED bit 24) were written as 16 bytes | 4-byte releases write 4 bytes (9f30683) |
| 53 | no draws: 3D macros skipped | NVN2 drives 3D/compute state and draws through Turing MME64 macros; compute dispatches through QMD 3.0 | changes.patch's MME interpreter (one macro context shared by the process's channels, also on the compute class), native compute and QMD scheduler (d7826ce) |
| 54 | draws hit Eden's Maxwell decoder ("Invalid insn") | shaders are SM86 SASS | changes.patch's Ampere-to-Maxwell lowering, widened until all 32 of title 4's programs lower; `EDEN_NX2_SHADER_DUMP` + `C:\dev\nx2work\lowertest` test it offline (d7826ce) |
| 55 | GPU desyncs again after ~15-30 s | the game ran a full ring of command memory ahead of Eden's GPU thread: the vi2 229 reply reports all buffers free, so nothing paced it | vi2 229 waits (bounded) for the frame's GPU fence (325cac7) |
| 56 | macro arguments lost | NVN interleaves several command buffers, so more than one split method can be pending | several parked methods keyed by resume address, dropped at the end of the submission (325cac7) |
| 57 | frames all black although draws run | gpu 214 creates GPU-only memory (six 300 MB regions at 0x81_0000_0000) that the graphics server backs itself; the render targets live there and were unmapped | back each region with application-pool pages mapped into the alias code region (325cac7). Title 4 now shows its loading screen and a menu screen |
| 58 | (title 4) hang after confirming difficulty: macro 0x142 "parameter FIFO underflow" | NVN2's conditional-render macro reads two query counters with MME_DMA_READ_FIFOED (3D methods 0x154/0x155 address, 0x158 count) and loads them after a load barrier | MME DMA reads into the load FIFO / data RAM (8cbe854) |
| 59 | equal-value gate semaphore missed (~24 s) | the CPU side ran ~96 command buffers ahead when the 250 ms present wait timed out during shader compiles | present waits up to 4 s for the frame fence (8cbe854) |
| 60 | 3D queue waits forever on a word the compute queue should release | NVN2 patches a third QMD release slot at word 52 for queue-to-queue syncs | fire release slots 24, 28 and 52 (8cbe854) |
| 61 | 3D queue waits for N compute completions, 4 missing | compute jobs whose shader Eden rejected never completed, so their +1 releases never fired | rejected QMDs still complete and release (a84ec80); IMAD -UR addend, LDSM lowered |
| 62 | queue waits for ≥ a value that keeps moving backwards | QMD release slot 0 (flag bit 19) is a max release into a word a 3D macro also advances | bit-19 QMD releases write max(old, payload) (a84ec80) |
| 63 | in-game draws refused (~20k per run), black screen behind the dialogue | strong LDG (bit 77), 2D TLD4, RED MAX, AST of RZ unsupported | lowered (a84ec80); 210/214 post-menu shaders lower |
| 64 | intro scene black (DLSS-style pass) | title 4 renders at 2560x1440 and upscales to 3840x2160 with a tensor-core (HMMA/MOVM/LDSM) compute network; emulating it takes the host driver minutes to compile and would cost seconds per frame | bilinear blit from its input (cbuf0+0x90) to its output (cbuf0+0xa0) when the program is rejected (22333bd); full emulation opt-in (EDEN_NX2_TENSOR) |
| 65 | fullscreen blit passes read constant bank 24 | console-provided blit parameter block, never bound by the game | bank 24 -> 17, synthesized from the render area by BlitConstantsPass in vertex and fragment shaders (22333bd, d75e57d) |
| 66 | synthesized blit sizes were 0 | the game never programs surface_clip, which feeds Eden's render area | fall back to colour target 0's size (d75e57d) |
| 67 | tonemap output black | the 3D colour-grading LUT is drawn once per scene by a pixel shader using CALL.REL.NOINC; capture treated it as a real call, never found the end and rejected the shader | NOINC calls count as forward branches in capture (4d58a74). **Intro cutscene renders correctly** |
| 68 | crash ~5 s after the save that follows character customization (garbage pointers in two threads) | nx2:gpu 213 (unmap) and 102 (destroy GPU memory) were no-ops; the game reuses the freed memory as heap and stale GPU write-backs overwrote it | unmap from every address space, drop freed semaphore regions (bed817c); not yet confirmed past customization: frame-locked replays desync across loading screens |
| 69 | hatched edges on UI portraits, buttons, squid cursor | UI textures are ASTC; the GPU ASTC decoder ignored the Switch 2 GOB swizzle | apply it in astc_decoder.comp (7fa2b46) |
| 70 | freeze ~15 s after the post-customization save; channels 2/4 wait on 0x2115c98000 (equal 0x5293, now 0x5310) | a 3D queue ticks a counter two queues acquire value by value; Eden wrote several ticks between their re-checks | releases latch the waiters they meet (8e37f58) |
| 71 | audio broken, game at 90-160 fps | vi2 presents were never paced | 60 Hz vblank pacing |
| 72 | stuck loading the next scene ("too many files open" for route_wallClimb.pack.zs, save.dat) | VFS open-file cache cap equal to the CRT limit (8192) | cap 512 below _getmaxstdio() |
| 73 | loading screen hangs, host GPU TDR | macro parameters split across GP entries were parked; method data continues across entries | continuous split methods (e8d060b) |
| 74 | rejected ATOMG/STL/LDL with small immediate offsets | lowering only took a zero offset | s20 ATOM offset, 24-bit STL/LDL offset (6302be8) |
| 75 | same crash 3 runs (415 s, 429 s, 9960 s): read 0x29 at pc 0x85096d10 (recursive mutex lock), lr 0x82ada310 | worker job runs main+0x2aba2b0 on object C = [A+0x18] whose manager [C+0x720] is null; C's destructor (main+0x2ab8e20, vtable main+0x4b660f8) nulls it before A's teardown (main+0x2b37b50) clears A+0x18; teardown and per-frame update are both jobs (via 0x8365c634); NFC Finalize comes 14-22 s before each crash; LSE atomics (exclusive pairs) and DMB (mfence) checked fine | open: catch all threads at the abort with the gdbstub (`nx2work/dbg_crash_threads.rsp`) |
| 76 | debugger run, crash at 406 s in the same scene load: writes through a null result at pc 0x82b471fc (main+0x2b271fc) | the scene's frame heap "PaintSceneProcHeap.26" (fixed 90 MiB + 0xfc8 parent, head 0x12c9525d18, tail 0x12c9800208) had 2.98 MB left for a 0x302088 allocation (A init main+0x2b37de0 -> main+0x2b27020); about 160 KB short, so some allocation of that scene is bigger than on the console (only 11.7 MB of the heap is nx2:gpu memory); the main thread was waiting in svc 0x34 for job workers | open: log that heap's allocations to find the oversized one |
| 77 | why the scene heap overflowed (run 6) | the heap size is a constant (mov w1, #0x5a00000 at main+0x2b32070); its fixed contents are 71.3 MiB (main+0x2ab9218) + 1.05 MiB + 1.25 MiB + 3 MiB; in runs 3, 4, 5 and the save test the scene's two 0x514000 GPU pools come from the root heap (next to the 0x2ee000 pool), in run 6 from the scene heap (current heap of the loading thread), 10 MiB it has no room for; which heap a job allocates from depends on the thread running it. The game's npdm asks for cores 0..5 (mask 0x3f, reported as is by GetInfo CoreMask) and Eden folds 4, 5 onto 0, 1, so two cores run two workers each | candidate: report or emulate the real core layout; ruled out: memory layout (12 GB), controllers (one), GPU pool alignment |
| 78 | six emulated cores (4fbfd5b), run 7 with the gdbstub: same scene-heap overflow at 292 s, the 0x514000 pools again in the scene heap | both overflow runs (6, 7) had the gdbstub on (memory accesses checked, slower JIT); the three runs without it (3, 4, 5) put those pools in the root heap and crashed later on the torn-down object instead, so the debugger itself shifts this timing | open: a six-core run without the gdbstub |
| 79 | run 8 (six cores, no gdbstub): original crash at 338.6 s, 11 s after the paint scene set up (327 s, pools in the root heap) | C's init (main+0x2ab9120) creates "ColPaintBuild.%u" (main+0x147d3c0 -> main+0x34f3850) sized by [C+0x738], 100 MiB by default (ctor main+0x2ab874c); when that returns null, init returns before setting [C+0x720], and the next state request (main+0x2aba230) crashes the update. So every crash at this load is a memory shortage: 100 MiB (no gdbstub) or 3 MiB in the scene heap (gdbstub). Root heap is 0x19506f000 at 0x1145e10020 | open: why memory is short (reported total memory for Switch 2, frees delayed behind the GPU) |
| 80 | the memory shortage behind rows 75-79 | title 4 is told 9.55 GiB (GetInfo 21) and plans all of it: SetHeapSize 3.125 GiB plus a 6.33 GiB VAMM-backed sead root heap in the alias region; Eden backed the six 300 MB gpu 214 regions from the application pool, 1.8 GB the kernel still reported as free | 214 backing from the non-secure system pool, +8 GiB DRAM on the 12 GB layout, device memory tracker 35 physical bits (5cff620); run 9 passed the crash spot (NFC Finalize 272.8 s, ran to 388 s), no allocation failures |
| 81 | title 1 hangs before its first frame (~17 s): main logic stops, two cores spin, only audio and hid keep running, almost no GPU submits | it asks gpu 214 for ~20 GiB (two 0x16ad80000 regions, these two still fail with the new pool); gets further than before (save data, friend list) | open: thread snapshot at the hang; 214 regions this large must be sparse on the console |
| 82 | title 1 hang (row 81) | `monitor handles` (a4051e4) showed the main thread waiting forever on the nx2:vi2 32 event (loop at main+0x25ec5e4 until [obj+0x88] >= 2); the stub never signalled it | 31 and 32 events signalled at 60 Hz (6196a11); title 1 now presents (241 frames by 95 s) |
| 83 | title 1 after the hang: ~3 fps with "fence not reached in 4 s" stalls | unimplemented vi2 104/105 (every frame) and gpu 215, gpu 214 still short for the two 5.7 GiB regions, one Nx2Map "out of SMMU space", rejected shaders (formatted surface store, IPA mode, LEA carry) | open |
| 84 | title 1 GPU stalls (row 83): channels wait on semaphores 0x21f6093140/3040 until the 4 s timeout | NVN2 queue syncs run through a ring of QMDs (dependency 2 each: the previous entry's launch and macro 0xda, which patches release slot 52 and schedules/decrements via PCAS 3/7); macro 0xda reads the ring base from MME data RAM word 0x27, which each channel sets (0x21098d39, 0x213e8ff8, 0x213ea168) but Eden shared one data RAM for all channels, so every channel drove the same ring | per-channel MME data RAM, one QMD table for all channels (6c20e37); no stalls, title 1 reaches its title screen (world map with logo) |
| 85 | title 1 address space | 214 = CreateSparseRegion, 215 = DestroySparseRegion (SDK symbols); SMMU space (16 GiB) ran out | 64 GiB device space (71b38ab), 215 handled (cc0c400) |
| 86 | title 1 after its title screen: black frames | tens of thousands of rejected shaders: IPA interpolation modes, integer carry/third source, uniform comparison, formatted surface load/store, uniform constant load, unsupported opcodes | open |
| 87 | title 1 loading crash: Vulkan out of memory | an 8192x8192x68-layer virtual texture (~23 GB) | images over 2^28 texels are created sparse with no memory (`Unbacked`: no uploads/downloads) (96a9b81); the game reaches its real title screen ("Press L + R"), ground black (that texture has no data) |
| 88 | title 1 shader coverage | checked against nvdisasm (`nx2work/nvd.py` disassembles one instruction at a time, since a single illegal word stops nvdisasm); calls: RET may restore the link through IMAD.MOV with code between, library functions may EXIT, `@Pg EXIT Pc`; PLOP3/UPLOP3 any table, UP2UR, ULOP3 predicate output; UIMAD.WIDE, uniform USHF/FMNMX/I2F/F2I, LDS `[Ra.Xn+URb]`, ATOMG/RED/STG forms, PRMT, FSWZADD, REDUX.MAX | 1351c73, 52cda01, 49396b7: 1699 of 1712 dumped programs lower; left: 12 task/mesh programs (ISBERD/ISBEWR, row 89) and one SURED |
| 89 | task/mesh shaders | title 1 draws some geometry with NV-style task + mesh programs: the task program culls a batch and writes its payload with `ISBEWR.O.ATTR [Rn+0x10]` and the mesh count with `ISBEWR [0x4]`; mesh programs read it back with `ISBERD.ATTR.SKEW`. The most-drawn rejected program (0x114c459500) is a task program. Eden has no mesh pipeline | open: needs a task/mesh pipeline (VK_EXT_mesh_shader) or a compute + indirect-draw emulation |
| 90 | each rejected program was lowered and logged again on every draw (1.7M log lines in 13 min) | no negative cache | the shader cache keeps the first exception per address until that memory is written; logs deduplicated (bcb4ff0) |
| 91 | a compute program with a jump table: "Failed to track indirect branch" | Eden's tracker wants IMNMX imm -> SHL imm -> LDC -> BRX; IMAD.SHL was lowered to an XMAD multiply | IMAD by a power of two lowers to SHL (49396b7+) |
| 92 | with rows 88's compute forms, title 1 stays on its loading map (file reads stop ~46 s in) | two compute programs failing in Eden's compiler: a jump table whose targets lie past the first EXIT, and a 128-bit store of RZ from unaligned scratch registers | programs with BRX end at the compiler's closing branch-to-self; R248-R251 for 128-bit zero stores (3df6a81) |
| 93 | title 1 colour cast (magenta/cyan, red/green, changing per course) | the sky/atmosphere 32^3 volume shader starts with an unguarded R2UR; R2UR/MATCH/HFMA2.RELU saved P6 with the uniform-guard mask `1 << (6 - 7)` (undefined), which shifted into the opcode bits | always save/restore P6 (2ed0c48). Found with `EDEN_NX2_DUMP` (flush GPU ranges at a present to files) and `nx2work/vol3d.py` / `volslices.py` (untile Switch 2 3D textures) |
| 94 | near terrain black | 8192x8192x68 ASTC 8x8 virtual textures: 1.5 GB in guest memory, 23 GB each as Eden's RGBA8 decode | huge GPU-decoded ASTC arrays keep levels from the first that fits 512 MiB (here 3) on the host (60c4791) |
| 95 | wind/foliage shards | SM70+ MUFU.SIN/COS take turns (sin(2pi x)); Eden's Maxwell MUFU takes radians | scale by 2pi (245e6ad); jungle canopy not re-checked yet |
| 96 | title 1 aborts after ~1.5-4 min at the title or when a race starts (nn::diag::Abort from main+0x19e5df0) | the movie player (main+0x19e5860) lays out its decoder frames with nn::codec::CalculateDecoderImagePoolBuffer, which sends nx2:gsv 991; auto-stubbed, the reply kept the request words, so the frame size read as 0x167027f'00000001 and the budget check at main+0x19e5b58 failed. (Not the "video-frame-count" < 9 check: the game sets that from a two-slot registry, at most 2.) | gsv 991 answers a block-linear NV12 size (640x360 -> 0x64000) |
| 97 | pop-in, objects flicker in and out, low-quality meshes | high-detail LODs drawn through task/mesh programs (bound in stage slot 1, header type VS, ISBERD/ISBEWR), which are rejected and skipped | open: task/mesh support |
| 98 | ~12 fps at the title | host GPU ~20% busy, emulated cores ~0.9 host cores, GPU thread 45-70% busy, much of it waiting on the host GPU: semaphore acquires wait for host fences under GPU accuracy High; ASTC uploads end in Finish() (a Flush() there crashed at boot); a 1 ms poll sleep while channels are suspended | gpu_accuracy=Low: ~16 fps; SDL main loop capped at 240 Hz (88f04ad). Open: ASTC upload without a full stall (and per-layer re-upload of the trimmed virtual textures), native MME interpreter cost, draw path |
| 99 | VRAM over the card (8.4-9.8 GB used on an 8 GB RTX 4070 Laptop), texture cache always collecting | ~700 ASTC 8x8 textures decoded to RGBA8 (16x their size, ~4.5 GB) | astc_recompression=Bc3 + accelerate_astc=CpuAsynchronous in the config; huge ASTC arrays stay on the GPU decoder as trimmed RGBA8 (ed0a3c9): 4.6-5.6 GB, ~10 -> 16 fps |
| 100 | 8 ImageTranscode threads at ~44% each (3.5 cores) | Eden's default pool is half the hardware threads | an eighth, at least 2 (`EDEN_NX2_TRANSCODE_THREADS` overrides): ~19 fps on the attract loop |
| 101 | Medal's capture hook inside Eden (`medal_hook64!OBS_Negotiate` in GPU-thread stacks) | Medal registers a per-user Vulkan implicit layer that loads into every Vulkan app | launches set `DISABLE_VULKAN_MEDAL_OBS_CAPTURE=1` (the layer's own off switch) |
| 102 | host crash (log stops, segfault) at 113 s in one run; under cdb "Free Heap block modified after it was freed" at 243 s, twice | Eden's TickAsyncDecode erased a finished decode while still holding its mutex (unlocked freed memory); only reached with accelerate_astc=CpuAsynchronous (the BC3 setup) | release the lock first, cancel decodes of deleted images (8766a72); clean past 330 s. Also made nx2 service callbacks (vsync timer, syncpt actions, capture/dump) independent of their sessions' lifetime (190747a) |
| 103 | (title 1) texture churn: the collector evicted ~1,250 decoded ASTC textures per 600 frames and they were decoded again (lag and pop-in turning the camera) | device memory sat above the collector's critical line (5.4 GB of a 7.0 GB budget): ASTC 8x8 is 4x larger as BC3, and a bloom chain at unmapped GPU 0xf0000-0x15b600 was re-created on every use (each lookup got a fresh stand-in CPU address), ~5,400 images per 600 frames on the map | unmapped GPU addresses get a stand-in address fixed by the GPU address, so images are found again; ASTC whose blocks carry no alpha (exact endpoint-mode scan, ~70% of them) goes to BC1 (`EDEN_NX2_ASTC_BC1=0` turns it off). Attract loop: 4.8-5.1 GB, no costly evictions, ~17 -> 22-23 fps |
| 104 | garbage textures (title logo, road patches) and a host crash 90-120 s in that skipped the crash handler (fail-fast) | an image re-decoded while a decode was in flight has two; the first to finish cleared IsDecoding, so deleting the image did not cancel the second, which was then uploaded into the next image to take the slot (another texture's data; a staging overrun when that image was smaller) | every pending decode of a deleted image is cancelled; a decode larger than its image is dropped and logged. Frames clean; crash: verifying with long runs |
| 105 | (title 1) movie abort moved on: `nn::codec::InitializeVp9Decoder` failed with 2001-0114 (InvalidHandle) | gsv 901/902/1001/1002 (`nn::codec::video::detail::ipc::Initialize`) were auto-stubbed; 902/1002 must return sessions (handle word 3, result word 4) whose TIPC 1 (CreateVideoDecoderService) returns another | sessions returned (2970100); the movie plays with no picture. Open: decode VP9 (vdecsvc 200/101/3/311 are auto-stubbed) |
| 106 | GPU thread ~50% in the Turing macro interpreter | every instruction re-decoded on every run, register file copied per instruction | decode once per PC, operands read up front (a97d26d): GPU thread ~68% -> ~45-50% busy |
| 107 | streamed textures shown as garbage/black for seconds; 2 transcode threads pegged | ASTC was decoded and re-compressed to BC on the CPU, ~40 textures/s on the attract loop | accelerate_astc=Gpu (config value 1) + BC recompression: GPU ASTC decode into a temporary RGBA8 image and `bc_encoder.comp` (69cc66b); textures ready the frame they are used |
| 108 | frame time is CPU work + GPU-thread work in series | nx2:vi2 229 blocks the game until Eden's GPU thread has processed the frame (WaitForFrame), and RequestComposite waits for the GPU thread to drain its queue too; the present timeline log shows 8-56 ms of each 34-105 ms frame blocked | `EDEN_NX2_FRAMES_AHEAD=1` (wait for the previous frame) deadlocks at the first scene change: the main thread waits on syncpoint 1 (present fences) for a value no queued work reaches. Open |
| 109 | ~1,400 SVCs per frame | the job system's WaitForAddress (0x34) / SignalToAddress (0x35) pairs and yields, all through Eden's global scheduler lock (3 of 6 cores seen blocked on it) | idle-core spinning (`EDEN_NX2_IDLE_SPIN_US`) gave no gain. Open |
| 110 | (title 1) GPU thread saturated, 3D attract scenes ~22 fps | profiled with the new GPU zone timers (`EDEN_NX2_GPU_PROFILE`) and the in-process sampler (`EDEN_NX2_SAMPLE`): per second, ~355 ms command parsing outside macros, ~210 ms in 745k macro calls, ~325 ms draw setup (33k draws), ~89 ms compute (4.4k dispatches) | see 111-115; now ~27-33 ms per frame in the same scenes |
| 111 | the guest clock (CNTPCT) cost ~10% of the main core | AMD CPUs have no CPUID leaf 0x15, so the TSC frequency stayed 0 and every read went through QueryPerformanceCounter | measure the TSC when the leaf is missing (948a391) |
| 112 | C++ exceptions on the GPU thread (~1-2%) | a rejected native shader (mesh/task programs) was rethrown twice every time it was bound | return false for remembered rejections (948a391) |
| 113 | debug aids on the hot path (~4% of the GPU thread) | semaphore watch per command word, FNV hash of every pushbuffer entry, first-write trace, two debug logs per macro call | flag checks unless their variable is set (948a391) |
| 114 | texture churn: 1,400-2,600 textures deleted and recreated per 600 frames, one purge of 4,300 | ~38% of device memory was empty space in 256 MiB VMA blocks; usage came from the driver, so evictions did not lower it | 64 MiB blocks and usage minus VMA block slack (7b0ff0b): 0-224 per 600 frames, 6.24 of 6.97 GB used |
| 115 | main core: ~12% slow-path reads of GPU-tracked pages, GPU thread and CPU fighting over the caches' mutex | `use_reactive_flushing` (on by default on desktop) protects tracked pages against reads too | off in the test config: the slow reads are gone, no visual differences seen in the attract loop. Kernel lock spins before sleeping (5c99031): ~20-25% less time blocked |
| 116 | macro interpreter is ~11-13% of the GPU thread | operand fetch is not the cost: a slot-based rewrite matched the old interpreter on 200,000 recorded calls (`EDEN_NX2_MME_RECORD`, `C:\dev\nx2work\mme_test`) but was no faster | open: needs a JIT or pre-built handlers |
| 117 | macro calls: ~1,000,000 a second, 7 instructions each, 96% exact repeats | the per-call interpreter cost was the time, not the instructions | memo of calls that depend only on their parameters and the data RAM words they read (1cd4a3c, b75fdf8, c28ea9a for compute macros): 97% replayed, verified on 200,000 recorded calls |
| 118 | crashes ~26 s in (fencing thread walking a corrupt operation list; query cache Barriers; Vulkan worker SubmitExecution) | `async_presentation=true` (config): 4 crashes in 8 runs with it, 0 in 6 without; a race between the present thread and the Vulkan scheduler | keep it off (test build 3b). The fence manager's SyncOperation also appended to a list SignalFence moved from another thread: own lock now |
| 119 | (tester, RTX 4070 SUPER) 50 MB log; textures evicted with 3 GB VRAM unused | nx2:vi2 104/105 called ~75 times a second through the logging auto stub; collector thresholds from Eden's ~7 GB device memory cap | quiet stubs (649fb56); thresholds from the driver budget (49225dd) |
| 120 | GPU thread asleep ~8% on the texture/buffer cache mutexes | holders: the fencing thread (it took both caches for every fence, one per compute dispatch, and holds the buffer cache while it copies GPU-written buffers back) and CPU write traps; std::recursive_mutex slept at once | spinning recursive mutex, per-cache fence flush flags, separate locks (f10767f): asleep 3.7%. Buffer downloads at every fence stay: the native MME DMA reads of indirect arguments depend on them (no synchronous flush logged with them on) |
| 121 | first-run freezes from shader translation (one 28,102-instruction compute shader: 3.4 s on the GPU thread) | BreadthFirstSearch's visited list was a linear scan (quadratic), and GlobalMemoryToStorageBufferPass searched again for every access | hash set past 64, per-start memo (ba173b6): 0.85 s; runtime builds now ~2-4 ms each. Translating on workers was not done: environments read live GPU state (texture types, constant buffer values) while translating |
| 122 | final clean run (attract loop): 21-25 ms per frame in most 3D stretches, 26-38 ms in compute-heavy courses (industrial, highway), host GPU ~38% busy | the GPU thread is still the limit: draw setup ~25%, command parsing and register writes ~15%, compute dispatches/inline uploads (heavy courses) | open: per-draw binding cache, uniform buffer copies (staging slices are recycled by tick, so reuse needs care), buffer download interval sets |
| 123 | (title 1) 5,500-7,000 fences a second, one per compute release, each made in full (both cache locks twice, a fence object, a fencing-thread wake-up) even with nothing to wait for: SignalFence 7.6% of the GPU thread | with fast GPU fences the release callback runs at once; ~20% of the fences had no downloads, queries or operations. OverlapRangeSet::Subtract restarted its search after every erase | empty fences skipped (one in 64 still made), one lock per cache per fence, single-pass Subtract (9905916): SignalFence 3.9% |
| 124 | GPU thread ~9% spinning on, sleeping on and waking for the cache locks | the fencing thread took the buffer cache once per fence to write 250-330 MB/s of downloaded buffer data back (13,000-20,000 ranges a second) | signaled fences released together (31cd7b2, `EDEN_NX2_FENCE_BATCH=0` off): spinning 7.0% -> 5.5%; fence cost per call unchanged within noise (longer holds). Open: the downloads themselves |
| 125 | crash every few minutes: BindHostGraphicsUniformBuffer -> ReadBlockUnsafe, access violation at 0x600 (2 of 4 runs) | DeviceMemoryManager's translation cache was shared by all threads; Map on a service thread (nx2:gpu 211 job mappings, logged at both crashes) cleared it between the GPU thread's null check and its copy | per-thread cache invalidated by a generation counter (1a497fa): 0 crashes in 6 runs |
| 126 | macro parameters appended to a std::vector per call (~2% of the GPU thread) | ~1,000,000 calls a second, nearly all whole in one method run | stack copy (3caf8ea, `EDEN_NX2_MACRO_COPY=1` off). Passing the command list words in place is unsafe: they are often guest memory read in place, and the memo stores the parameters after the call |
| 127 | main core: ~23% of its samples read CNTPCT (~17 million reads a second) | the game polls the counter while it waits for the GPU; a spinning thread slows its SMT sibling (8-core/16-thread Ryzen) | PAUSE between counter reads less than ~1.7 us apart (8abeb0b, `EDEN_NX2_CNTPCT_PAUSE=0` off); effect not separable from run-to-run noise |
| 128 | Vulkan worker woken with a system call every 8 draws (~2% of the GPU thread) | DispatchWork notified even when the worker was busy | every 32 draws, notify only a waiting worker (d9d7e94) |
| 129 | macro memo arena (16 MiB) cleared every ~2 s | records stored again with another shape left their old copies | compaction (9a37142); live records are ~12 MiB, so most compactions still end in a clear; hits stay ~97% (46% on the last call's slot) |
| 130 | measurements | a stuck `nvidia-smi -l` loop had spun one core for 12 h (every run since the evening of 10-28 had one core less); the attract loop shows different courses each run (frame means vary +-15%); stopped test loops kept running and overlapped | compare per-call costs (`C:\dev\nx2work\zones.py` on `EDEN_NX2_GPU_PROFILE` logs); `ab_env.sh` and `hangcheck.sh` refuse to start while eden-cli runs; `stopall.ps1` stops every loop |
| 131 | (title 1) 250-330 MB/s of buffer data copied back at fences (13,000-20,000 ranges a second) | Eden downloads every GPU-written range at the next fence. Measured: the native MME's DMA reads never touched those pages (`EDEN_NX2_GPU_PROFILE` counter), and with reactive flushing the guest CPU read GPU-written memory at most 64 times in 200 s | off for Switch 2 titles (d530bc4, `EDEN_NX2_BUFFER_DOWNLOADS=0/1` forces): fences ~1.1 us (were ~5), draw setup -7%, cache-lock spinning 1.8% |
| 132 | inline uploads 10.5% of the GPU thread | IsRegionGpuModified visited every GPU-written range in the query (they stay fragmented without downloads); std::hash of integers and pointers is FNV-1a on MSVC in every Common::unordered_map | RangeSet::Intersects (one lookup) and a pass-through hash for integer/pointer keys (0fe581f): 7.5% |
| 133 | some courses (desert/canyon, dense city, highway) run 30-45 ms with presents blocked 9-15 ms on the previous frame's fence (also before today: t14) | not the host GPU (nvidia-smi: 20-40% busy, ~22 W) but the GPU thread: ~100% busy there, its command processing 370-420 ms/s instead of 300-350 (more commands per frame). The 8 GB laptop GPU is also near its budget (6.6-6.9 of 7.1 GB): the texture collector evicts hundreds of decoded ASTC images per 600 frames, the buffer collector GPU-written buffers every frame | `EDEN_NX2_TC_MIN_AGE=120` (0f0fbb9) made no difference over three runs; buffer evictions batched into one GPU wait (db4c539), still ~2.7% on those courses (open: asynchronous eviction) |
| 134 | each inline write invalidated the buffer, texture, shader and query caches on its own | thousands of small writes between draws | guest memory written at once, pending query results dropped at once, GPU-written buffer data uploaded in place; the other caches get merged ranges at the next FlushCaching (c06ac80, `EDEN_NX2_INLINE_DEFER=0` off): inline writes 4.4% |
| 135 | ~6 million 3D register writes a second, ~60% writing the value the register already holds (mostly macro memo replays) | each went through shadow RAM, dirty flags (re-checked at the next draw) and the method switch | state registers (not in execution_mask) written with their value are dropped after shadow RAM processing (f265911, `EDEN_NX2_SKIP_SAME=0` off): macro calls 169 ns instead of 185, draw setup 5.80 us instead of 5.95. For non-executable registers ProcessMethodCall is skipped as well (it acts only on executable ones) |
| 136 | native compute: ~20% of the GPU thread under KeplerCompute::ProcessNativeMethod (dispatches 8%, inline QMD/uploads 4%) | its compute macros replay ~64 QMD register writes per dispatch, each through the macro/QMD/upload/semaphore checks | plain native registers stored directly from CallMethod (f7a82b2): its own time 2.2% -> 1.6% |
| 137 | 2D blits (Fermi2D) 3.6% on heavy courses | mip chains 64x64..4x1 copied near the GPU-only regions; neither image is GPU-written, so the software blitter (un)swizzles them on the CPU | open (expected for small copies; many of them) |
| 138 | `async_presentation=true` crashed 4 of 8 runs ~26 s in (row 118) | retested after the fence-operation race (row 118) and translation cache race (row 125) were fixed: 0 crashes in 4 runs, screenshots correct | on in the dev config (`async_presentation\default=false`, `async_presentation=true`); every later run is a crash test for it |
| 139 | buffer collector: up to 64 buffers evicted every frame, GPU-written ones downloaded with a GPU wait (~5% of the GPU thread on 8 GB) | it compared the whole device's use (textures, ~6.8 GB) with thresholds from the device size (6.14 GiB reported): aggressive every frame, while the buffers were ~200 MiB | Switch 2 titles collect on the buffers' own size (1 GiB, aggressive 2 GiB) and keep GPU-written buffers outside aggressive mode (c4eaa4b): 0%, presents no longer blocked on those courses (22-26 ms in t23) |
| 140 | macros ~20% of the GPU thread; the memo's lookups (~76 ns a call in the game, 48 ns in a warm harness) missed the CPU caches | memo records and its map fall out of cache between calls; 4-way last-call slots raised slot hits 48% -> 63% but saved nothing | x86-64 compiler for the Turing MME (a95cedd, `macro_turing_jit.h`, `EDEN_NX2_MME_JIT=0` off): identical to the interpreter on 200,000 recorded calls (`C:\dev\nx2work\mme_test\jit_test.cpp`); in the game 148 ns a call instead of 174, 0 failures, 0 crashes in 3 runs |
| 141 | draw setup ~6.4 us a draw (55,000 a second) | spread over ~12 steps: uniform buffer copies (2.4%, the fast path streams them every draw and bypasses write tracking), texture views (8.5% incl. ~3.6% streaming new course textures), handle reads, vertex/storage buffer lookups | no single cache-able step; open: small per-step trims. The way past it is a second thread (NX2_GPU_THREAD_SPLIT.md) |
| 142 | GPU thread ~100% busy, host GPU 20-40%: one thread parses, runs macros, sets state and records draws | the work is sequential but only a few points need a result from what came before (semaphore acquires, MME DMA reads) | GPU thread split (babd395, `gpu_split.h`, `EDEN_NX2_SPLIT=1`): engine methods run in order on a second thread fed by a lock-free ring; releases, syncpoints, wait-for-idle, flushes and channel binds are deferred into the ring, only acquires wait (~2,400 drains a second). ~47,400 draws a second against ~44,000 (3 A/B rounds), same per-call costs |
| 143 | crashes with the split: CommandChunk::Record (null chunk), vkCmdEndRenderPass on the Vulkan worker | two threads recording into the scheduler: (a) the DMA pusher's split-method debug log read guest memory with the cache-flushing ReadBlock on the front; (b) with async presentation, Swapchain::AcquireNextImage -> Scheduler::Wait flushed the main scheduler from the present thread whenever its tick was not submitted yet (pre-existing race, the split made it frequent) | (a) ReadBlockUnsafe; (b) e4523ab: the present thread waits up to 20 ms for the GPU thread to submit, never flushes. Found with `EDEN_NX2_SPLIT_CHECK=1` (stack of any other thread using the scheduler while the back runs) |
| 144 | macros on the split's front (`EDEN_NX2_SPLIT=2`) need the registers they read before the back has them | the front keeps copies: every write it sends plus what the back changes on its own (constant buffer offset, shadow_scratch[0], macro bind pointer, shadow RAM). `EDEN_NX2_SPLIT_VERIFY=1` found three differences: writes still in the back's method sink at the handover; a multi-word run that changes the shadow RAM mode mid-run; and shadow_state.shadow_ram_control being the shadow copy of register 0x49 itself (a Replay write reads the mode back) | all three modelled: 0 differences in 40,193 checks (~165 million macro calls). Front: macros 128 ns a call (180 on one thread, cache-hot), drops ~55% of 3D state writes as same-value. Back: 50,000-66,500 draws a second against 41,500-47,000 without the split; frames 20.5 / 25.5 ms on two courses where the unsplit runs measured 25.6-40.7 ms. Default stays off until a clean A/B |
| 145 | is the split worth having on by default | clean A/B, nothing else running, 3 rounds each (profiler on in both): one thread 41,100 / 48,600 / 37,000 draws a second, frames 30.5 / 35.0 / 27.1 ms; `EDEN_NX2_SPLIT=2` 59,200 / 62,400 / 62,700 draws a second, frames 28.6 / 23.3 / 25.1 ms. Throughput 1.46x, frames 17% shorter (the back is now ~92% busy, ~53% of it draw setup); 0 crashes, 0 register differences in 57,089 checks after the front's fast paths (no guarded statics, inline ring writes) | on by default for Switch 2 titles (`EDEN_NX2_SPLIT=0` one thread, `=1` macros on the back). Default run without the profiler: 16.9 ms frames (~59 fps) on a light course |
| 146 | title 4 aborts with InvalidHandle (0xE401) at its first movie, ~30 min into a new game | newer SDKs (`nn::codec::video::detail::InitializeServerContext`) open the decoder with gsv 906(u8, u64 library version 0x3'00030000) instead of 902; auto-stubbed, no handle came back | 906 hands out the 902 session (d8d11ec) |
| 147 | title 1: Grand Prix cup-opening screen stays black, decoder registered but never started; then an abort with 2111-3604 (0x1c286f) when it did start | `nn::syncpt::sdk::CancelSyncPointEvent` sends 202 and always waits on the slot's event; we signalled it only for pending requests. Then the movie player reused slot 0, which NVN's fence waits had left signalled, read it as a decoded frame and found none | 202 always signals, 201 clears the slot's event first (b09629f) |
| 148 | title 1/4 movies have no picture | the video decoder service (nn::codec::video, gsv 902/906/1002 -> vdec -> vdecsvc) completed frames without decoding | FFmpeg decodes on a thread per decoder service (`video_core/host1x/nx2_video_decoder`): 101 names the codec ("video/vp9") in an A buffer, 109 submits frames (A buffer, pts us). 311/314 reply NativeSyncHandles {~syncpoint, 0, value} on a syncpoint the service allocates, incremented per decoded frame, frames filling images in register/release order. gsv 911 creates the pool's images, 961 returns the CpuUserImageDescriptor that `ReadDecoderOutputImage*` copy out with (pitch-linear NV12/RGBA/BGRA, kinds 10/8/9). Title 1's Mushroom Cup opening (2560x1440 VP9, 1423 frames) decodes in real time and the game goes on to the race. The picture still stayed on the first frame: see 149. `EDEN_NX2_VDEC_DUMP=<dir>` saves decoded frames |
| 149 | title 1: the cup-opening movie's picture stays on its first frame although every frame decodes and reaches GPU memory | not the texture cache: the movie is drawn into a 2560x1440 render target whose NVN2 texture has GPU address 0 (also the 640x360 and 512x512 passes), so the draws go nowhere. NVN2's texture init (sdk+0x5ca240) places a texture in a compressible pool (flags 0x080000a1) only after `NvRmGpuMappingModify` switched the pool's mapping (0x815f0d0000, a gpu 214 region) to the texture's kind; NvRm fails that call with 2 unless the characteristics flags have NVGPU_GPU_FLAGS_SUPPORT_MAPPING_MODIFY (bit 34; ours were 0x4400155), and the texture keeps address 0 | report bit 34; the modify then runs through nn::gpu::sdk's VirtualMachine as gpu sub 212 (NormalModifyMemoryKind4K: words 3-6 range, 7 wanted kind, 8 fallback, 10 job slot; reply word 3 the kind applied), answered with the wanted kind and the slot signalled. The movie plays on screen |
Run settings used (in `build/bin/user/config/sdl2-config.ini`): `memory_layout_mode=4`
(12 GB). Debug knobs: `EDEN_NX2_TC_EXPECTED_MIB` / `EDEN_NX2_TC_CRITICAL_MIB` (collector thresholds),
`EDEN_NX2_ASTC_FLUSH=1` (GPU ASTC decode without the full wait), `EDEN_NX2_CRASH_DUMP=<dir>` (prints the
faulting stack as module+offset, then a minidump), `EDEN_NX2_AUTO_STUB`, `EDEN_NX2_IPC_TRACE`, `EDEN_NX2_SVC_TRACE`,
`EDEN_NX2_PULLER_TRACE`, `EDEN_NX2_METHOD_TRACE`, `EDEN_NX2_HIDE_WINDOW`,
`EDEN_NX2_GPU_PROFILE` (GPU-thread zone timers every 10 s), `EDEN_NX2_SAMPLE=<threads>` with
`_START`/`_SECONDS`/`_OUT` (sampling profiler; `C:\dev\nx2work\folded.py` reads its .folded stacks),
`EDEN_NX2_MME_RECORD=<file>` (macro calls for `C:\dev\nx2work\mme_test`),
`EDEN_NX2_FENCE_BATCH=0`, `EDEN_NX2_MACRO_COPY=1`, `EDEN_NX2_CNTPCT_PAUSE=0` (turn rows 124, 126, 127 off),
`EDEN_NX2_MME_MEMO_STATS` (memo counters),
`EDEN_NX2_BUFFER_DOWNLOADS=0/1`, `EDEN_NX2_INLINE_DEFER=0`, `EDEN_NX2_TC_MIN_AGE=<frames>` (rows 131, 134, 133),
`EDEN_NX2_SKIP_SAME=0` (row 135),
`EDEN_NX2_MME_JIT=0` (row 140),
`EDEN_NX2_SPLIT=0/1` (default 2), `EDEN_NX2_SPLIT_VERIFY=1`, `EDEN_NX2_SPLIT_CHECK=1`, `EDEN_NX2_SPLIT_SKIP=0` (rows 142-145),
`EDEN_NX2_FRAME_CAPTURE=<dir>` (every 120th presented frame as TGA), `EDEN_NX2_PB_CHECK`
(pushbuffer entries changed between submit and execution), `EDEN_NX2_QUERY_TRACE` (every GPU
report/semaphore write). `C:\dev\crashrun.sh` runs eden-cli under cdb and prints the crash stack;
`C:\dev\stackdump.sh <main> <seconds>` dumps host thread stacks of a running test.
## Where it stops now
All runs: `EDEN_NX2_AUTO_STUB=1`, 12 GB memory layout, loose ExeFS with the sibling `romfs/`.
**Title 4** (`fourth switch game\exefs`, 0400380020056000) **plays its intro cutscene correctly
rendered**: boot, loading screen, menu (difficulty select, keyboard input), confirm, then the intro
(plane cockpit, newspaper/map) lit, colour-graded and subtitled, matching console footage. The
DLSS-style upscale is a bilinear stand-in. In a played session the game went on to a later loading
screen that did not finish; open there: a compute shader with a 64-bit add (UIADD3 carry-out to
UP0), TMML (0xb69) in two graphics shaders, "reserved attribute 23" (a[0x5c]) reads. Scripted runs
(`C:\dev\run.sh`) move any title-4 save to `C:\dev\save_backups` first so they start a new game.
Scripted input: `EDEN_NX2_KEY_SCRIPT="50:4,50.5:-4,55:79,55.4:-79,57:4,57.5:-4"`.
**Title 3**: last checked before the graphics work; should now track title 4.
**Title 1** (`first switch game\exefs`, 0400C3F00006E000, `Drive.nss`): boots, shows the world-map
loading screen and reaches its title screen ("Press L + R to start") with sky, buildings, props,
kart and logo drawn. Wrong: magenta/cyan colour cast, black ground (unbacked virtual texture, row
87), task/mesh geometry missing (row 89), ~10-17 fps. Launch with frame capture:
`EDEN_NX2_FRAME_CAPTURE=C:/dev/frames_t1x EDEN_NX2_AUTO_STUB=1 EDEN_NX2_SHADER_DUMP=C:/dev/shaders_title1 ./eden-cli.exe -g "C:/switch/first switch game/exefs/main"`.
Offline shader check: `nx2work/lowertest/lowertest.exe C:/dev/shaders_title1 [-all [op]]`.
**Title 2**: main/sdk/subsdk0 fail LZ4 and SHA-256 checks; can't be run from these files.
## Prioritized next steps
1. **Title 4 past the intro**: lower the 64-bit add with carry (UIADD3/IADD3 carry-out and .X
carry-in), TMML, and attribute 0x5c; then whatever the next area needs (dump + lowertest).
Squid/paint-splat hatch on the menu. A faster tensor-core emulation (shared-memory HMMA) would
replace the bilinear upscale stand-in.
2. **Title 1 / title 3** retests with the current tree; dump their shaders and widen the lowering.
3. Replace guesses with real protocol knowledge: syncpt page layout, the semaphore mirror page,
gsv 214 (GPU-only memory) and vi2 present buffer availability (real buffer release instead of
the fence wait), gsv 1510 formats, svcGetInfo 29-33.
4. **Six cores** instead of folding 4/5 onto 0/1; rare SIMD encodings still undecoded.
5. Upstream-worthy pieces independent of Switch 2: MSVC 14.51 build fixes, eden-cli SDL3
lifetime fixes and `_setmaxstdio`, ExeFS loader fixes, FAC/FAH offset fix, GetPseudoDeviceId
without NACP.

View file

@ -67,7 +67,7 @@ For information on provided development tooling, see the [Tools directory](./too
## Download
You can download the latest releases from [our release page](https://git.eden-emu.dev/eden-emu/eden/releases).
You can download the latest releases from [our release page](https://git.eden-emu.dev/Takoroka/eden-nx2/releases).
Save us some bandwidth! We have [mirrors available](./docs/user/ThirdParty.md#mirrors) as well.

View file

@ -588,7 +588,11 @@ Result InfoUpdater::UpdateBehaviorInfo(BehaviorInfo& behaviour_) {
return Service::Audio::ResultInvalidUpdateInfo;
}
if (in_params->revision != behaviour_.GetUserRevision()) {
// Compare revision numbers: Switch 2 opens the renderer with 'REV?' but tags its updates
// 'REv?' (lowercase v) for the same revision.
if (GetRevisionNum(in_params->revision) != GetRevisionNum(behaviour_.GetUserRevision())) {
LOG_ERROR(Service_Audio, "BehaviorInfo revision {:08X} does not match the renderer's {:08X}",
in_params->revision, behaviour_.GetUserRevision());
return Service::Audio::ResultInvalidUpdateInfo;
}

View file

@ -4,8 +4,10 @@
// SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include <algorithm>
#include <array>
#include <atomic>
#include <cstdlib>
#include <memory>
#include <span>
#include <vector>
@ -213,6 +215,19 @@ void SinkStream::ProcessAudioOutAndRender(std::span<s16> output_buffer, std::siz
std::memcpy(last_frame.data(), &output_buffer[(frames_written - 1) * frame_size], frame_size_bytes);
// EDEN_NX2_AUDIO_LEVEL: log the output peak about every two seconds (is a game audible?).
if (static const bool meter = std::getenv("EDEN_NX2_AUDIO_LEVEL") != nullptr; meter) {
for (size_t i = 0; i < num_frames * frame_size; ++i) {
level_peak = std::max<s32>(level_peak, std::abs(static_cast<s32>(output_buffer[i])));
}
level_frames += num_frames;
if (level_frames >= TargetSampleRate * 2) {
LOG_INFO(Audio, "Audio output peak over {} frames: {} of 32767", level_frames, level_peak);
level_frames = 0;
level_peak = 0;
}
}
{
std::scoped_lock lk{sample_count_lock};
last_sample_count_update_time = system.CoreTiming().GetGlobalTimeNs();

View file

@ -245,6 +245,9 @@ private:
SinkBuffer playing_buffer{};
/// The last played (or received) frame of audio, used when the callback underruns
std::array<s16, MaxChannels> last_frame{};
/// EDEN_NX2_AUDIO_LEVEL meter state.
s32 level_peak{};
u64 level_frames{};
/// Number of buffers waiting to be played
std::atomic<u32> queued_buffers{};
/// The ring size for audio out buffers (usually 4, rarely 2 or 8)

View file

@ -96,6 +96,7 @@ add_library(
${CMAKE_CURRENT_BINARY_DIR}/scm_rev.cpp
scm_rev.h
scope_exit.h
spin_recursive_mutex.h
scratch_buffer.h
settings.cpp
settings.h
@ -151,7 +152,7 @@ add_library(
set_source_files_properties(zbic_compression.cpp PROPERTIES
INCLUDE_DIRECTORIES "${ZBIC_INCLUDE_DIR}"
COMPILE_OPTIONS "$<$<CXX_COMPILER_ID:Clang,GNU>:-Wno-unused-function;-Wno-missing-declarations;-Wno-shadow>")
COMPILE_OPTIONS "$<$<CXX_COMPILER_ID:Clang,GNU>:-Wno-unused-function;-Wno-missing-declarations;-Wno-shadow>;$<$<CXX_COMPILER_ID:MSVC>:/wd4505;/wd5245>")
if(WIN32)
target_sources(common PRIVATE windows/timer_resolution.cpp

View file

@ -9,7 +9,7 @@
namespace Common {
template <class Key, class T, class Hash = std::hash<Key>, class Pred = std::equal_to<Key>,
template <class Key, class T, class Hash = FastHash<Key>, class Pred = std::equal_to<Key>,
class Allocator = std::allocator<std::pair<const Key, T>>>
using unordered_map = boost::unordered::unordered_flat_map<Key, T, Hash, Pred, Allocator>;

View file

@ -9,7 +9,7 @@
namespace Common {
template <class Key, class Hash = std::hash<Key>, class Pred = std::equal_to<Key>,
template <class Key, class Hash = FastHash<Key>, class Pred = std::equal_to<Key>,
class Allocator = std::allocator<Key>>
using unordered_set = boost::unordered::unordered_flat_set<Key, Hash, Pred, Allocator>;

View file

@ -19,6 +19,28 @@
namespace Common {
/// The default hash of Common::unordered_map/set. std::hash of an integer or a pointer is a
/// byte-wise FNV-1a on MSVC, and boost's flat containers mix a hash that does not say it
/// avalanches anyway: pass those keys through (the GPU thread's page maps looked keys up
/// tens of thousands of times a second with it). Other keys keep std::hash.
template <class Key>
struct FastHash : std::hash<Key> {};
template <class Key>
requires std::is_integral_v<Key>
struct FastHash<Key> {
std::size_t operator()(Key key) const noexcept {
return static_cast<std::size_t>(key);
}
};
template <class T>
struct FastHash<T*> {
std::size_t operator()(T* pointer) const noexcept {
return reinterpret_cast<std::size_t>(pointer);
}
};
namespace detail {
template <typename T>

View file

@ -225,6 +225,13 @@ const CPUCaps g_cpu_caps = [] {
}
}
// AMD CPUs have no leaf 0x15 (Zen 4 stops at 0x10): measure their invariant TSC too, or every
// guest clock read (CNTPCT, which games read thousands of times a frame) goes through the
// host's much slower QueryPerformanceCounter.
if (caps.invariant_tsc && caps.tsc_frequency == 0) {
caps.tsc_frequency = X64::EstimateRDTSCFrequency();
}
if (max_std_fn >= 0x16) {
__cpuid(cpu_id, 0x16);
caps.base_frequency = cpu_id[0];

View file

@ -26,6 +26,9 @@ public:
void Clear();
bool Empty() const;
/// Whether any range overlaps [base_address, base_address + size): one lookup.
bool Intersects(AddressType base_address, size_t size) const;
template <typename Func>
void ForEach(Func&& func) const;

View file

@ -116,28 +116,25 @@ struct OverlapRangeSet<AddressType>::OverlapRangeSetImpl {
}
AddressType end_address = base_address + static_cast<AddressType>(size);
IntervalType interval{base_address, end_address};
bool any_removals = false;
m_split_ranges_set += std::make_pair(interval, -amount);
do {
any_removals = false;
auto it = m_split_ranges_set.lower_bound(interval);
if (it == m_split_ranges_set.end()) {
return;
// One pass, erasing the segments that dropped to zero or below in address order (the
// search restarted after every erase: quadratic in the segments of the range, and the
// GPU fencing thread does this for each downloaded buffer range while holding the buffer
// cache's lock). Erasing a node leaves the other iterators valid.
auto it = m_split_ranges_set.lower_bound(interval);
const auto end_it = m_split_ranges_set.upper_bound(interval);
while (it != end_it) {
if (it->second > 0) {
++it;
continue;
}
auto end_it = m_split_ranges_set.upper_bound(interval);
for (; it != end_it; it++) {
if (it->second <= 0) {
if constexpr (has_on_delete) {
if (it->second == 0) {
on_delete(it->first.lower(), it->first.upper());
}
}
any_removals = true;
m_split_ranges_set.erase(it);
break;
if constexpr (has_on_delete) {
if (it->second == 0) {
on_delete(it->first.lower(), it->first.upper());
}
}
} while (any_removals);
m_split_ranges_set.erase(it++);
}
}
template <typename Func>
@ -223,6 +220,18 @@ bool RangeSet<AddressType>::Empty() const {
return m_impl->m_ranges_set.empty();
}
template <typename AddressType>
bool RangeSet<AddressType>::Intersects(AddressType base_address, size_t size) const {
const auto& set = m_impl->m_ranges_set;
if (set.empty() || size == 0) {
return false;
}
const AddressType end_address = base_address + static_cast<AddressType>(size);
// The first range not entirely below the query; it overlaps when it starts before its end.
const auto it = set.lower_bound(typename RangeSetImpl::IntervalType{base_address, end_address});
return it != set.end() && it->lower() < end_address;
}
template <typename AddressType>
template <typename Func>
void RangeSet<AddressType>::ForEach(Func&& func) const {

View file

@ -26,7 +26,7 @@ static std::mutex vector_regions_mutex {};
// Workaround for handling non-commited memory accessed by Dynarmic; usually result of an error
static LONG WINAPI FakePageFaultHandler(PEXCEPTION_POINTERS info) {
DWORD code = info->ExceptionRecord->ExceptionCode;
u64 exception_addr = reinterpret_cast<u64>(info->ExceptionRecord->ExceptionInformation[1]);
u64 exception_addr = static_cast<u64>(info->ExceptionRecord->ExceptionInformation[1]);
if (code != EXCEPTION_ACCESS_VIOLATION || info->ExceptionRecord->ExceptionInformation[0] == 1) {
// Not our problem

View file

@ -0,0 +1,98 @@
// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
// SPDX-License-Identifier: GPL-3.0-or-later
#pragma once
#include <atomic>
#include <cstdint>
#include <thread>
#if defined(_M_X64) || defined(_M_IX86) || defined(__x86_64__) || defined(__i386__)
#include <immintrin.h>
#endif
namespace Common {
/// A recursive mutex that spins briefly before sleeping. The GPU caches' locks are taken tens of
/// thousands of times a second by the GPU thread, the fencing thread and CPU write traps, and
/// held for microseconds: std::recursive_mutex put a waiting thread to sleep at once, and its
/// wake-up took longer than the critical section (title 1's GPU thread spent ~8% asleep there).
class SpinRecursiveMutex {
public:
void lock() {
const std::thread::id self = std::this_thread::get_id();
if (owner.load(std::memory_order_relaxed) == self) {
++depth;
return;
}
std::uint32_t expected = 0;
if (!state.compare_exchange_strong(expected, 1, std::memory_order_acquire,
std::memory_order_relaxed)) {
LockContended();
}
owner.store(self, std::memory_order_relaxed);
depth = 1;
}
bool try_lock() {
const std::thread::id self = std::this_thread::get_id();
if (owner.load(std::memory_order_relaxed) == self) {
++depth;
return true;
}
std::uint32_t expected = 0;
if (!state.compare_exchange_strong(expected, 1, std::memory_order_acquire,
std::memory_order_relaxed)) {
return false;
}
owner.store(self, std::memory_order_relaxed);
depth = 1;
return true;
}
void unlock() {
if (--depth != 0) {
return;
}
owner.store(std::thread::id{}, std::memory_order_relaxed);
if (state.exchange(0, std::memory_order_release) == 2) {
state.notify_one();
}
}
private:
// About 10 microseconds of spinning before sleeping.
static constexpr int SpinCount = 1024;
static void Pause() {
#if defined(_M_X64) || defined(_M_IX86) || defined(__x86_64__) || defined(__i386__)
_mm_pause();
#elif defined(__aarch64__)
asm volatile("yield");
#endif
}
void LockContended() {
for (int spin = 0; spin < SpinCount; ++spin) {
Pause();
if (state.load(std::memory_order_relaxed) == 0) {
std::uint32_t expected = 0;
if (state.compare_exchange_weak(expected, 1, std::memory_order_acquire,
std::memory_order_relaxed)) {
return;
}
}
}
// 2: held with (possible) sleepers; taken this way it stays marked, so its unlock wakes
// the next sleeper.
while (state.exchange(2, std::memory_order_acquire) != 0) {
state.wait(2, std::memory_order_relaxed);
}
}
std::atomic<std::uint32_t> state{0}; ///< 0 free, 1 held, 2 held with sleepers
std::atomic<std::thread::id> owner{};
std::uint32_t depth{}; ///< Recursion depth, touched only by the owner
};
} // namespace Common

View file

@ -457,6 +457,7 @@ add_library(core STATIC
hle/service/am/service/application_functions.cpp
hle/service/am/service/application_functions.h
hle/service/am/service/application_proxy.cpp
hle/service/am/service/graphics_scale_controller.h
hle/service/am/service/application_proxy.h
hle/service/am/service/application_proxy_service.cpp
hle/service/am/service/application_proxy_service.h
@ -526,6 +527,7 @@ add_library(core STATIC
hle/service/audio/audio_out.h
hle/service/audio/audio_renderer_manager.cpp
hle/service/audio/audio_renderer_manager.h
hle/service/audio/nx2_audio_user_client.h
hle/service/audio/audio_renderer.cpp
hle/service/audio/audio_renderer.h
hle/service/audio/audio.cpp
@ -865,6 +867,7 @@ add_library(core STATIC
hle/service/nvdrv/devices/nvhost_vic.h
hle/service/nvdrv/devices/nvmap.cpp
hle/service/nvdrv/devices/nvmap.h
hle/service/nvdrv/devices/nx2_stub_device.h
hle/service/nvdrv/nvdata.h
hle/service/nvdrv/nvdrv.cpp
hle/service/nvdrv/nvdrv.h
@ -1033,6 +1036,8 @@ add_library(core STATIC
hle/service/server_manager.h
hle/service/service.cpp
hle/service/service.h
hle/service/nx2_stub_interface.h
hle/service/nx2_gsv.h
hle/service/services.cpp
hle/service/services.h
hle/service/set/factory_settings_server.cpp

View file

@ -4,6 +4,11 @@
// SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include <atomic>
#include <cstdlib>
#if defined(_M_X64) || defined(__x86_64__)
#include <immintrin.h>
#endif
#include "common/settings.h"
#include "core/arm/dynarmic/arm_dynarmic.h"
#include "core/arm/dynarmic/arm_dynarmic_64.h"
@ -24,13 +29,24 @@ DynarmicCallbacks64::DynarmicCallbacks64(ArmDynarmic64& parent, Kernel::KProcess
u64 DynarmicCallbacks64::MemoryRead(u64 vaddr, size_t size) {
CheckMemoryAccess(vaddr, size, Kernel::DebugWatchpointType::Read);
u64 value{};
switch (size) {
case sizeof(u64): return m_memory.Read64(vaddr);
case sizeof(u32): return m_memory.Read32(vaddr);
case sizeof(u16): return m_memory.Read16(vaddr);
case sizeof(u8): return m_memory.Read8(vaddr);
case sizeof(u64): value = m_memory.Read64(vaddr); break;
case sizeof(u32): value = m_memory.Read32(vaddr); break;
case sizeof(u16): value = m_memory.Read16(vaddr); break;
case sizeof(u8): value = m_memory.Read8(vaddr); break;
default: UNREACHABLE();
}
// With memory aborts ignored the guest keeps running on zeros: say where it read from.
if (auto& unmapped = Core::Memory::UnmappedReadFlag(); unmapped != 0) {
unmapped = 0;
static std::atomic<u32> logged{};
if (logged.fetch_add(1) < 32) {
LOG_ERROR(Core_ARM, "Unmapped read {:#x} ({} bytes) at pc {:#x} lr {:#x} sp {:#x}", vaddr, size,
m_parent.m_jit->GetPC(), m_parent.m_jit->GetRegister(30), m_parent.m_jit->GetSP());
}
}
return value;
}
Dynarmic::A64::Vector DynarmicCallbacks64::MemoryRead128(u64 vaddr) {
CheckMemoryAccess(vaddr, 16, Kernel::DebugWatchpointType::Read);
@ -154,7 +170,26 @@ u64 DynarmicCallbacks64::GetTicksRemaining() {
}
u64 DynarmicCallbacks64::GetCNTPCT() {
return m_parent.m_system.CoreTiming().GetClockTicks();
const u64 ticks = m_parent.m_system.CoreTiming().GetClockTicks();
// Guest code that waits by polling the counter (title 1's main thread reads it ~17 million
// times a second while waiting for the GPU) keeps a host core busy at full speed, and the
// GPU thread is often on the other half of that SMT core. Two reads within ~1.7 us of each
// other are such a loop: pause, which hands the core's resources to the sibling thread for a
// moment. The guest sees the same counter. EDEN_NX2_CNTPCT_PAUSE=0 turns it off.
static const bool pause_spins = [] {
const char* value = std::getenv("EDEN_NX2_CNTPCT_PAUSE");
return value == nullptr || value[0] != '0';
}();
if (pause_spins && ticks - m_last_cntpct < 32) {
#if defined(_M_X64) || defined(__x86_64__)
_mm_pause();
_mm_pause();
#elif defined(__aarch64__)
asm volatile("yield");
#endif
}
m_last_cntpct = ticks;
return ticks;
}
bool DynarmicCallbacks64::CheckMemoryAccess(u64 addr, u64 size, Kernel::DebugWatchpointType type) {

View file

@ -65,6 +65,7 @@ public:
Core::Memory::Memory& m_memory;
u64 m_tpidrro_el0{};
u64 m_tpidr_el0{};
u64 m_last_cntpct{}; ///< The previous counter read (GetCNTPCT spin detection)
Kernel::KProcess* m_process{};
const bool m_debugger_enabled{};
const bool m_check_memory_access{};

View file

@ -24,6 +24,7 @@
#include "core/device_memory.h"
#include "core/file_sys/fs_filesystem.h"
#include "core/file_sys/patch_manager.h"
#include "core/file_sys/program_metadata.h"
#include "core/file_sys/registered_cache.h"
#include "core/file_sys/romfs_factory.h"
#include "core/file_sys/savedata_factory.h"
@ -68,6 +69,8 @@
namespace Core {
std::atomic<u64> g_nx2_present_count{};
FileSys::VirtualFile GetGameFileFromPath(const FileSys::VirtualFilesystem& vfs,
const std::string& path) {
// To account for split 00+01+etc files.
@ -107,6 +110,19 @@ FileSys::VirtualFile GetGameFileFromPath(const FileSys::VirtualFilesystem& vfs,
return vfs->OpenFile(path, FileSys::OpenMode::Read);
}
namespace {
// Switch 2 titles boot as a loose ExeFS, with main.npdm next to main.
bool IsSwitch2Program(const FileSys::VirtualFile& file) {
const auto dir = file != nullptr ? file->GetContainingDirectory() : nullptr;
const auto npdm = dir != nullptr ? dir->GetFile("main.npdm") : nullptr;
FileSys::ProgramMetadata metadata;
return npdm != nullptr && metadata.Load(npdm) == Loader::ResultStatus::Success &&
metadata.IsSwitch2();
}
} // namespace
struct System::Impl {
explicit Impl(System& system)
: kernel{system}, fs_controller{system}, hid_core{kernel}, cpu_manager{system},
@ -297,6 +313,16 @@ struct System::Impl {
fs_controller.InitTempStorage();
}
// Switch 2 titles size their heap from the total memory and need the 12 GiB layout. Set it
// as a per-game value, so RestoreGlobalState brings back the configured layout afterwards.
if (virtual_filesystem != nullptr &&
Settings::values.memory_layout_mode.GetValue() != Settings::MemoryLayout::Memory_12Gb &&
IsSwitch2Program(GetGameFileFromPath(virtual_filesystem, filepath))) {
LOG_INFO(Core, "Switch 2 program: using the 12 GiB memory layout");
Settings::values.memory_layout_mode.SetGlobal(false);
Settings::values.memory_layout_mode.SetValue(Settings::MemoryLayout::Memory_12Gb);
}
InitializeKernel(system);
if (params.applet_type == Service::AM::AppletType::Application) {

View file

@ -6,6 +6,7 @@
#pragma once
#include <atomic>
#include <cstddef>
#include <deque>
#include <functional>
@ -115,6 +116,9 @@ class RenderdocAPI;
namespace Core {
/// Frames the guest has presented (Switch 2 vi2 present calls); input replay is locked to it.
extern std::atomic<u64> g_nx2_present_count;
class CpuManager;
class Debugger;
class DeviceMemory;

Some files were not shown because too many files have changed in this diff Show more