Compare commits

..
Author SHA1 Message Date
crueter
a2aa4d97e1
[docs] Remove AI restrictions, and adapt to new open registration (#4585)
Some checks failed
tx-src / sources (push) Has been cancelled
Check Strings / check-strings (push) Has been cancelled
Registration is now open. Enjoy

Signed-off-by: crueter <crueter@eden-emu.dev>

- [x] I have read and followed the [Contribution Guidelines](https://git.eden-emu.dev/eden-emu/eden/src/branch/master/CONTRIBUTING.md#code-contributions).
- [x] I have read and followed the [Coding Guidelines](https://git.eden-emu.dev/eden-emu/eden/src/branch/master/docs/policies/Coding.md) to the best of my ability.

-------------------

Reviewed-on: https://git.eden-emu.dev/eden-emu/eden/pulls/4585
Reviewed-by: lizzie <lizzie@eden-emu.dev>
2026-10-29 17:10:37 +01:00
211 changed files with 418 additions and 16483 deletions

View file

@ -1,5 +1,4 @@
- [ ] I have read and followed the [Contribution Guidelines](https://git.eden-emu.dev/eden-emu/eden/src/branch/master/CONTRIBUTING.md#code-contributions).
- [ ] I have read and followed the [AI Policy](https://git.eden-emu.dev/eden-emu/eden/src/branch/master/docs/policies/AI.md).
- [ ] I have read and followed the [Coding Guidelines](https://git.eden-emu.dev/eden-emu/eden/src/branch/master/docs/policies/Coding.md) to the best of my ability.
-------------------

View file

@ -16,15 +16,14 @@ Eden is free, open-source, copyleft software, licensed under the terms of the [G
### Policies
- We are trialing a slightly relaxed AI policy to allow high-quality contributors to use AI *responsibly*. Please read our [detailed AI policy](docs/policies/AI.md) for more info.
- Usage of any form of profanity or otherwise unsavory language is generally discouraged.
- This is primarily because it rarely helps to actually understand what's going on.
- Remember that your comments should be focused and actually address what's happening. With very few exceptions, expletives are actively detrimental at best.
- New code must follow the same general style as the surrounding codebase. Exceptions may be granted in certain cases.
- AI usage is allowed, but we reserve the right to close and rework your PR if you don't fully understand the changes you're making.
- As a common courtesy, please disclose your AI usage under the PR description, as `Assisted-by: <AI model> (<how it helped>)`.
- Usage of any form of profanity or otherwise unsavory language is discouraged but not strictly disallowed.
- New code must follow the same style as the surrounding codebase. Exceptions may be granted in certain cases.
- Maintainers reserve the right to change your patches and pull requests at will. We will try to avoid this.
- You should generally respect all decisions made by the [code owners](docs/CODEOWNERS) in your particular subsystem.
- However, if you feel they are overstepping or are incorrect, don't be afraid to stand your ground! Maintainers are not always correct.
- While we do *not* adhere to the terms of a formal code of conduct, you will generally be expected to respect other developers, contributors, and community members.
- While we do *not* adhere to the terms of a formal code of conduct, you are expected to treat other developers, contributors, and community members respectfully.
- You **must** have basic knowledge of [Git](https://git-scm.com/learn). Knowing how to manage your branches and follow proper fork policies is a necessity.
### Where do I start?
@ -38,7 +37,8 @@ You may also search for open issues on [GitHub](https://github.com/eden-emulator
There are five primary ways to contribute.
- Submit pull requests directly to our source tree
- You must first request an account on our Forgejo. See [Signup](#how-do-i-sign-up)
- You must first [create an account](https://git.eden-emu.dev/user/sign_up).
- You may use an internal account, or you may use a GitHub/Codeberg account linked to our instance.
- Submit pull requests to our [GitHub mirror](https://github.com/eden-emulator/Issue-Reports)
- Note that this is subject to being removed by DMCA, so don't rely on this.
- Submit pull requests to our [Codeberg mirror](https://codeberg.org/eden-emu/eden)
@ -50,31 +50,8 @@ There are five primary ways to contribute.
To test your changes, ensure to read the [build documentation](./docs/Build.md) for your specific platform(s) to ensure everything compiles and works properly. You should also make sure that your branch is up-to-date with the upstream `master` branch before opening a pull request.
### How do I sign up?
<details>
<summary>To sign up and begin contributing... (click to open)</summary>
Email [crueter@crueter.xyz](mailto:crueter@crueter.xyz) with the following format:
```txt
Subject: [Eden Git] Registration Request
Username: <Your Desired Username>
Email: <Your Desired Email>
I wish to sign up because... <your reason here>
```
Received mail that does not follow this format will be ignored.
Once your request is processed, you will receive a confirmation email with your temporary password and some information on Git access and policies. If you do not receive a response within 48 hours, feel free to send another email.
> [!WARNING]
> Some email providers may place the response email in your spam/junk folder; notable offenders include Gmail and Outlook. *Always* ensure to check your Spam or Junk folder.
</details>
## Non-code Contributions
Alongside the other contribution methods listed up top, you can also choose to contribute through documentation, organization, or community guides. These can be done either through the code contribution methods described above, or created externally and shared via our Discord community.
If you have an external tool/page that you believe would be handy to integrate/link into Eden, please additionally email our developers at [`developers@eden-emu.dev`](mailto:developers@eden-emu.dev). **Do not submit vibe-coded or AI generated tools or applications**.
If you have an external tool/page that you believe would be handy to integrate/link into Eden, please additionally email our developers at [`developers@eden-emu.dev`](mailto:developers@eden-emu.dev).

View file

@ -1,63 +0,0 @@
# 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.

View file

@ -1,264 +0,0 @@
# 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/Takoroka/eden-nx2/releases).
You can download the latest releases from [our release page](https://git.eden-emu.dev/eden-emu/eden/releases).
Save us some bandwidth! We have [mirrors available](./docs/user/ThirdParty.md#mirrors) as well.

View file

@ -1,203 +0,0 @@
# AI Policy
> [!WARNING]
> We are trialing a slightly relaxed AI policy to allow high-quality contributors to use AI *responsibly*.
> This policy is subject to change at any point in time, so check back frequently.
>
> This banner will be removed once we have finalized our revised policy.
- [AI Policy](#ai-policy)
- [TL;DR](#tldr)
- [Definitions](#definitions)
- [Scope](#scope)
- [Introduction](#introduction)
- [Requirements](#requirements)
- [Sign-off, Authorship, and Copyright](#sign-off-authorship-and-copyright)
- [Permitted Usage](#permitted-usage)
- [Recommended Prior Discussion](#recommended-prior-discussion)
- [Forbidden Usage](#forbidden-usage)
- [Human Effort Denial of Service](#human-effort-denial-of-service)
- [Disclosure](#disclosure)
- [Codebase Ingestion](#codebase-ingestion)
- [Enforcement](#enforcement)
- [Addendum: AI Commit Messages](#addendum-ai-commit-messages)
- [References](#references)
## TL;DR
AI-assisted contributions are welcome, provided you take responsibility for them.
- You **must** disclose any AI or LLM usage in your pull request description *and* commit message.
- You **must** be able to explain, defend, and revise every line of code you submit.
- Commit messages, PR descriptions, issues, and comments **must** be written by you. Verbatim AI output is not acceptable.
- Undisclosed AI usage will result in your contribution being rejected.
- Low-quality contributions, disclosed or not, will be reworked or rejected.
- Spamming automated PRs, comments, or issues is treated as a Denial of Service (DoS) attack on our reviewers; see [Human Effort Denial of Service](#human-effort-denial-of-service).
## Definitions
An *AI or LLM-based tool* is any piece of software that generates code, text, or other content through machine learning models. This includes, but is not limited to: chat assistants, agentic coding tools/harnesses, inline code completion suggestions you have accepted into the code, and AI-driven explanation/refactoring features.
For the purposes of this policy, basic deterministic/static development tools (such as clang-tidy, git, compilers, etc.) are not considered AI tools.
Terms such as "substantially", "large-scale", "majority", and "low-quality" are intentionally left vague, and will be interpreted at the maintainers' discretion. When in doubt, ask first.
## Scope
This policy applies to all code contributions, alongside issues and pull requests surrounding this codebase. This includes:
- Commit messages
- PR descriptions
- Code review comments
- Issue/PR comments
- Patches submitted via e-mail
Areas that are not directly related to code--such as our surrounding communities--are not covered by the terms of this policy. Additionally, spelling/grammar checking via AI or LLM-based tools is not covered under this document.
## Introduction
In the right hands, AI and LLM-based tools are incredibly useful for getting features and bugfixes out the door as fast as possible. At the same time, reviewer bandwidth is a scarce resource. We are therefore open to AI-assisted contributions, provided they are produced and disclosed properly.
The goal of this document is *not* to set arbitrary restrictions, but rather to keep our development healthy and stable, while still allowing contributors to use the tools they find helpful. A contribution must be worth more to the project than the time it takes to review it.
Throughout this document, **must** and **must not** are inflexible, binding requirements. Refusal to follow these is treated as a violation of the terms outlined in this document.
## Requirements
These apply to all contributions, whether AI-assisted or not. AI-assisted contributions also carry additional disclosure requirements as outlined in [Disclosure](#disclosure).
1. *Disclosure*. Any and all usage of AI or LLM-based tools within any contribution **must** be explicitly marked and disclosed as appropriate.
2. *Responsibility*. The human submitter of any contribution is held wholly responsible for the accuracy, completeness, and intent of what is submitted. You **must** be able to explain, defend, and revise your code during review.
3. *Review*. All code **must** be thoroughly reviewed by the author before opening a PR or marking it ready for review.
4. *Testing*. All code changes **must** be tested directly by the human author.
5. *Quality*. All submitted code is held to a strict standard of excellence. Low-quality AI-generated content does not uphold this standard.
6. *Licensing*. All code **must** be compatible with the GPLv3 or later. You **must** warrant that you have the right to submit it under those terms, and that it was not copied from work published under an incompatible license.
## Sign-off, Authorship, and Copyright
AI agents **must not** add `Signed-off-by:` tags. The human submitter **must** sign off on all changes.
Similarly, AI agents cannot hold authorship or copyright over code, and **must not** add `Co-authored-by:` or `Co-developed-by:` tags. The human submitter **must** author all changes.
## Permitted Usage
This is a non-exhaustive list of permitted uses of AI and LLM-based tools, provided they are disclosed properly.
- Drafting issues and pull requests, provided the final text and submitted code are substantially your own.
- Debugging issues, bugs, glitches, and crashes, whether pre-existing or introduced by your own changes.
- Reviewing your own changes before submitting a PR.
- Research, brainstorming, and codebase searching/understanding.
- Applying tedious but not particularly large changes, such as mass-migrating specific classes to polyfills or abstractions.
## Recommended Prior Discussion
Some usage cases are not necessarily forbidden nor permitted. These should be discussed with maintainers in an issue or e-mail *before* submission. Maintainers and reviewers may also comment on this during the review/submission phase.
- Generating boilerplate/scaffolding for new code.
- Writing unit tests.
- Large-scale refactors primarily driven by an AI agent.
- Any contribution where an AI agent or model was responsible for the majority of the implementation.
## Forbidden Usage
This is a non-exhaustive list of strictly forbidden uses of AI and LLM-based tools.
- Vibecoding, or changes entirely authored by AI agents without human intervention or understanding.
- If you don't choose to put in the necessary effort to understand every line of code you submitted, neither will we.
- Fully AI-generated commit messages, PR descriptions, issues, or comments. See [Addendum: AI Commit Messages](#addendum-ai-commit-messages).
- Submission of AI-assisted work that is not disclosed per the [Disclosure](#disclosure) requirements.
- Low-quality submissions that fail to satisfy our standard of code excellence.
- Submitting secrets, credentials, private user data, or unpublished security reports to third-party AI services.
- This includes console keys, encryption material, and private vulnerability disclosures.
- Verbatim AI code reviews (e.g. Codex/Copilot reviews).
## Human Effort Denial of Service
Using automated tools and AI to submit PRs, issues, and comments that we have to carefully review and handle is considered a [Denial-of-service attack](https://en.wikipedia.org/wiki/Denial-of-service_attack) on our human effort.
It takes very little effort on your side (an LLM prompt) to generate a large amount of effort on ours (careful and thorough code review). Please respect our time and don't do that.
## Disclosure
How you disclose AI usage is up to you. However, it **must** be explicit, easy to identify, and placed both within the commit message *and* the surrounding PR description. For patches submitted by e-mail, the cover letter counts as the PR description. We recommend an `Assisted-by:` footer:
```txt
Assisted-by: <model> (<what it was used for>)
```
For instance:
```txt
Assisted-by: Claude Opus 5.5 (initial draft and code review)
```
Contributions without a commit message (e.g. issues/comments) should instead use a plain sentence in the text body such as `I used <model> to help with <X>`. PRs and commit messages may opt to use this format as well.
Multi-commit series/squashes **must** maintain the trailer. Maintainers will handle this for merges to master, but you are responsible for ensuring the disclosure survives into the final commit/PR body of squashes within your branches.
Regardless, what matters is that reviewers can tell *at a glance* where and how AI was involved in your contribution.
## Codebase Ingestion
Ingestion of this codebase by LLMs for any purpose--such as training or prompt context--is permitted *for now*. This may change at any point as litigation and court rulings around AI and code copyright proceed.
AI models may have ingested AGPLv3 or otherwise copyleft-incompatible code and may regurgitate it. You remain responsible for ensuring that all code you submit is compatible with the GPLv3 or later.
## Enforcement
- If we suspect undisclosed AI usage in a contribution, we will ask for disclosure.
- Refusal to provide disclosure will result in rejection of the contribution, and your PR or issue will be closed. Maintainers may choose to rework it themselves if the change is worth having.
- Accounts that spam us with repeated automated PRs or comments will be blocked or permanently suspended from our Git instance.
- Maintainers reserve the right to close your issues and pull requests if you don't follow the guidelines and procedures outlined in this document, and refuse or fail to make changes as necessary to comply.
- Repeated failures to comply may result in loss of contribution privileges.
## Addendum: AI Commit Messages
The patchset for pull request [#3422](https://git.eden-emu.dev/eden-emu/eden/pulls/3422) was fed into a dedicated code AI model, and it produced the following:
```txt
profile_manager: Refactor user deletion to use index instead of UUID
This commit improves the user profile deletion flow in the following ways:
1. **UI layer (ConfigureProfileManager)**:
- Replaces UUID-based deletion with index-based deletion
- Passes the selected row index from the tree view to the confirmation dialog
- Emits deleteUser(int index) signal from ConfigureProfileManagerDeleteDialog
- Removes lambda capture and std::function callback in favor of direct signal/slot connection
2. **Delete dialog**:
- Changes SetInfo() signature to accept index instead of accept callback
- Stores index internally and emits it via deleteUser signal on acceptance
3. **Core ProfileManager**:
- Refactors RemoveUser(UUID) → RemoveProfileAtIndex(index)
- Removes duplicated stable_partition + save logic from RemoveUser()
- Centralizes profile removal and save-file writing in RemoveProfileAtIndex()
4. **Minor follow-up**:
- Adds profile_manager.ResetUserSaveFile() call in PopulateUserList() to ensure consistent state when refreshing the user list
These changes eliminate redundant code, reduce the risk of UUID → index mismatch bugs, and make the deletion path more straightforward and maintainable.
Fixes: potential inconsistency when user order changes between UI and backend
```
This:
- Is needlessly verbose (nobody cares about most of these details)
- Doesn't address the actual purpose of the PR (fixing a double-deletion bug in the profile manager)
- Has unicode arrows (this is bad for command-line editors)
- Uses corporate and word-salad language that pads out the message to make it look like a lot changed
A good commit message should be concise, descriptive, and convey the *intent* behind the change, such as the bug being fixed or the feature being added. It should *not* be a detailed changelog of every individual byte that's changed.
## References
This policy document has been adapted from the AI policies of numerous other free and open-source projects, whether restrictive or permissive.
- [Gitea](https://github.com/go-gitea/gitea/blob/main/CONTRIBUTING.md#ai-contribution-policy)
- The [Linux kernel](https://kernel.org/doc/html/next/process/coding-assistants.html)
- [Gentoo Linux](https://wiki.gentoo.org/wiki/Project:Council/AI_policy)
- [cURL](https://curl.se/dev/contribute.html#on-ai-use-in-curl)
- [FastAPI](https://tiangolo.com/open-source/contributing/#automated-code-and-ai)

View file

@ -588,11 +588,7 @@ Result InfoUpdater::UpdateBehaviorInfo(BehaviorInfo& behaviour_) {
return Service::Audio::ResultInvalidUpdateInfo;
}
// 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());
if (in_params->revision != behaviour_.GetUserRevision()) {
return Service::Audio::ResultInvalidUpdateInfo;
}

View file

@ -4,10 +4,8 @@
// 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>
@ -215,19 +213,6 @@ 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,9 +245,6 @@ 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,7 +96,6 @@ 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
@ -152,7 +151,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>;$<$<CXX_COMPILER_ID:MSVC>:/wd4505;/wd5245>")
COMPILE_OPTIONS "$<$<CXX_COMPILER_ID:Clang,GNU>:-Wno-unused-function;-Wno-missing-declarations;-Wno-shadow>")
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 = FastHash<Key>, class Pred = std::equal_to<Key>,
template <class Key, class T, class Hash = std::hash<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 = FastHash<Key>, class Pred = std::equal_to<Key>,
template <class Key, class Hash = std::hash<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,28 +19,6 @@
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,13 +225,6 @@ 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,9 +26,6 @@ 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,25 +116,28 @@ 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);
// 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;
do {
any_removals = false;
auto it = m_split_ranges_set.lower_bound(interval);
if (it == m_split_ranges_set.end()) {
return;
}
if constexpr (has_on_delete) {
if (it->second == 0) {
on_delete(it->first.lower(), it->first.upper());
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;
}
}
m_split_ranges_set.erase(it++);
}
} while (any_removals);
}
template <typename Func>
@ -220,18 +223,6 @@ 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 = static_cast<u64>(info->ExceptionRecord->ExceptionInformation[1]);
u64 exception_addr = reinterpret_cast<u64>(info->ExceptionRecord->ExceptionInformation[1]);
if (code != EXCEPTION_ACCESS_VIOLATION || info->ExceptionRecord->ExceptionInformation[0] == 1) {
// Not our problem

View file

@ -1,98 +0,0 @@
// 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,7 +457,6 @@ 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
@ -527,7 +526,6 @@ 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
@ -867,7 +865,6 @@ 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
@ -1036,8 +1033,6 @@ 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,11 +4,6 @@
// 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"
@ -29,24 +24,13 @@ 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): 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;
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);
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);
@ -170,26 +154,7 @@ u64 DynarmicCallbacks64::GetTicksRemaining() {
}
u64 DynarmicCallbacks64::GetCNTPCT() {
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;
return m_parent.m_system.CoreTiming().GetClockTicks();
}
bool DynarmicCallbacks64::CheckMemoryAccess(u64 addr, u64 size, Kernel::DebugWatchpointType type) {

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