Donkey Kong Bananza

A Switch 2 game: update v458752 over the base game, SDK 20.5.19 built for the Switch 2 (sdk names its spec Ounce), NVN2. Nintendo EPD's engine (sead / agl / gsys, Phive), modules main + sdk. Plays on the Mac (M4 Pro): three play sessions of about 5 minutes each and scripted runs of up to 15,900 frames, through the opening, the mining tunnel and the areas after it, with the guide movies in their panel, at 60 fps; in busy scenes the game drops itself to 30 fps for a few seconds (Open). The first start with no pipeline archive builds the pipelines for about 4 minutes, once. Nothing past the first level's opening areas is tested; iOS is untested.

What it needs

Files

• Built from a dump folder (base/, updates/458752/, dlc/040074a01bf13001/): build_ipa_from_nsp.py <folder> --platform mac-app packs it with scripts/pack_dump.py, with no keys; the app has no icon, as a dump has no control data. The update's RomFS is the patched one (80,273 files, 8.6 GB; it drops 42 base files, among them the Product.100 tables its Product.200 ones replace), so it is packed on its own. The DLC (add-on content index 1, 8,000 files, 0.9 GB) mounts as EmeraldRush.
• Shaders: Shader/*.Nin_Ounce_NVN2.bfsha.zs (BNSH) and Shader/Project.Nin_Ounce_NVN2.release.sarc.zs. The code is Volta / Ampere's 128-bit encoding (0x794d EXIT, 0x7918 NOP), not Maxwell's.
• NVN2 control sections (dumped with DM_NVN_SHDUMP): magic 0x98761234, version 0xc9; +0x38 the instructions' length (SPH included), +0x40 that rounded to 0x100, +0x44 the code area's (the c1 segment lies between the two), +0x4c the register count, +0x58 compute's thread group. Their own length is at +0x30 in BNSH files only (0xd8; 0xd0 in effects, 0 in agl's BAHS and in memory), so the shader key hashes the first 0x48 bytes with +0x30 zeroed, plus the code area (hle_nvn.c nvn_ctl_lens, shader_split.py key_ctl).
• agl's upscaler is an NDSR pack, agl_ndsr.pck in Lib/agl/agl_resource.Nin_Ounce_NVN2.release.sarc.zs: 16-byte records from +0x10 (control section offset, control length << 4 | type, code offset, code length / 16 << 8 | type) up to a zero one; 30 shaders (3 vertex, 3 pixel, 24 compute: a small network on the tensor cores, HMMA, LDSM, MOVM). shader_split.py ndsr_shaders reads it, and pairs each vertex record with the pixel record after it.
• The split reads the bfsha / BNSH files, agl's BAHS archives (agl_common.sharcb and the others: 1,555 shader records, one more than the archive counts, so the record check allows 1 % of strays) and sead's primitive drawer: 22,852 shaders (with the 30 of the NDSR pack).
• geo_mesh.tsv: the game's one vertex / geometry pair (f9c807930abd1234 with the passthrough kernel 5d537f7630ec3671), which draws every geometry draw, about 600 a frame; built as a mesh library of 16 triangles a group (Mesh geometry groups sized to the vertex output).
• The prewarm seed list, pso_manifest.tsv (3,793 lines: 1,726 shaders, 2,030 render pipelines, 37 compute pipelines), is the Mac's device list of the play sessions before the host read device lists without a seed list (development.md): boot, the title, the opening and its guide movies, the mining tunnel and the areas after it, recorded with prewarm off. A start with no device list builds it in about 8 s before the first frame.

Shaders

The translator decodes Ampere code natively (job arch: sm86; Ampere shader code): all 22,852 shaders translate (chain 6f41726ef36d). NVN2 keeps NVN's texture handle table in c[2] (0x20 + 8 slot) and the graphics stages' storage buffer tables in c[0] (0x110 vertex, 0x510 pixel); compute lays out c[0] per shader from its control section, push constants first (Ampere shader code). The upscaler pushes (1 / w, 1 / h, w, h) of its target for its graphics passes and the source and destination rectangles for its compute network.

Rendering

The scene renders into 1676×943 targets (R11G11B10F colour), agl's upscaler takes it to 2304×1296, and that is copied to the 2880×1620 window. On the way to the first picture:

1. The upscaler's network did not translate (HMMA and the other tensor instructions): the window kept the output's clear colour (0.3, 0.3, 1).
2. Six effect compute shaders (d076142721d836a5 and others) did not load: a storage address was a c[0] entry plus an offset from c[1], and the translator took the c[1] term as the base.
3. Black frames: the composite was cleared before the upscaler read it (render_cond, What it needs).
4. Black speckles and white blotches over a dark frame: NaN from the environment cube's mip chain (Mip-chain pass sampling the level it renders). Found with raw dumps (DM_NVN_SNAPRAW) of the R11G11B10F targets, DM_NVN_HILITE to narrow the shaders (cca896efc6bb7c21 and c8637178e2818d18 wrote most of them), then those shaders' MSL instrumented to record their first NaN temporary: a fetch of the prefiltered cube, back through the prefilter to the cube's own mip passes.
5. Woven rock surfaces and holes in the rock: the textures the game loads were untiled in the Tegra X1's order inside a GOB (The Switch 2's order inside a GOB). The rock material's albedo (ASTC 8×8) written straight to its output looked right and its BC7 normal map did not, so the BC7 file (BaseRock000_Pnrm) was decoded offline in every order of the GOB's address bits.
6. The geometry shader 5d537f7630ec3671 (no vertex output: it writes its primitives to memory with ATOMG / STG) needed LEPC, OUT and ISBERD, and 7529699c7b7e49c0 (ActionLibraryPass compute) MATCH; with them all 468 shaders of the first gameplay minute translate.
7. Black tunnel walls and floor, with only the loose rocks and Donkey Kong drawn (compared with a YouTube capture of the same spot): the terrain's colour pass, about 400 indirect draws, drew nothing.
   • How the terrain is drawn: each chunk's depth prepass runs the culling geometry shader 5d537f7630ec3671. It writes the surviving triangles' indices and adds their count to the chunk's indirect draw record (64 bytes, bound as its storage buffer 0). The colour pass draws from that record.
   • What was wrong: the records held instance count 0, and their counts grew frame after frame. Every frame gsys's vertex shader c9a56f796ea4a4cd (96 points or more) writes {count 0, one instance} into each record whose address its storage buffer lists. The translator had dropped those stores (Global memory at an address loaded from memory).
   • What NVN2 does: decoding NVN2's indirect draw macro showed that instance count 0 draws nothing and the draw never resets the record.
   • Now: the walls, floor and lit passage draw as in the capture, at 60 fps.

Boot (Mac)

Each step below stopped the boot before its fix:

1. sdk: no SDK version string: the version reads nnSdk-Generic_20_5_19 (loader.c dm_sdk_major skips the variant name).
2. GetInfo 0x21, a Switch 2 info type: the kernel's page size. nn::ro::ProtectRelro takes value − 1 as the RELRO alignment mask; 0 aborted (hle_svc.c answers 0x1000). Type 28 is Horizon's AliasRegionExtraSize (0, named with 25–27 so that Switch 1 games print no new red line); 29–32 are the Switch 2's remote region: nnosInitialize builds nn::os::detail::RemoteRegionManager's address space allocator over [type 29, type 29 + type 30), and nn::os::GetRemoteMemoryUsage returns types 31 and 32. Only the GPU driver's boot maps into it (nv::InitializeGraphics → nn::gs::sdk::BootGraphicsSystem → BootWithRemoteMemory), and the host replaces nv::InitializeGraphics (the overload with a pool, the one the game calls), so all four are 0: no remote region.
3. NVN2 reports API 231: the game requires major 231, minor ≥ 30 (main+0x5c61a4) and creates no graphics system otherwise (null read in its frame loop). dm_sdk_ounce (the sdk's spec: Ounce) selects 231.30 (hle_nvn.c).
4. 197 NVN2 entry points (ray tracing, mesh shaders, DLSS, device commands, HDR window metadata, a ProgramBuilder) were unknown names, so lookups returned null; they are in names.txt now, the unimplemented ones going to the red stub. Implemented: the ProgramBuilder (nvnProgramInitializeFromBuilder = Initialize + SetShaders), QueueBuilderSetPrimaryQueue, CommandBufferSetRecordingMode / SetGeometryProcessingMode / SetComputeOverlapMode.
5. Switch 2 SDK calls through services the host does not serve: hid::InitializeJcms and the Joy-Con mouse sensor (hle_hid.c), oe::GetCurrentGraphicsScale (hle_oe.c), audio::IsDeviceOutputVolumeAutoTuneModeEnabled (hle_audren.c), streamplay (GameShare, hle_misc.c), vi HDR / VRR / extended swap interval (dm/gfx/hle_vi.c).
6. nv::InitCheck (from main+0xe19d04, the start of a video decoder, after the first punch) tested the SDK's own flag, which nv::InitializeGraphics sets and the host's version of it does not: hle_nv.c answers it from the host's record. The same function then starts a video decoder: the Switch 2's nn::codec video API and sync point events, served on the host (nn::codec video decoders taken over).

Open

• Video decoding: the guide movies are VP9 and play (a 180-picture clip on loop in the tutorial's panel, read back as NV12). The player also imports the VP8, H.264 and AV1 decoders: VP8 and H.264 are served too (checked by t_codec_video only), AV1 is not (nn::codec video decoders taken over).
• The gameplay picture: the mining tunnel matches a YouTube capture of the same spot (Rendering 7); nothing past the tunnel's first minute has been compared.
• Full shader build (about 75 minutes on the Mac from an empty cache): 22,852 of 22,852 shaders on chain 6f41726ef36d (SUATOM decoded; 22,850 on chain 3c66991eb439). Chain 9a0726336e9d had 22,834: the Metal compiler refused 11 (too deeply nested jump tables) and 5 used BRXU or REDUX (Ampere shader code); 2 still do not translate (0x593).
• oe::GetCurrentGraphicsScale: which tier the console reports in which mode is not measured (the host answers 0 handheld, 1 console). The game uses it as a quality tier, not for correctness: a byte of its graphics settings and a renderer parameter of 0, 9 or 13 (hle_oe.c).
• The 9 GB pool.
• The terrain's stores through loaded pointers keep reaching memory pools the host learns only afterwards (12 in a 17,700-frame session); each such miss is checked against the uniform shadows that relied on it, and none so far covered a written page (Global memory at an address loaded from memory, Open).
• Frame rate in busy scenes: the game picks its own present interval: 3 (20 fps) in its cutscenes, and 1 ↔ 2 (60 ↔ 30 fps) by its frame-rate controller (main+0x204a18; settings at 0x703f72b200: interval 1 to 3, margin -0.2, 2 late frames to go up, 6 to come back). It measures each frame from its start (+0xf0) to the later of its submission and the GPU's end (+0x128, ticks) and wants interval 2 from 1.2 vsyncs (20 ms). On the Mac (M4 Pro) that measure sits at 17 ms a frame (p90 24): the host's translation (CPU, 5-9 ms) before the GPU's 10-13 ms. The upscaler's HMMA through simdgroup_matrix (Ampere HMMA through simdgroup_matrix) cut the switches by about a fifth; the geometry runs' dependency check, quadratic in the run's length (Geometry dependency checks skip read-only pairs), was about 8 % of the submitting thread, and making it linear cut them from 30 to 12 on the same scripted route. Entering a level stalled for 1.3 s on GPU drains, one or two per geometry draw (CPU reads after a drain wait only for later writes); one CPU-bound build frame of 195-280 ms remains. Measured without sample attached: it suspends the process and makes the submitting thread look stalled. Its geometry draws are mesh draws (Mesh geometry groups sized to the vertex output): translation 9.9 → 7.2 ms a frame, GPU p50 11.8 → 9.0 ms on the scripted route. Open: the rest of the translation time in the busiest views (argument encoding, render pass setup), where translation and the GPU together still miss a vsync now and then.
• The NVN2 device values the driver reads from the GPU at run time (17 items: scratch scale factors, L2 size, texture layers and sizes, linear stride alignment, pool page size, …) are still the Switch 1 ones; the constant ones are NVN2's (The Switch 2's device values).
• iOS untested.

Found on this game

• Everything under Boot is new: the first Switch 2 game here.
• Ampere shader code (Switch 2): the translator's Ampere frontend, its scoreboard model and the IPA / texture / load semantics found on this game's menu (green boxes instead of text; no text); the tensor instructions, NVN2 compute's c[0] layout and push constants, MUFU SIN / COS in turns, and the storage base rule found in its gameplay.
• The Switch 2's order inside a GOB: the textures from its files.
• nn::codec video decoders taken over: the guide movies.
• The Switch 2's system tick is 31.25 MHz: the frame-rate controller held 20 fps.
• Ampere HMMA through simdgroup_matrix: agl's upscaler network, 2-6.5 ms a frame emulated with shuffles, 1.1-1.6 ms on the matrix hardware.
• The Switch 2's device values: NVN2's own nvnDeviceGetInteger answers, read from the game's sdk.
• The device's prewarm list unread without a seed list: compile stalls every session until the device list was read.
• Global memory at an address loaded from memory: the terrain's indirect draw records, reset by a gsys vertex shader's stores through loaded pointers, and the host's learned write set for such stores.
• Mip-chain pass sampling the level it renders: the environment cube's NaN texels.
• Geometry dependency checks skip read-only pairs: runs of 56 geometry draws, checked pairwise, took 8 % of the submitting thread.
• CPU reads after a drain wait only for later writes: the level opening drained the GPU for each geometry draw of an indirect-count multi-draw, a hitch of 1.3 s.
• Mesh geometry groups sized to the vertex output: 388-byte vertices, over Metal's mesh limit at 32 triangles a group.
• Ampere shader code (Switch 2): SUATOM (0x593) in two agl technique compute shaders, the last two of the package's 22,852 that did not translate.
• Shader libraries loaded when the program is made: a new area's 3 to 7 libraries loaded on the submit thread, up to 4 s with the drive slowed.