MacKwnoX/NX2navi
1
0
Fork 0
mirror of https://github.com/tangerineturtle1/NX2navi.git synced 2026-11-15 19:19:58 -04:00
Code Issues Packages Projects Releases Wiki Activity

Fetch the lazy write-back's ranges without the freeze lock, so a touch of GPU-written memory holds up only the threads that touch that range, with wb_lazy_unlocked = 0 to go back.

Browse Source 
A CPU touch of an armed range ran its fetch (vk_memory.c lazy_fetch: a GPU copy after everything handed over, and a wait of up to 10 s) inside the fault handler under frz_lock. Every other fault, thaw, freeze and arm waited for the whole copy meanwhile: the replay's copies of guest memory (mem_copy_readable, vk_memory.c mirror_capture) and its freeze pass, the completion thread's arms, file reads and pool finalizes. A write that lost the race to the fetch went to thaw_fault after it and released every other armed range of its 2 MiB region, a second fetch. mirror_evict and hot_check fetched their victims' ranges under the exclusive pool_lock, which the completion thread's arms and every pool lookup wait for. This is AUDIT F-4-05 (C-12); the design's probe of the unchanged code had an unrelated mem_copy_readable, arm and freeze each wait 299 ms behind a 300 ms fetch, and the losing write 378.8 ms with two fetches.

A range now leaves the armed list into a table of ranges in flight (memory.c "Ranges in flight", lz_fly) in the same lz_lock hold, and stays there, its pages inaccessible, until its bytes are down and its pages read-write again (lazy_fly_run), without frz_lock. A fault on a range in flight waits for that entry alone on lz_lock's condition and runs the access again, the losing write included; a thaw, a mapping change, a release of its buffer, an arm, a freeze and mem_lazy_touches see it. Every store to lzr is made under lz_lock from the list, the entries and the region an arm is protecting (lz_arm_r, which replaces the previous commit's restore). An entry starts under frz_lock and lz_lock and ends under lz_lock alone, so a holder of frz_lock may wait for one: its fetch takes only the queue lock, which no frz_lock taker holds. mem_thaw releases what is armed in its regions before it takes frz_lock (lazy_release_regions), so file reads and pool finalizes fetch without it; a mapping change's thaw still fetches under it, as its page calls follow. The holds a fault sets (lz_hot, lz_hots) are now set as the range goes into flight, so an arm meanwhile finds them. The fetch takes several copies in one submission and one wait (MemLazyCopy), and the back end has four slots of command pool, command buffer and fence, so fetches of different ranges run at once; a release takes up to 16 armed ranges out together.

mirror_evict and hot_check now fetch what only their victims' mirrors hold before taking pool_lock exclusively, and then, under the lock, release only what was armed since (usually nothing) before the mirror leaves the list and the device address table names the imported buffer again. The design moved the release after the lock instead, which would have had the table change before the fetch: a shader of a submission still running reads through the table as it changes, and the fetch waits for every submission handed over. A victim that vk_pool_fini took off the list meanwhile is left to it; hot_check copies its log line's fields under the lock, as its records may be freed once the lock drops.

The new setting wb_lazy_unlocked (NX2NAVI_WB_LAZY_UNLOCKED, default 1, restart) set to 0 keeps frz_lock from the take-out to the end of each flight and leaves mem_thaw's releases to the thaw, as in 0.10.1; the table, the slots and the eviction order stay. The "mem: lazy write-back" line counts segments fetched; a second line gives the touches that waited for a range in flight, the longest such wait and the most ranges in flight at once; "mem: sizes" counts fetches that found every slot busy. The frz_lock and lz_lock comments say what holds what now (B-12). mem_lazy_touches has callers outside memory.c (camera_speed.c, nvn_api_cmd.c ctl_fingerprint), which now skip a range in flight instead of faulting on it and waiting.

Not reproduced in the game before the change. tests\lazy_fetch_check.py now runs 20 scenarios as their own processes: arm races, no stall of freezes, copies, arms, thaws and releases of other ranges during a fetch, joins with frz_lock held elsewhere, a losing write that keeps the region's other range armed, two fetches at once, a thaw, a mapping change and a release meeting a range in flight, an arm refused over one, the foreign give-back, the chunk mode, the old wait with wb_lazy_unlocked = 0 (joins and losing writes as well), the report lines, 80 touching threads past the 64 entries, and a freeze refused; all pass, in six runs. Taking the join before frz_lock out fails the join scenario. cfg_rows_check reports 0 problems; every runtime file compiles, and memory.c and vk_memory.c add no -Wall -Wextra warning; the test matrix passes 75 of 77, installer_test and worktree_test failing for the environment as before. vk_memory.c's slots and the eviction's order have no check without a device.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
This commit is contained in:
chrg6
2026-11-13 19:33:58 -05:00
parent f586c291fa
commit 2e077edc9d
8 changed files with 806 additions and 213 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
docs/building.md
View File

@@ -688,12 +688,18 @@ A change to where the lazy write-back is on by default (`runtime/wb_lazy_choice.
`vk_mem_lazy_decide`) passes `python tests\wb_lazy_check.py`, which builds the header alone and checks the choice and
its reason by platform, by GPU driver and by what the player set. No game files, no GPU.
A change to the lazy write-back's arms, faults and releases (`runtime/memory.c` "Lazy write-back") passes
`python tests\lazy_fetch_check.py`, which builds memory.c with the Windows fault handler and a fetch that only waits,
and runs each scenario as its own process, so an access violation fails it: a thread that reads a range over and over
while it is armed and given back again never reaches the crash handler, also when the arm's clean test takes 100 ms;
and the "mem: lazy write-back" line gives the longest fetch and the longest wait of a fault for the freeze lock since
the line before. Windows only; no game files, no GPU.
A change to the lazy write-back's arms, faults, releases or fetches (`runtime/memory.c` "Lazy write-back" and "Ranges
in flight", `runtime/vk_memory.c` `lazy_fetch`, `mirror_evict`, `hot_check`) passes `python tests\lazy_fetch_check.py`,
which builds memory.c with the Windows fault handler and a fetch that only waits, and runs each scenario as its own
process, so an access violation fails it. A thread that reads a range over and over while it is armed and given back
again never reaches the crash handler, also when the arm's clean test takes 100 ms. A fetch holds up nothing but the
touches of its range: freezes, thaws, arms, releases and copies of other ranges return at once meanwhile, two ranges
fetch at once, and more touching threads than the flight table holds all end. The touches of a range in flight wait
for it and run again, a write among them without releasing the region's other ranges; thaws, mapping changes,
releases of its buffer, arms and freezes see it. The same with `wb_lazy_chunk = 1`, and with `wb_lazy_unlocked = 0`
the fetch holds the freeze lock again. The "mem: lazy write-back" lines give the longest fetch, the longest wait of a
fault for the freeze lock and the touches that waited for a range in flight. `vk_memory.c`'s fetch slots and the
order of an eviction against a pool's finalize have no check without a device. Windows only; no game files, no GPU.
A change to when the freezing of idle guest regions waits (`runtime/freeze_hold.h`, `runtime/vk_memory.c`
`freeze_pass`) passes `python tests\freeze_hold_check.py`, which builds the header alone: the backoff after a slow

6
nx2navi.ini.example
View File

@@ -1184,6 +1184,12 @@
# at the start of buffers of several MiB about twice a second. 0: the whole buffer comes back and everything near it
# is copied back every frame for up to 32 seconds, as before (a tenth fewer frames a second in a village).
# wb_lazy_chunk = 0
# With wb_lazy: 1 (default): when the game's CPU code touches memory the video card wrote, the copy that brings it
# back holds up only the threads that touch that memory; the game's other threads, the video card's bookkeeping and
# file reads go on meanwhile, and several such copies can run at once. 0: every such copy holds them all up until it
# is done, as in 0.10.1. Set 0 if the game ends or the graphics device is lost where it did not before, and send the
# log.
# wb_lazy_unlocked = 1
# Barriers between the game's draw passes and compute dispatches. 1: only where the game itself asks for one, or where
# the work changes between drawing and compute; this takes about a tenth off the graphics card's time for a frame.
# 0: a full barrier before every pass and every dispatch. Default: 1 on NVIDIA graphics cards, 0 on all others, where it

6
runtime/cfg.c
View File

@@ -77,7 +77,7 @@ static const char *const known[] = {
"SWKBD_ANSWER",
"SET_AT", "SWKBD_TEST", "TAA", "DBG_TEST", "DEBUG_MENU", "MENU_TEST_DEBUG", "FSR4", "FSR4_FG", "FG4_VERSION", "FG4_SWAPCHAIN", "TEST_PAD", "TEST_PAD_AT", "TEST_PREPARE_SHADERS", "TEX_ALIAS", "DEPTH_INIT_INLINE", "DEPTH_ALIAS_LOG", "TEX_DISUSE", "TEX_DISUSE_EVICT", "TEX_GEN_UNDO", "TEX_MEMO", "TEX_VRAM_PRESSURE", "TEX_RELEASED_MB", "TEX_RELEASED_MAX", "TEX_RELEASED_HOLD", "TEX_REUSE", "TEX_SPLIT_FILL", "TEX_COMMIT_FREE_MB", "TEX_PROC_COMMIT_MB", "TEX_EVICT", "TEX_COMMIT_EVICT", "MIRROR_EVICT", "VRAM_BUDGET_MB", "TIMER_1MS", "TRACE", "TXD_TRANSPOSE", "UBO_SNAPSHOT", "EVO_TRACE", "EVENT_AT_SYNC", "UPSCALER", "VK_DYN_A2C", "VK_DYN_VERTEX", "VK_NAN_PRESERVE", "VK_NOOPT", "VK_NO_FLOAT64",
"VK_PRESENT", "VK_RECONVERGE", "VK_NAN_BITS", "VK_RECORDER", "VK_SSBO_ALIGN", "VK_SUBGROUP", "VK_SUBGROUP_PIN", "VK_VALIDATION", "VOLUME", "VRAM_HEADROOM", "VS_OUTS", "WATER_PROJECTIVE",
"WAIT_WATCH_S", "WALK_KEY", "WALK_SCALE", "WB_EXACT", "WB_FENCE_LIMIT", "WB_MERGE", "WB_NO_REUP", "WB_MERGE_MAX", "WB_PROBE", "WB_SPLIT", "WB_LAZY", "WB_LAZY_CHUNK", "WB_LAZY_TOUCH", "WINDOW_SIZE", "XLATE", "XLATE_ID", "XLATE_KEEP", "XLATE_PROCS",
"WAIT_WATCH_S", "WALK_KEY", "WALK_SCALE", "WB_EXACT", "WB_FENCE_LIMIT", "WB_MERGE", "WB_NO_REUP", "WB_MERGE_MAX", "WB_PROBE", "WB_SPLIT", "WB_LAZY", "WB_LAZY_CHUNK", "WB_LAZY_TOUCH", "WB_LAZY_UNLOCKED", "WINDOW_SIZE", "XLATE", "XLATE_ID", "XLATE_KEEP", "XLATE_PROCS",
"UPLOAD_BATCH_MB", "VRAM_HEAP_LIMIT", "VRAM_RESERVE_MB", "TEX_CLOBBER", "LOW_END_GPU",
"ZSTD_CHECK", "ZSTD_NATIVE",
"BIND_REUSE", "DYN_DELTA", "STG_ARENA", "SUBMIT_BATCH", "SUBMIT_THREAD", "ASTC_GPU_DESWIZZLE", "ASTC_SCRATCH_KEEP", "ASTC_ARENA", "ASTC_BC1",
@@ -172,7 +172,7 @@ static const struct { const char *key, *section; } sections[] = {
{"VK_NAN_PRESERVE", "performance"}, {"VK_NAN_BITS", "performance"},
{"DRAW_CRUMBS", "diagnostics"}, {"TRACE", "diagnostics"}, {"FLIGHT", "diagnostics"}, {"PASS_CRUMBS", "diagnostics"}, {"FLIGHT_SPIKE_MS", "diagnostics"}, {"STUTTER_MAP", "diagnostics"}, {"PAR_JOBS", "performance"}, {"PIPE_RELEASE", "performance"},
{"FLIGHT_FLASH", "diagnostics"}, {"FLIGHT_FLASH_DIFF", "diagnostics"}, {"FLIGHT_FLASH_CELLS", "diagnostics"},
{"WB_FENCE_LIMIT", "diagnostics"}, {"WB_MERGE", "diagnostics"}, {"WB_EXACT", "diagnostics"}, {"MEM_CENSUS", "diagnostics"}, {"WB_NO_REUP", "diagnostics"}, {"WB_PROBE", "diagnostics"}, {"WB_SPLIT", "performance"}, {"SHADOW_BULK_WRITE", "performance"}, {"WB_LAZY", "performance"}, {"HEAP_LAZY", "performance"}, {"BARRIER_GAME", "performance"}, {"MESH_AHEAD", "performance"}, {"MESH_FALLBACK", "performance"}, {"MESH_FALLBACK_MB", "performance"}, {"MESH_FALLBACK_WORKERS", "performance"}, {"WB_LAZY_TOUCH", "diagnostics"}, {"WB_LAZY_CHUNK", "performance"}, {"FLUSH_UPLOAD", "diagnostics"}, {"GPA_TWIN_MAX_KB", "diagnostics"}, {"MIRROR_CAPTURE_KB", "diagnostics"}, {"WB_SKIP_MB", "diagnostics"},
{"WB_FENCE_LIMIT", "diagnostics"}, {"WB_MERGE", "diagnostics"}, {"WB_EXACT", "diagnostics"}, {"MEM_CENSUS", "diagnostics"}, {"WB_NO_REUP", "diagnostics"}, {"WB_PROBE", "diagnostics"}, {"WB_SPLIT", "performance"}, {"SHADOW_BULK_WRITE", "performance"}, {"WB_LAZY", "performance"}, {"HEAP_LAZY", "performance"}, {"BARRIER_GAME", "performance"}, {"MESH_AHEAD", "performance"}, {"MESH_FALLBACK", "performance"}, {"MESH_FALLBACK_MB", "performance"}, {"MESH_FALLBACK_WORKERS", "performance"}, {"WB_LAZY_TOUCH", "diagnostics"}, {"WB_LAZY_UNLOCKED", "diagnostics"}, {"WB_LAZY_CHUNK", "performance"}, {"FLUSH_UPLOAD", "diagnostics"}, {"GPA_TWIN_MAX_KB", "diagnostics"}, {"MIRROR_CAPTURE_KB", "diagnostics"}, {"WB_SKIP_MB", "diagnostics"},
{"WB_MERGE_MAX", "diagnostics"}, {"MV_SURVEY", "diagnostics"}, {"TEXUP_LOG_MS", "diagnostics"},
{"WAIT_WATCH_S", "diagnostics"},
{"HANDLE_CENSUS", "diagnostics"},
@@ -433,7 +433,7 @@ static const CfgMeta metas[] = {
{"PIPE_RELEASE", "1", "0|1", RESTART},
// flight.c "The flash trigger": on or off, and flight_flash.h flash_test's two thresholds
{"FLIGHT_FLASH", "1", "0|1", RESTART}, {"FLIGHT_FLASH_DIFF", "24", "4..255", RESTART}, {"FLIGHT_FLASH_CELLS", "3", "1..144", RESTART},
{"WB_FENCE_LIMIT", "65536", "0..1073741824", RESTART}, {"WB_MERGE", "1", "0|1", RESTART}, {"WB_EXACT", "1", "0|1", RESTART}, {"MEM_CENSUS", "0", "0|1", RESTART}, {"WB_NO_REUP", "1", "0|1", RESTART}, {"WB_PROBE", "0", "0|1", RESTART}, {"WB_SPLIT", "1", "0|1|2|3", RESTART}, {"WB_LAZY", WB_LAZY_DEFAULT, "0|1", RESTART}, {"HEAP_LAZY", "0", "0|1", RESTART}, {"WB_LAZY_TOUCH", "0", "0..16", RESTART}, {"WB_LAZY_CHUNK", "0", "0|1", RESTART}, {"FLUSH_UPLOAD", "0", "0|1", RESTART}, {"GPA_TWIN_MAX_KB", "0", "0..1048576", RESTART},
{"WB_FENCE_LIMIT", "65536", "0..1073741824", RESTART}, {"WB_MERGE", "1", "0|1", RESTART}, {"WB_EXACT", "1", "0|1", RESTART}, {"MEM_CENSUS", "0", "0|1", RESTART}, {"WB_NO_REUP", "1", "0|1", RESTART}, {"WB_PROBE", "0", "0|1", RESTART}, {"WB_SPLIT", "1", "0|1|2|3", RESTART}, {"WB_LAZY", WB_LAZY_DEFAULT, "0|1", RESTART}, {"HEAP_LAZY", "0", "0|1", RESTART}, {"WB_LAZY_TOUCH", "0", "0..16", RESTART}, {"WB_LAZY_CHUNK", "0", "0|1", RESTART}, {"WB_LAZY_UNLOCKED", "1", "0|1", RESTART}, {"FLUSH_UPLOAD", "0", "0|1", RESTART}, {"GPA_TWIN_MAX_KB", "0", "0..1048576", RESTART},
{"MIRROR_CAPTURE_KB", "0", "0..1048576", RESTART}, {"WB_SKIP_MB", "0", "0..4096", RESTART}, // vk_memory.c capture_cap, wb_skip
{"WB_MERGE_MAX", "32", "1..256", RESTART}, {"TEXUP_LOG_MS", "0", "0..60000", RESTART}, // vk_texture.c vk_texstat_frame
{"WAIT_WATCH_S", "0", "0..3600", RESTART}, // kernel.c wait_watch: a lock or thread-end wait this long is logged; 0 off

7
runtime/host.h
View File

@@ -285,8 +285,9 @@ void mem_freeze_stats(uint64_t *frozen, uint64_t *freezes, uint64_t *fault_thaws
uint64_t *hot_holds); // hot_holds: quick thaws that held their region thawed (memory.c "Hot regions")
// NX2NAVI_WB_PROBE=1 (memory.c, "Probed regions"): who reads what the GPU wrote back into guest memory.
// NX2NAVI_WB_LAZY (memory.c, "Lazy write-back"): pages the GPU wrote that come down only when the CPU touches them.
int mem_lazy_enable(void (*fetch)(uint64_t h0, uint64_t h1, uint64_t off, uint64_t size),
int (*clean)(uint64_t lo, uint64_t hi, void *ctx));
typedef struct { uint64_t h0, h1, off, size; } MemLazyCopy; // [off, off + size) of buffer h0 into buffer h1
// fetch: the n copies after all the GPU work handed over so far, and a wait for them; called on several threads at once
int mem_lazy_enable(void (*fetch)(const MemLazyCopy *c, int n), int (*clean)(uint64_t lo, uint64_t hi, void *ctx));
uint64_t mem_lazy_arm(uint64_t a, uint64_t n, uint64_t h0, uint64_t h1, uint64_t off, void *ctx);
int mem_lazy_skip(uint64_t a, uint64_t n, uint64_t h0, int64_t delta, uint64_t out[][2], int max);
int mem_lazy_armed(uint64_t a, uint64_t n);
@@ -294,7 +295,7 @@ int mem_lazy_armed(uint64_t a, uint64_t n);
int mem_copy_readable(uint8_t *const *dst, const uint64_t (*src)[2], int n, uint8_t *ok);
int mem_lazy_touches(uint64_t a, uint64_t n); // 0: the CPU can read [a, a + n) without a fault now
void mem_lazy_window(int open);
void mem_lazy_release(uint64_t a, uint64_t n);
void mem_lazy_release(const uint64_t (*rg)[3], int n); // rg[i] = {lo, hi, buffer or 0 for any}: armed parts come down
void mem_lazy_report(uint64_t frame, uint64_t skipped);
void mem_lazy_note(char *out, size_t n); // its state as a phrase, without a lock (a fatal error's line)
int mem_lazy_pick(uint64_t *addr, int max); // test: addresses of armed pages

460
runtime/memory.c
View File

@@ -70,9 +70,10 @@ static int reserve_at(uint64_t at, uint64_t size, int watch) { return plat_mem_r
// read) fails on a read-only page instead of faulting, so file reads thaw their buffer first (mem_thaw), and so does
// every commit, decommit and protection change here. A region stays thawed for FRZ_HOLD_MS after a mem_thaw: a file
// read still in flight writes its buffer later. Nothing writes guest memory under frz_lock. It is not a leaf: lz_lock
// and vm_lock come inside it, a strict thaw's page calls run under it (thaw_pages_vm), and a fetch of lazily written
// pages (lazy_gone_locked, lazy_piece_fetch_locked: vk_memory.c lazy_fetch) takes the queue lock and waits for the GPU
// under it. So no holder of the queue lock may touch guest memory or take frz_lock.
// and vm_lock come inside it, a strict thaw's page calls run under it (thaw_pages_vm), and a holder may wait for a
// range of the lazy write-back in flight, whose fetch (vk_memory.c lazy_fetch) takes the queue lock and waits for the
// GPU ("Ranges in flight"; in a mapping change's thaw, and with NX2NAVI_WB_LAZY_UNLOCKED=0, the fetch itself runs
// under frz_lock). So no holder of the queue lock may touch guest memory or take frz_lock.
#define FRZ_SHIFT 21
#define FRZ_SIZE (1ull << FRZ_SHIFT)
#define FRZ_HOLD_MS 5000u
@@ -182,7 +183,8 @@ static void unprobe_locked(uint64_t r) {
// every part (mem_lazy_armed, mem_lazy_window). Under frz_lock it waited for freezes, thaws and arms on other threads,
// whose page queries and protection changes take that lock for tens to hundreds of microseconds: 8% of the replay
// thread's samples standing in Hyrule Field on an RX 580 (0.7.1, lazy write-back on by default). A change of the list
// holds both locks (frz_lock first); a read holds either.
// holds both locks (frz_lock first); a read holds either. lz_lock also guards the ranges in flight (below); a wait for
// one waits on its condition, lz_fly_cv, which releases lz_lock meanwhile.
typedef struct { uint64_t lo, hi, h0, h1; int64_t delta; uint8_t stale; } LzSeg; // pages [lo, hi); buffer offset = address + delta
static PlatMutex lz_lock = PLAT_MUTEX_INIT; // inside frz_lock, never the other way round
static LzSeg *lzs; // changed under frz_lock and lz_lock, read under either
@@ -192,7 +194,7 @@ static uint32_t lz_hot[RT_ADDRESS_LIMIT >> FRZ_SHIFT]; // frz_lock: plat_
static uint8_t lz_back[RT_ADDRESS_LIMIT >> FRZ_SHIFT]; // frz_lock: seconds of the region's last hold
static volatile int32_t lz_on;
static int lz_want; // the setting, read in mem_init
static void (*lz_fetch)(uint64_t h0, uint64_t h1, uint64_t off, uint64_t size);
static void (*lz_fetch)(const MemLazyCopy *c, int n);
// The back end's test that the mirror holds every CPU write to [lo, hi) (vk_memory.c lazy_clean): a fetch copies the
// mirror's bytes over the whole range, so a CPU write the mirror does not have yet would be lost to it.
static int (*lz_clean)(uint64_t lo, uint64_t hi, void *ctx);
@@ -236,27 +238,145 @@ static uint64_t lz_foreign; // frz_lock, since the start: segments o
#define LZ_NOTES 8
static struct { uint64_t addr, lo, hi, out; uint32_t tid; uint8_t write, stale, hold; } lz_note[LZ_NOTES];
static int lz_nnote;
// frz_lock: since the last report: faults by a read, by a write, other releases, fetches, their bytes and time, the
// longest fetch, and the longest wait of a fault for frz_lock (behind other threads' fetches, thaws, freezes and arms)
static uint64_t lz_faults[2], lz_other, lz_fetches, lz_fetch_bytes, lz_fetch_ns, lz_fetch_max_ns, lz_lock_wait_max_ns;
static uint64_t lz_gaps;
// frz_lock: since the last report: faults by a read, by a write, other releases, faults inside a part's window, and
// the longest wait of a fault for frz_lock (behind other threads' thaws, freezes and arms, and their fetches where
// those run under it)
static uint64_t lz_faults[2], lz_other, lz_gaps, lz_lock_wait_max_ns;
// lz_lock: since the last report: segments fetched, their bytes, the time in fetches and the longest fetch (one call of
// lz_fetch, which may carry several segments); touches that waited for a range in flight and the longest such wait; the
// most ranges in flight at once
static uint64_t lz_fetches, lz_fetch_bytes, lz_fetch_ns, lz_fetch_max_ns, lz_joins, lz_join_max_ns;
static int lz_fly_most;
static volatile uint32_t lz_win_tid; // lz_lock: the replay thread's id while its part is between check and hand-over
// For a fatal error's summary line (mem_lazy_note): written under frz_lock and read without it, since the thread that
// fails may be the one inside a fetch.
static volatile uint64_t lz_fetch_total; // fetches since the start
static volatile uint32_t lz_fetch_ms; // plat_now_ms() | 1 at the end of the last one; 0: none yet
static volatile uint32_t lz_fetch_tid; // the thread inside a fetch now; 0: none
static volatile uint64_t lz_fetch_now; // that fetch's bytes
// For a fatal error's summary line (mem_lazy_note), read without a lock, since the thread that fails may be the one
// inside a fetch: the segments fetched since the start and the end of the last fetch. Those in progress are in lz_fly.
static volatile uint64_t lz_fetch_total; // segments fetched since the start
static volatile uint32_t lz_fetch_ms; // plat_now_ms() | 1 at the end of the last fetch; 0: none yet
// Ranges in flight. A segment leaves the armed list with its pages still inaccessible; then its bytes come down when
// the mirror's are newer (lz_fetch: a GPU copy after everything handed over, and a wait for it) and its pages get their
// access back (lazy_fly_run), without frz_lock. Under frz_lock a fetch held up every other fault, thaw, freeze and arm,
// the replay's copies of guest memory (mem_copy_readable) and the completion thread's arms for the length of the copy;
// and a write that lost the race to it released every other armed range of its 2 MiB region (thaw_fault). From the
// moment it leaves the list until its pages are read-write again the range is an entry of lz_fly:
// - a fault on it waits for that entry alone, on lz_fly_cv and never on frz_lock, and runs the access again
// (lazy_fly_join_locked);
// - a thaw or a release that reaches it waits for it (lazy_drop_region_locked, mem_lazy_release), so the kernel's
// access, a mapping change or the end of the buffer comes after its copy and its page calls;
// - an arm leaves its pages alone (mem_lazy_arm), a freeze its region (lzr), and mem_lazy_touches counts it.
// Every segment that leaves the list with inaccessible pages becomes an entry in the same lz_lock hold, and an entry
// ends only once its pages are read-write; lists and entries never overlap. So a fault that finds its address in
// neither, holding frz_lock (no arm runs then), has an address that is not ours or no longer inaccessible. An entry
// starts under frz_lock and lz_lock and ends under lz_lock alone, so a holder of frz_lock may wait for one: its fetch
// takes the queue lock, and no holder of the queue lock takes frz_lock. Every store to lzr is made under lz_lock, from
// the list, the entries and the region an arm is protecting (lz_arm_r). Fetches of different ranges run at once (the
// back end has a slot for each, vk_memory.c lazy_fetch). NX2NAVI_WB_LAZY_UNLOCKED=0 keeps frz_lock from the moment a
// fault or a release takes a range out until its entry ends, as before.
#define LZ_FLY 64 // more than the threads that touch armed pages at once; when full, a touch waits for room
#define LZ_BATCH 16 // segments a release takes out together and fetches in one call
typedef struct { uint64_t lo, hi, h0, bytes; uint32_t seq, tid; uint8_t used; } LzFly; // bytes: to fetch, 0 none
static LzFly lz_fly[LZ_FLY]; // lz_lock
static volatile int32_t lz_nfly; // lz_lock: entries in use (read without it as a hint)
static PlatCond lz_fly_cv = PLAT_COND_INIT; // with lz_lock: an entry ended
static int64_t lz_arm_r = -1; // lz_lock: the region whose pages mem_lazy_arm is protecting, -1 none
static int lz_unlocked = 1; // NX2NAVI_WB_LAZY_UNLOCKED, read in mem_init
// lz_lock held: lzr of the regions [lo, hi) reaches, from the armed segments, the ranges in flight and the arm in
// progress.
static void lazy_flags_locked(uint64_t lo, uint64_t hi) {
for (uint64_t r = lo >> FRZ_SHIFT; r <= (hi - 1) >> FRZ_SHIFT; r++) {
int any = 0;
for (int j = 0; j < lz_n && !any; j++)
any = lzs[j].lo < ((r + 1) << FRZ_SHIFT) && lzs[j].hi > (r << FRZ_SHIFT);
uint64_t b0 = r << FRZ_SHIFT, b1 = b0 + FRZ_SIZE;
int any = (int64_t)r == lz_arm_r;
for (int j = 0; j < lz_n && !any; j++) any = lzs[j].lo < b1 && lzs[j].hi > b0;
for (int k = 0; k < LZ_FLY && lz_nfly && !any; k++)
any = lz_fly[k].used && lz_fly[k].lo < b1 && lz_fly[k].hi > b0;
__atomic_store_n(&lzr[r], (uint8_t)any, __ATOMIC_RELEASE);
}
}
// lz_lock held: an entry in flight that meets [lo, hi), of buffer h0 (0: any), -1 none.
static int lazy_fly_find_locked(uint64_t lo, uint64_t hi, uint64_t h0) {
for (int k = 0; k < LZ_FLY && lz_nfly; k++)
if (lz_fly[k].used && lz_fly[k].lo < hi && lz_fly[k].hi > lo && (!h0 || lz_fly[k].h0 == h0)) return k;
return -1;
}
// lz_lock held (frz_lock too, or not): returns once entry k has ended; the time that took. The timeout only looks
// again; an entry ends on every path of lazy_fly_run.
static uint64_t lazy_fly_wait_entry_locked(int k) {
uint32_t seq = lz_fly[k].seq;
uint64_t t0 = plat_now_ns();
while (lz_fly[k].seq == seq) plat_cond_wait_ns(&lz_fly_cv, &lz_lock, 100000000ull);
return plat_now_ns() - t0;
}
// lz_lock held: a fault at a waits for the range in flight that holds it. 1 when there was one: the access runs again.
static int lazy_fly_join_locked(uint64_t a) {
int k = lazy_fly_find_locked(a, a + 1, 0);
if (k < 0) return 0;
uint64_t ns = lazy_fly_wait_entry_locked(k);
lz_joins++;
if (ns > lz_join_max_ns) lz_join_max_ns = ns;
return 1;
}
// lz_lock held: returns once no range in flight meets any of rg[j] = {lo, hi, buffer or 0 for any}.
static void lazy_fly_wait_ranges_locked(const uint64_t (*rg)[3], int nr) {
for (int j = 0; j < nr; j++)
for (int k; (k = lazy_fly_find_locked(rg[j][0], rg[j][1], rg[j][2])) >= 0;) lazy_fly_wait_entry_locked(k);
}
// lz_lock held: returns once an entry is free, having released lz_lock meanwhile if it waited.
static void lazy_fly_room_locked(void) {
while (lz_nfly >= LZ_FLY) plat_cond_wait_ns(&lz_fly_cv, &lz_lock, 100000000ull);
}
// frz_lock and lz_lock held, an entry free: s, out of the list in this hold, is in flight from here on.
static int lazy_fly_start_locked(LzSeg s) {
int k = 0;
while (lz_fly[k].used) k++;
LzFly *f = &lz_fly[k];
f->lo = s.lo, f->hi = s.hi, f->h0 = s.h0, f->bytes = s.stale ? s.hi - s.lo : 0, f->tid = plat_thread_id();
f->used = 1;
lz_nfly++;
if (lz_nfly > lz_fly_most) lz_fly_most = lz_nfly;
return k;
}
// No lock held (frz_lock in a mapping change's thaw, in mem_lazy_skip's give-back and with NX2NAVI_WB_LAZY_UNLOCKED=0):
// the n segments s[] in flight in the entries k[] get the mirror's bytes, the stale ones in one call of lz_fetch, then
// their pages' access, and leave flight. n is at most LZ_BATCH.
static void lazy_fly_run(const int *k, const LzSeg *s, int n) {
MemLazyCopy cp[LZ_BATCH];
int nc = 0;
uint64_t bytes = 0, ns = 0;
for (int i = 0; i < n; i++)
if (s[i].stale && lz_fetch) {
cp[nc++] = (MemLazyCopy){s[i].h0, s[i].h1, (uint64_t)((int64_t)s[i].lo + s[i].delta), s[i].hi - s[i].lo};
bytes += s[i].hi - s[i].lo;
}
if (nc) {
uint64_t t0 = plat_now_ns();
lz_fetch(cp, nc);
ns = plat_now_ns() - t0;
__atomic_store_n(&lz_fetch_ms, (uint32_t)plat_now_ms() | 1, __ATOMIC_RELAXED);
__atomic_add_fetch(&lz_fetch_total, (uint64_t)nc, __ATOMIC_RELAXED);
}
for (int i = 0; i < n; i++)
for (uint64_t a = s[i].lo, b; a < s[i].hi; a = b)
if (plat_mem_state(a, s[i].hi, &b) == PLAT_PROT_NONE && !plat_mem_protect(a, b - a, PLAT_PROT_RW))
rt_fatal("memory: giving back the lazily written pages %#llx+%#llx failed", (unsigned long long)a,
(unsigned long long)(b - a));
plat_mutex_lock(&lz_lock);
for (int i = 0; i < n; i++) lz_fly[k[i]].used = 0, lz_fly[k[i]].seq++;
lz_nfly -= n;
for (int i = 0; i < n; i++) lazy_flags_locked(s[i].lo, s[i].hi);
lz_fetches += (uint64_t)nc, lz_fetch_bytes += bytes, lz_fetch_ns += ns;
if (ns > lz_fetch_max_ns) lz_fetch_max_ns = ns;
plat_cond_broadcast(&lz_fly_cv);
plat_mutex_unlock(&lz_lock);
}
// frz_lock and lz_lock held: segment i leaves the list, with the stale mark mem_lazy_skip may have set until now. A
// later mem_lazy_skip no longer finds it and writes its range back as ever.
static LzSeg lazy_take_locked(int i) {
@@ -265,41 +385,12 @@ static LzSeg lazy_take_locked(int i) {
return s;
}
static void lazy_gone_locked(LzSeg s, int why);
// frz_lock held exclusively. Segment i goes: its bytes come down first when the mirror's are newer, then its pages
// are read-write again. why: 0 a read fault, 1 a write fault (both hold the regions back), else a thaw or the mirror
// going away.
static void lazy_drop_locked(int i, int why) {
plat_mutex_lock(&lz_lock);
LzSeg s = lazy_take_locked(i);
plat_mutex_unlock(&lz_lock);
lazy_gone_locked(s, why);
}
// frz_lock held exclusively: the pages of s, out of the list already, get their bytes and their access back.
static void lazy_gone_locked(LzSeg s, int why) {
if (s.stale && lz_fetch) {
uint64_t t0 = plat_now_ns();
__atomic_store_n(&lz_fetch_now, s.hi - s.lo, __ATOMIC_RELAXED);
__atomic_store_n(&lz_fetch_tid, plat_thread_id(), __ATOMIC_RELEASE);
lz_fetch(s.h0, s.h1, (uint64_t)((int64_t)s.lo + s.delta), s.hi - s.lo);
__atomic_store_n(&lz_fetch_ms, (uint32_t)plat_now_ms() | 1, __ATOMIC_RELAXED);
__atomic_store_n(&lz_fetch_total, lz_fetch_total + 1, __ATOMIC_RELAXED);
__atomic_store_n(&lz_fetch_tid, 0, __ATOMIC_RELEASE);
uint64_t ns = plat_now_ns() - t0;
lz_fetches++, lz_fetch_bytes += s.hi - s.lo, lz_fetch_ns += ns;
if (ns > lz_fetch_max_ns) lz_fetch_max_ns = ns;
}
for (uint64_t a = s.lo, b; a < s.hi; a = b)
if (plat_mem_state(a, s.hi, &b) == PLAT_PROT_NONE && !plat_mem_protect(a, b - a, PLAT_PROT_RW))
rt_fatal("memory: giving back the lazily written pages %#llx+%#llx failed", (unsigned long long)a,
(unsigned long long)(b - a));
lazy_flags_locked(s.lo, s.hi);
if (why > 1) { lz_other++; return; }
lz_faults[why]++;
// frz_lock held: a fault took [lo, hi) out whole, so its regions are held back from arms (lz_hot: 1 s, doubling up to
// 32 s while they are touched again within a minute of their last hold). Set as the range goes into flight, so an arm
// meanwhile finds the hold.
static void lazy_hold_locked(uint64_t lo, uint64_t hi) {
uint32_t now = (uint32_t)plat_now_ms();
for (uint64_t r = s.lo >> FRZ_SHIFT; r <= (s.hi - 1) >> FRZ_SHIFT; r++) {
for (uint64_t r = lo >> FRZ_SHIFT; r <= (hi - 1) >> FRZ_SHIFT; r++) {
int again = lz_hot[r] && (int32_t)(now - lz_hot[r]) < 60000; // touched again within a minute of its last hold
uint8_t b = again ? (uint8_t)(lz_back[r] < 32 ? lz_back[r] * 2 : 32) : 1;
lz_back[r] = b ? b : 1;
@@ -338,11 +429,10 @@ static void lazy_push_locked(LzSeg s) {
lzs[lz_n++] = s;
}
// A fault at a in segment i: the LZ_CHUNK piece around a goes (its bytes come down first when the mirror's are newer,
// then its pages are read-write again) and is held; what is left of the segment on either side stays armed. In two
// steps, as a whole segment goes (lazy_fault): the piece leaves the list under lz_lock (lazy_piece_take_locked), and
// its fetch and page calls run without that lock (lazy_piece_fetch_locked), which the replay asks at every draw that
// writes a mirrored pool (mem_lazy_skip) and at the end of every part.
// A fault at a in segment i: the LZ_CHUNK piece around a goes into flight and is held (lazy_piece_hold_locked); what
// is left of the segment on either side stays armed. Its fetch and page calls run in flight, as a whole segment's do,
// without lz_lock, which the replay asks at every draw that writes a mirrored pool (mem_lazy_skip) and at the end of
// every part.
typedef struct { LzSeg s; uint64_t c0, c1; } LzPiece; // the segment the piece came from, and the piece [c0, c1)
// frz_lock held exclusively and lz_lock held: the piece leaves the list, the segment's rest on either side goes back.
@@ -356,41 +446,57 @@ static LzPiece lazy_piece_take_locked(int i, uint64_t a) {
return (LzPiece){s, c0, c1};
}
// frz_lock held exclusively, lz_lock not: the piece's bytes and its access come back, and it is held. Returns its size.
static uint64_t lazy_piece_fetch_locked(LzPiece p, uint64_t a, int write) {
LzSeg s = p.s;
uint64_t c0 = p.c0, c1 = p.c1;
if (s.stale && lz_fetch) {
uint64_t t0 = plat_now_ns();
__atomic_store_n(&lz_fetch_now, c1 - c0, __ATOMIC_RELAXED);
__atomic_store_n(&lz_fetch_tid, plat_thread_id(), __ATOMIC_RELEASE);
lz_fetch(s.h0, s.h1, (uint64_t)((int64_t)c0 + s.delta), c1 - c0);
__atomic_store_n(&lz_fetch_ms, (uint32_t)plat_now_ms() | 1, __ATOMIC_RELAXED);
__atomic_store_n(&lz_fetch_total, lz_fetch_total + 1, __ATOMIC_RELAXED);
__atomic_store_n(&lz_fetch_tid, 0, __ATOMIC_RELEASE);
uint64_t ns = plat_now_ns() - t0;
lz_fetches++, lz_fetch_bytes += c1 - c0, lz_fetch_ns += ns;
if (ns > lz_fetch_max_ns) lz_fetch_max_ns = ns;
}
for (uint64_t x = c0, y; x < c1; x = y)
if (plat_mem_state(x, c1, &y) == PLAT_PROT_NONE && !plat_mem_protect(x, y - x, PLAT_PROT_RW))
rt_fatal("memory: giving back the lazily written pages %#llx+%#llx failed", (unsigned long long)x,
(unsigned long long)(y - x));
lazy_flags_locked(s.lo, s.hi);
lz_faults[write != 0]++, lz_pieces++, lz_piece_bytes += c1 - c0;
// frz_lock held: the piece around a is held back from arms (lz_hots, 1 s doubling up to 32 s like a region's hold), as
// it goes into flight.
static void lazy_piece_hold_locked(uint64_t a) {
uint32_t now = (uint32_t)plat_now_ms();
LzHot *e = lazy_hot_locked(a >> LZ_CHUNK_SHIFT, 1);
int again = e->until && (int32_t)(now - e->until) < 60000; // touched again within a minute of its last hold
e->back = again ? (uint8_t)(e->back < 32 ? e->back * 2 : 32) : 1;
if (!e->back) e->back = 1;
e->until = (now + (uint32_t)e->back * 1000u) | 1;
return c1 - c0;
}
// frz_lock held exclusively: every segment that reaches into region r goes.
static void lazy_drop_region_locked(uint64_t r, int why) {
for (int i = lz_n; i-- > 0;)
if (lzs[i].lo < ((r + 1) << FRZ_SHIFT) && lzs[i].hi > (r << FRZ_SHIFT)) lazy_drop_locked(i, why);
// frz_lock and lz_lock held: up to LZ_BATCH armed segments that meet any of rg[j] = {lo, hi, buffer or 0 for any} leave
// the list into flight, into k[] and s[]. Returns how many; -1 when some meet but no entry is free (wait for one with
// lazy_fly_room_locked, then again).
static int lazy_take_batch_locked(const uint64_t (*rg)[3], int nr, int *k, LzSeg *s) {
int n = 0;
for (int i = lz_n; i-- > 0 && n < LZ_BATCH;) {
int hit = 0;
for (int j = 0; j < nr && !hit; j++)
hit = lzs[i].lo < rg[j][1] && lzs[i].hi > rg[j][0] && (!rg[j][2] || lzs[i].h0 == rg[j][2]);
if (!hit) continue;
if (lz_nfly >= LZ_FLY) return n ? n : -1;
s[n] = lazy_take_locked(i);
k[n] = lazy_fly_start_locked(s[n]);
n++;
}
return n;
}
// frz_lock held exclusively (a thaw: a mapping change, a commit of the lazy heap, a write fault in a frozen region):
// every segment that reaches into region r goes, its fetch under frz_lock, since the caller's page calls follow under
// it; then the ranges other threads have in flight there end first. Not a fault on an armed page: no hold. mem_thaw
// releases its regions before it takes the lock (lazy_release_regions), so what is left here is rare.
static void lazy_drop_region_locked(uint64_t r) {
uint64_t rg[1][3] = {{r << FRZ_SHIFT, (r + 1) << FRZ_SHIFT, 0}};
for (;;) {
int k[LZ_BATCH];
LzSeg s[LZ_BATCH];
plat_mutex_lock(&lz_lock);
int n = lazy_take_batch_locked((const uint64_t (*)[3])rg, 1, k, s);
if (n <= 0) {
if (n < 0) lazy_fly_room_locked();
else lazy_fly_wait_ranges_locked((const uint64_t (*)[3])rg, 1);
plat_mutex_unlock(&lz_lock);
if (n < 0) continue;
return;
}
plat_mutex_unlock(&lz_lock);
lz_other += (uint64_t)n;
lazy_fly_run(k, s, n);
}
}
static int lazy_covered_locked(uint64_t lo, uint64_t hi) {
@@ -474,14 +580,22 @@ uint64_t mem_lazy_arm(uint64_t a, uint64_t n, uint64_t h0, uint64_t h1, uint64_t
// The region is marked before its pages become inaccessible. A fault on them from the protect on then finds
// the mark, waits in lazy_fault for frz_lock, held here, and finds the segment or its pages given back.
// Marked only after the push, a read in between had lazy_fault return 0 at once, and access_fault passed
// the read of an inaccessible page to the crash handler. had: a segment of the region, or an earlier turn
// of this loop, marked it already; when the arm is refused the mark goes back to what it was (every store
// to lzr is made under frz_lock).
int had = lzr[r];
if (!had) __atomic_store_n(&lzr[r], 1, __ATOMIC_RELEASE);
if (!lazy_protect_locked(x, e, ctx)) {
// the read of an inaccessible page to the crash handler. lz_arm_r keeps the mark while a range in flight
// in the region ends meanwhile (lazy_flags_locked); when the arm is refused the mark is made again from
// what is there. Pages of a range in flight are inaccessible until it ends and are no segment's: such a
// range is left alone.
plat_mutex_lock(&lz_lock);
int busy = lazy_fly_find_locked(x, e, 0) >= 0;
if (!busy) lz_arm_r = (int64_t)r, __atomic_store_n(&lzr[r], 1, __ATOMIC_RELEASE);
plat_mutex_unlock(&lz_lock);
if (busy || !lazy_protect_locked(x, e, ctx)) {
no = LZ_NO_PAGES;
if (!had) __atomic_store_n(&lzr[r], 0, __ATOMIC_RELEASE);
if (!busy) {
plat_mutex_lock(&lz_lock);
lz_arm_r = -1;
lazy_flags_locked(x, e);
plat_mutex_unlock(&lz_lock);
}
}
}
if (no >= 0) {
@@ -500,6 +614,7 @@ uint64_t mem_lazy_arm(uint64_t a, uint64_t n, uint64_t h0, uint64_t h1, uint64_t
j = -1;
}
lazy_push_locked((LzSeg){slo, shi, h0, h1, delta, stale});
lz_arm_r = -1; // the mark stays: the segment is in the list
plat_mutex_unlock(&lz_lock);
armed += e - x;
}
@@ -518,8 +633,9 @@ int mem_lazy_skip(uint64_t a, uint64_t n, uint64_t h0, int64_t delta, uint64_t o
int any = 0, k = 0;
for (uint64_t r = a >> FRZ_SHIFT; r <= (a + n - 1) >> FRZ_SHIFT && !any; r++) any = lzr[r];
if (!any) return 0;
// A segment armed for another buffer goes back under frz_lock (lazy_gone_locked), with arms held off until the
// rest is marked, as before lz_lock; that is rare, and the replay's per-draw check holds lz_lock only.
// A segment armed for another buffer goes back under frz_lock, through flight without a fetch (stale 0: its pages
// only), with arms held off until the rest is marked, as before lz_lock; that is rare, and the replay's per-draw
// check holds lz_lock only.
int foreign = 0;
plat_mutex_lock(&lz_lock);
for (int i = 0; i < lz_n && !foreign; i++)
@@ -529,17 +645,20 @@ int mem_lazy_skip(uint64_t a, uint64_t n, uint64_t h0, int64_t delta, uint64_t o
plat_mutex_lock(&frz_lock);
for (;;) {
plat_mutex_lock(&lz_lock);
int i = lz_n;
lazy_fly_room_locked();
int i = lz_n, f = -1;
while (i-- > 0 && !(lzs[i].lo < a + n && lzs[i].hi > a && (lzs[i].h0 != h0 || lzs[i].delta != delta))) {}
LzSeg s;
if (i >= 0) {
lzs[i].stale = 0;
lz_foreign++;
s = lazy_take_locked(i);
f = lazy_fly_start_locked(s);
}
plat_mutex_unlock(&lz_lock);
if (i < 0) break;
lazy_gone_locked(s, 4);
lz_other++;
lazy_fly_run(&f, &s, 1);
}
plat_mutex_lock(&lz_lock);
}
@@ -556,8 +675,8 @@ int mem_lazy_skip(uint64_t a, uint64_t n, uint64_t h0, int64_t delta, uint64_t o
return k;
}
// 0 when no page of [a, a + n) can be inaccessible right now (armed by the lazy write-back or probed), so the CPU can
// read it without a fault; 1 when one may be.
// 0 when no page of [a, a + n) can be inaccessible right now (armed by the lazy write-back, in flight, or probed), so
// the CPU can read it without a fault; 1 when one may be.
int mem_lazy_touches(uint64_t a, uint64_t n) {
if (!n) return 0;
if (a >= RT_ADDRESS_LIMIT || a + n > RT_ADDRESS_LIMIT || __atomic_load_n(&prb_on, __ATOMIC_RELAXED)) return 1;
@@ -568,6 +687,7 @@ int mem_lazy_touches(uint64_t a, uint64_t n) {
int hit = 0;
plat_mutex_lock(&lz_lock);
for (int i = 0; i < lz_n && !hit; i++) hit = lzs[i].lo < a + n && lzs[i].hi > a;
if (!hit) hit = lazy_fly_find_locked(a, a + n, 0) >= 0;
plat_mutex_unlock(&lz_lock);
return hit;
}
@@ -585,8 +705,9 @@ int mem_lazy_armed(uint64_t a, uint64_t n) {
// without a fault; ok[i] says which were. An armed range, a probed region or a page the game made inaccessible is left
// out instead of faulting on the caller's thread (vk_memory.c mirror_capture runs inside a submission, where a fetch
// could not be sent). frz_lock is held over all of them, so nothing is armed or probed meanwhile; a range in regions
// with neither (lzr, prb: both change under that lock, with the pages' protection) is copied without asking the system
// for its pages' state. Returns how many were copied.
// with neither (lzr, prb: a page becomes inaccessible only under that lock, after its region is marked, and a mark goes
// only once its pages are read-write again) is copied without asking the system for its pages' state; a range in flight
// gives its pages back meanwhile, which only makes them readable. Returns how many were copied.
// With the lazy heap (NX2NAVI_HEAP_LAZY) a heap page nothing touched yet is still reserved, and reading it faults
// into heap_lazy_fault, which waits for frz_lock: held here, by the same thread, so the copy never returned and every
// later fault, thaw and freeze waited behind it. So a range that meets the heap is copied run by run as the system
@@ -623,19 +744,54 @@ int mem_copy_readable(uint8_t *const *dst, const uint64_t (*src)[2], int n, uint
return got;
}
// Releases what is armed in [a, a + n): before the buffers a fetch would read go away.
void mem_lazy_release(uint64_t a, uint64_t n) {
if (!__atomic_load_n(&lz_on, __ATOMIC_RELAXED) || !n) return;
// Releases what is armed in each range rg[j] = {lo, hi, buffer} (buffer 0: any), and returns once the ranges of those
// buffers in flight there have ended as well: before the buffers a fetch would read go away (vk_memory.c mirror_evict,
// hot_check), and before the kernel's access in a thaw (mem_thaw). Up to LZ_BATCH segments go out at a time and are
// fetched in one call, without frz_lock (with it when NX2NAVI_WB_LAZY_UNLOCKED=0).
void mem_lazy_release(const uint64_t (*rg)[3], int nr) {
if (!__atomic_load_n(&lz_on, __ATOMIC_RELAXED) || nr <= 0) return;
for (;;) {
int k[LZ_BATCH];
LzSeg s[LZ_BATCH];
plat_mutex_lock(&frz_lock);
plat_mutex_lock(&lz_lock);
int n = lazy_take_batch_locked(rg, nr, k, s);
if (n <= 0) {
plat_mutex_unlock(&frz_lock);
if (n < 0) lazy_fly_room_locked();
else lazy_fly_wait_ranges_locked(rg, nr);
plat_mutex_unlock(&lz_lock);
if (n < 0) continue;
return;
}
plat_mutex_unlock(&lz_lock);
lz_other += (uint64_t)n;
if (lz_unlocked) plat_mutex_unlock(&frz_lock);
lazy_fly_run(k, s, n);
if (!lz_unlocked) plat_mutex_unlock(&frz_lock);
}
}
// mem_thaw, before it takes frz_lock: what is armed in the 2 MiB regions of [a, a + n) is released here, its fetches
// without that lock, and the regions are held from arms first (thaw_ms, which thaw_begin_locked sets as well), so the
// thaw finds them settled under the lock. Nothing with NX2NAVI_WB_LAZY_UNLOCKED=0: the thaw releases them itself.
static void lazy_release_regions(uint64_t a, uint64_t n) {
if (!lz_unlocked || !n || a >= RT_ADDRESS_LIMIT || !__atomic_load_n(&lz_on, __ATOMIC_RELAXED)) return;
uint64_t r0 = a >> FRZ_SHIFT, r1 = (a + n - 1 < RT_ADDRESS_LIMIT ? a + n - 1 : RT_ADDRESS_LIMIT - 1) >> FRZ_SHIFT;
int any = 0;
for (uint64_t r = r0; r <= r1 && !any; r++) any = __atomic_load_n(&lzr[r], __ATOMIC_ACQUIRE);
if (!any) return;
plat_mutex_lock(&frz_lock);
for (int i = lz_n; i-- > 0;)
if (lzs[i].lo < a + n && lzs[i].hi > a) lazy_drop_locked(i, 4);
uint32_t now = (uint32_t)plat_now_ms() | 1;
for (uint64_t r = r0; r <= r1; r++) thaw_ms[r] = now;
plat_mutex_unlock(&frz_lock);
uint64_t rg[1][3] = {{r0 << FRZ_SHIFT, (r1 + 1) << FRZ_SHIFT, 0}};
mem_lazy_release((const uint64_t (*)[3])rg, 1);
}
// The back end is ready to fetch: arming starts. 0 when the setting is off (or the probe is on, which watches the
// same pages).
int mem_lazy_enable(void (*fetch)(uint64_t h0, uint64_t h1, uint64_t off, uint64_t size),
int (*clean)(uint64_t lo, uint64_t hi, void *ctx)) {
int mem_lazy_enable(void (*fetch)(const MemLazyCopy *c, int n), int (*clean)(uint64_t lo, uint64_t hi, void *ctx)) {
if (!lz_want) return 0;
lz_fetch = fetch, lz_clean = clean;
__atomic_store_n(&lz_on, 1, __ATOMIC_RELEASE);
@@ -655,10 +811,19 @@ void mem_lazy_window(int open) {
plat_mutex_unlock(&lz_lock);
}
// An access fault at a: 1 when a was armed or in flight and its pages are read-write again (the access runs again), 0
// when a is none of the lazy write-back's.
static int lazy_fault(uint64_t a, int write) {
if (a >= RT_ADDRESS_LIMIT || !__atomic_load_n(&lzr[a >> FRZ_SHIFT], __ATOMIC_ACQUIRE)) return 0;
uint32_t me = plat_thread_id();
for (int tries = 0;; tries++) {
// a range in flight (another thread fetches it or gives it back): wait for that one alone, without frz_lock
if (__atomic_load_n(&lz_nfly, __ATOMIC_ACQUIRE)) {
plat_mutex_lock(&lz_lock);
int joined = lazy_fly_join_locked(a);
plat_mutex_unlock(&lz_lock);
if (joined) return 1;
}
for (int tries = 0;;) {
uint64_t w0 = plat_now_ns();
plat_mutex_lock(&frz_lock);
uint64_t waited = plat_now_ns() - w0;
@@ -666,18 +831,26 @@ static int lazy_fault(uint64_t a, int write) {
plat_mutex_lock(&lz_lock); // the window's test and the segment's leaving the list in one step
int i = 0;
while (i < lz_n && !(a >= lzs[i].lo && a < lzs[i].hi)) i++;
if (i == lz_n) {
plat_mutex_unlock(&lz_lock);
if (i == lz_n) { // in flight since the look above; else not ours, or given back (no arm runs under frz_lock)
plat_mutex_unlock(&frz_lock);
return 0;
int joined = lazy_fly_join_locked(a);
plat_mutex_unlock(&lz_lock);
return joined;
}
uint32_t w = __atomic_load_n(&lz_win_tid, __ATOMIC_ACQUIRE);
if (w && w != me && tries < 400) { // up to 20 ms: after the hand-over the fetch has that part's writes
tries++;
plat_mutex_unlock(&lz_lock);
plat_mutex_unlock(&frz_lock);
plat_wait_ns(50000);
continue;
}
if (lz_nfly >= LZ_FLY) { // every entry in use: wait for one to end, without frz_lock, then look again
plat_mutex_unlock(&frz_lock);
lazy_fly_room_locked();
plat_mutex_unlock(&lz_lock);
continue;
}
if (w) lz_gaps++;
int k = lz_nnote < LZ_NOTES ? lz_nnote++ : -1;
if (k >= 0)
@@ -694,19 +867,26 @@ static int lazy_fault(uint64_t a, int write) {
}
}
if (reader) lz_heavy_ms[r] = now | 1;
if (lz_chunk && !reader) { // the piece leaves the list under lz_lock; its pages come back without it
LzPiece piece = lazy_piece_take_locked(i, a);
plat_mutex_unlock(&lz_lock);
out = lazy_piece_fetch_locked(piece, a, write != 0);
LzSeg s;
if (lz_chunk && !reader) { // the piece goes into flight and is held; the rest of the segment stays armed
LzPiece p = lazy_piece_take_locked(i, a);
s = (LzSeg){p.c0, p.c1, p.s.h0, p.s.h1, p.s.delta, p.s.stale};
out = p.c1 - p.c0;
lazy_piece_hold_locked(a);
lz_pieces++, lz_piece_bytes += out;
if (k >= 0) lz_note[k].hold = lazy_hot_locked(a >> LZ_CHUNK_SHIFT, 1)->back;
} else { // the segment leaves the list under lz_lock; its pages come back without it
LzSeg s = lazy_take_locked(i);
plat_mutex_unlock(&lz_lock);
lazy_gone_locked(s, write != 0);
} else { // the whole segment goes into flight, and its regions are held
s = lazy_take_locked(i);
lazy_hold_locked(s.lo, s.hi);
if (k >= 0) lz_note[k].hold = lz_back[r];
}
lz_faults[write != 0]++;
if (k >= 0) lz_note[k].out = out;
plat_mutex_unlock(&frz_lock);
int f = lazy_fly_start_locked(s);
plat_mutex_unlock(&lz_lock);
if (lz_unlocked) plat_mutex_unlock(&frz_lock);
lazy_fly_run(&f, &s, 1);
if (!lz_unlocked) plat_mutex_unlock(&frz_lock);
return 1;
}
}
@@ -725,19 +905,25 @@ int mem_lazy_pick(uint64_t *addr, int max) {
}
// The lazy write-back's state as a phrase for a fatal error's summary line, "" while nothing is armed by it. No lock:
// the count of armed ranges may be a moment old, and a fetch in progress is named with its thread, which tells a game
// that ended inside a fetch (the device lost under it, or its wait run out) from one that ended elsewhere.
// the count of armed ranges may be a moment old, and a fetch in progress is named with its thread (this thread's
// first, if it has one), which tells a game that ended inside a fetch (the device lost under it, or its wait run out)
// from one that ended elsewhere.
void mem_lazy_note(char *out, size_t n) {
*out = 0;
if (!__atomic_load_n(&lz_on, __ATOMIC_RELAXED)) return;
uint32_t tid = __atomic_load_n(&lz_fetch_tid, __ATOMIC_ACQUIRE), last = __atomic_load_n(&lz_fetch_ms, __ATOMIC_RELAXED);
uint32_t tid = 0, me = plat_thread_id(), last = __atomic_load_n(&lz_fetch_ms, __ATOMIC_RELAXED);
uint64_t bytes = 0;
for (int f = 0; f < LZ_FLY && __atomic_load_n(&lz_nfly, __ATOMIC_RELAXED) && tid != me; f++) { // a hint
uint32_t t = lz_fly[f].tid;
uint64_t b = lz_fly[f].bytes;
if (lz_fly[f].used && b && (!tid || t == me)) tid = t, bytes = b;
}
int k = snprintf(out, n, "%d ranges armed, %llu fetches so far", lz_n,
(unsigned long long)__atomic_load_n(&lz_fetch_total, __ATOMIC_RELAXED));
if (k < 0 || (size_t)k >= n) return;
if (tid)
snprintf(out + k, n - (size_t)k, ", one of %llu bytes in progress on thread %u%s",
(unsigned long long)__atomic_load_n(&lz_fetch_now, __ATOMIC_RELAXED), tid,
tid == plat_thread_id() ? " (this one)" : "");
snprintf(out + k, n - (size_t)k, ", one of %llu bytes in progress on thread %u%s", (unsigned long long)bytes,
tid, tid == me ? " (this one)" : "");
else if (last)
snprintf(out + k, n - (size_t)k, ", the last ended %u ms ago", (uint32_t)plat_now_ms() - last);
}
@@ -753,13 +939,18 @@ void mem_lazy_report(uint64_t frame, uint64_t skipped) {
}
if (!__atomic_load_n(&lz_on, __ATOMIC_RELAXED)) return;
plat_mutex_lock(&frz_lock);
uint64_t bytes = 0, fr = lz_faults[0], fw = lz_faults[1], other = lz_other, nf = lz_fetches, fb = lz_fetch_bytes,
fns = lz_fetch_ns, fmax = lz_fetch_max_ns, wmax = lz_lock_wait_max_ns, gaps = lz_gaps, unclean = lz_unclean,
no[LZ_NO_N];
uint64_t bytes = 0, fr = lz_faults[0], fw = lz_faults[1], other = lz_other, wmax = lz_lock_wait_max_ns,
gaps = lz_gaps, unclean = lz_unclean, no[LZ_NO_N];
int n = lz_n;
for (int i = 0; i < lz_n; i++) bytes += lzs[i].hi - lzs[i].lo;
lz_faults[0] = lz_faults[1] = lz_other = lz_fetches = lz_fetch_bytes = lz_fetch_ns = lz_gaps = lz_unclean = 0;
lz_fetch_max_ns = lz_lock_wait_max_ns = 0;
lz_faults[0] = lz_faults[1] = lz_other = lz_lock_wait_max_ns = lz_gaps = lz_unclean = 0;
plat_mutex_lock(&lz_lock);
uint64_t nf = lz_fetches, fb = lz_fetch_bytes, fns = lz_fetch_ns, fmax = lz_fetch_max_ns, joins = lz_joins,
jmax = lz_join_max_ns;
int most = lz_fly_most;
lz_fetches = lz_fetch_bytes = lz_fetch_ns = lz_fetch_max_ns = lz_joins = lz_join_max_ns = 0;
lz_fly_most = lz_nfly;
plat_mutex_unlock(&lz_lock);
uint64_t small = lz_no_small, calls = lz_arm_calls;
int nnote = lz_nnote;
__typeof__(lz_note) note;
@@ -783,6 +974,10 @@ void mem_lazy_report(uint64_t frame, uint64_t skipped) {
"(a CPU write not in the mirror yet)", (unsigned long long)frame, n, bytes / 1048576.0, skipped / 1048576.0,
(unsigned long long)fr, (unsigned long long)fw, (unsigned long long)gaps, (unsigned long long)other,
(unsigned long long)nf, fb / 1048576.0, fns / 1e6, fmax / 1e6, wmax / 1e6, (unsigned long long)unclean);
if (joins)
rt_log("mem: lazy write-back at frame %llu: %llu touches waited for a range another thread was fetching, the "
"longest %.2f ms; at most %d ranges in flight at once", (unsigned long long)frame,
(unsigned long long)joins, jmax / 1e6, most);
for (int i = 0; i < nnote && (rt_trace & RT_TRACE_PERF); i++)
rt_log("mem: lazy write-back: thread %u %s %#llx, %llu KiB into an armed range of %.2f MiB at %#llx; %llu KiB of "
"it given back%s and held for %u s", note[i].tid, note[i].write ? "wrote" : "read",
@@ -842,7 +1037,7 @@ static void thaw_pages_vm(uint64_t r, int strict) { // vm_lock held
// Returns 1 when the region is frozen.
static int thaw_begin_locked(uint64_t r, int why) {
if (why == 2) thaw_ms[r] = (uint32_t)plat_now_ms() | 1;
if (lzr[r]) lazy_drop_region_locked(r, why + 1); // not a fault on an armed page: no hold
if (lzr[r]) lazy_drop_region_locked(r); // armed ranges, and ranges in flight there end first
if (prb[r]) unprobe_locked(r);
return frz[r] == 1;
}
@@ -971,6 +1166,7 @@ static int heap_lazy_fault(uint64_t a) {
}
void mem_thaw(uint64_t a, uint64_t n) {
lazy_release_regions(a, n); // the lazy write-back's fetches there, without frz_lock
plat_mutex_lock(&frz_lock); // also when nothing is frozen: the hold keeps a read's buffer writable
if (heap_lazy || !frz_async || !n || a >= RT_ADDRESS_LIMIT) {
if (heap_lazy) {
@@ -1045,6 +1241,7 @@ int mem_freeze(uint64_t a) {
plat_mutex_lock(&frz_lock);
uint32_t now = (uint32_t)plat_now_ms();
if (frz[r]) ok = 1;
else if (__atomic_load_n(&lzr[r], __ATOMIC_ACQUIRE)) ok = 0; // armed pages, or a range in flight there
else if (thaw_ms[r] && (int32_t)(now - thaw_ms[r]) < (int32_t)FRZ_HOLD_MS) ok = 0; // |1 may be ahead
else if (frz_hot_until[r] && (int32_t)(frz_hot_until[r] - now) > 0) ok = 0; // hot: held thawed
else if (freeze_locked(r)) {
@@ -1276,6 +1473,7 @@ void mem_init(void) {
plat_mem_access_hook(access_fault);
lz_want = 1;
lz_chunk = cfg_bool("NX2NAVI_WB_LAZY_CHUNK");
lz_unlocked = cfg_bool("NX2NAVI_WB_LAZY_UNLOCKED"); // fetches outside frz_lock ("Ranges in flight")
}
const char *fa = cfg_get("NX2NAVI_FREEZE_ASYNC"); // thaws outside frz_lock (FRZ_THAWING); unset: on
frz_async = !fa || !*fa || atoi(fa) != 0;

172
runtime/vk_memory.c
View File

@@ -2809,17 +2809,41 @@ static uint64_t mirror_evict(uint64_t need, uint32_t min_idle, const char *why)
uint64_t can = 0; // what all the candidates would free: short of need, none goes (no mirror lost for nothing)
for (int i = 0; i < nv; i++) can += v[i]->import_hi - v[i]->import_lo;
if (can < need) nv = 0;
// The victims, as the loop below takes them: their ranges and mirrors. What only a victim's mirror holds (its
// armed ranges the replay left there, memory.c "Lazy write-back") comes down before the mirror goes, and before the
// device address table names the imported buffer again: a shader still running reads through the table as it
// changes, and the fetch waits for all the work handed over. That fetch, one call for all the victims, runs here
// without pool_lock, which the completion thread's arms and every pool lookup wait for. The victims are mirrored
// still, so an arm can land on them before the lock is taken again: the release under the lock below takes that,
// usually nothing. The mirrored list changes meanwhile only by vk_pool_fini, which takes a pool off it under the
// lock: mirror_create, hot_check and this run on the replay thread, under use_lock while chunks record.
uint64_t rel[MAXV][3], f = 0;
int nr = 0;
for (; nr < nv && f < need; nr++) {
PoolRec *r = v[nr];
rel[nr][0] = r->import_lo, rel[nr][1] = r->import_hi, rel[nr][2] = (uint64_t)(uintptr_t)r->mbuf;
f += r->import_hi - r->import_lo;
}
if (nr) {
plat_mutex_unlock(&pool_lock);
mem_lazy_release((const uint64_t (*)[3])rel, nr);
plat_mutex_lock(&pool_lock);
}
uint64_t freed = 0;
uint32_t idle_for[MAXV];
VkBuffer dead[MAXV];
VmaAllocation dead_a[MAXV];
int n = 0;
// All under the exclusive lock: vk_pool_fini, on another thread, takes a pool off the list under it, so a victim is
// still a live pool here, and a finalize that follows finds it imported, its table pages on the imported buffer.
for (; n < nv && freed < need; n++) {
PoolRec *r = v[n];
// All under the exclusive lock: vk_pool_fini, on another thread, takes a pool off the list under it, so a victim
// still on the list is a live pool here, and a finalize that follows finds it imported, its table pages on the
// imported buffer.
for (int k = 0; k < nr; k++) {
PoolRec *r = v[k];
int on = 0;
for (int i = 0; i < nm && !on; i++) on = mlist[i] == r;
if (!on || r->mbuf != (VkBuffer)(uintptr_t)rel[k][2]) continue; // finalized meanwhile: its mirror is fini's
idle_for[n] = msub - r->acc_sub;
mem_lazy_release(r->import_lo, r->import_hi - r->import_lo); // what only the mirror holds comes down first
mem_lazy_release((const uint64_t (*)[3])&rel[k], 1); // what was armed since the release above
mirror_forget(r); // off the list, mirror_gen bumped: no buffer use gets the mirror from here on
r->mstate = MS_NONE, r->acc_sub = 0, r->retry_sub = 0;
dead[n] = r->mbuf, dead_a[n] = r->malloc_;
@@ -2833,6 +2857,7 @@ static uint64_t mirror_evict(uint64_t need, uint32_t min_idle, const char *why)
if (m_evicted_n + n < 8)
rt_log("mem: mirror of pool %#llx+%#llx (%.1f MiB) evicted, idle for %u submissions: %s", (unsigned long long)r->lo,
(unsigned long long)r->size, (r->import_hi - r->import_lo) / 1048576.0, idle_for[n], why);
n++;
}
plat_mutex_unlock(&pool_lock);
for (int i = 0; i < n; i++) vk_defer_buffer(dead[i], dead_a[i]); // after the submission being recorded
@@ -3739,9 +3764,13 @@ static void snap_free(void *arg) { // completion thread: the submission is don
// from the prologue of a part split off earlier) it writes back, as ever, what it skipped and finds no longer armed;
// the snapshot for that merge is taken in the part's own prologue, before its GPU work. From that check to the part's
// hand-over (vk_mem_lazy_handed) memory.c's window keeps a fault from releasing a range unseen.
static VkCommandPool lz_pool;
static VkCommandBuffer lz_cb;
static VkFence lz_fence;
// Fetches run on several threads at once (memory.c "Ranges in flight": no memory.c lock is held around one), each in a
// slot of its own: a command pool (used by one thread at a time), its command buffer and a fence. LZ_SLOTS is more than
// the threads that touch armed pages together; a fetch that finds every slot busy waits 20 us at a time.
#define LZ_SLOTS 4
static struct { VkCommandPool pool; VkCommandBuffer cb; VkFence fence; } lz_slot[LZ_SLOTS];
static volatile uint32_t lz_slot_free; // bit i: slot i is free
static volatile int64_t lz_slot_waits; // fetches that found every slot busy, so far
static volatile int32_t lz_state; // 0 not asked, 1 on, 2 off
static PlatMutex lz_init_lock = PLAT_MUTEX_INIT;
typedef struct { VkMemSpan sp; uint64_t size; int shader; } LzSkip;
@@ -3753,34 +3782,57 @@ static uint64_t lz_skipped; // replay: bytes left in the mirror since th
static void mbarrier(VkCommandBuffer cb, VkPipelineStageFlags2 ss, VkAccessFlags2 sa, VkPipelineStageFlags2 ds,
VkAccessFlags2 da);
// memory.c's fetch (frz_lock held there, so one at a time): [off, off + size) of the mirror h0 into the imported
// buffer h1, after all the GPU work handed over so far.
static void lazy_fetch(uint64_t h0, uint64_t h1, uint64_t off, uint64_t size) {
if (vk_device_is_lost()) return; // the run is ending (vk_device.c vk_device_lost_check): the bytes stay as they are
static int lz_slot_take(void) {
for (int waited = 0;;) {
uint32_t m = __atomic_load_n(&lz_slot_free, __ATOMIC_ACQUIRE);
if (m) {
int i = __builtin_ctz(m);
if (__atomic_compare_exchange_n(&lz_slot_free, &m, m & ~(1u << i), 0, __ATOMIC_ACQ_REL, __ATOMIC_ACQUIRE))
return i;
continue;
}
if (!waited++) plat_inc64(&lz_slot_waits);
plat_wait_ns(20000);
}
}
// memory.c's fetch, with no memory.c lock held (frz_lock in a mapping change's thaw or with
// NX2NAVI_WB_LAZY_UNLOCKED=0): each copy cp[i], [off, off + size) of the mirror h0 into the imported buffer h1, after
// all the GPU work handed over so far, in one submission with one wait.
static void lazy_fetch(const MemLazyCopy *cp, int n) {
if (vk_device_is_lost() || n <= 0) return; // the run is ending (vk_device.c vk_device_lost_check): bytes stay
VkCtx *c = vk_ctx();
VkBufferCopy reg = {off, off, size};
int s = lz_slot_take();
VkCommandBuffer cb = lz_slot[s].cb;
VkCommandBufferBeginInfo bi = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO,
.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT};
vkResetCommandPool(c->device, lz_pool, 0);
vkBeginCommandBuffer(lz_cb, &bi);
mbarrier(lz_cb, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT, VK_ACCESS_2_MEMORY_WRITE_BIT, VK_PIPELINE_STAGE_2_COPY_BIT,
vkResetCommandPool(c->device, lz_slot[s].pool, 0);
vkBeginCommandBuffer(cb, &bi);
mbarrier(cb, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT, VK_ACCESS_2_MEMORY_WRITE_BIT, VK_PIPELINE_STAGE_2_COPY_BIT,
VK_ACCESS_2_TRANSFER_READ_BIT | VK_ACCESS_2_TRANSFER_WRITE_BIT);
vkCmdCopyBuffer(lz_cb, (VkBuffer)(uintptr_t)h0, (VkBuffer)(uintptr_t)h1, 1, &reg);
mbarrier(lz_cb, VK_PIPELINE_STAGE_2_COPY_BIT, VK_ACCESS_2_TRANSFER_WRITE_BIT,
uint64_t bytes = 0;
for (int i = 0; i < n; i++) {
VkBufferCopy reg = {cp[i].off, cp[i].off, cp[i].size};
vkCmdCopyBuffer(cb, (VkBuffer)(uintptr_t)cp[i].h0, (VkBuffer)(uintptr_t)cp[i].h1, 1, &reg);
bytes += cp[i].size;
}
mbarrier(cb, VK_PIPELINE_STAGE_2_COPY_BIT, VK_ACCESS_2_TRANSFER_WRITE_BIT,
VK_PIPELINE_STAGE_2_HOST_BIT | VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
VK_ACCESS_2_HOST_READ_BIT | VK_ACCESS_2_MEMORY_READ_BIT | VK_ACCESS_2_MEMORY_WRITE_BIT);
vkEndCommandBuffer(lz_cb);
VkSubmitInfo si = {.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO, .commandBufferCount = 1, .pCommandBuffers = &lz_cb};
vkEndCommandBuffer(cb);
VkSubmitInfo si = {.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO, .commandBufferCount = 1, .pCommandBuffers = &cb};
vk_queue_lock();
vk_sync_flush_held(); // the held submissions wrote first
VkResult r = vkQueueSubmit(c->gfx_queue, 1, &si, lz_fence);
VkResult r = vkQueueSubmit(c->gfx_queue, 1, &si, lz_slot[s].fence);
vk_queue_unlock();
if (r == VK_SUCCESS) r = vkWaitForFences(c->device, 1, &lz_fence, VK_TRUE, 10000000000ull);
if (r == VK_SUCCESS) r = vkWaitForFences(c->device, 1, &lz_slot[s].fence, VK_TRUE, 10000000000ull);
if (r != VK_SUCCESS) {
vk_device_lost_check(r, "the submit or the wait that brings GPU writes down to guest memory");
rt_fatal("vk: bringing %llu bytes the GPU wrote down to guest memory failed (%d)", (unsigned long long)size, (int)r);
rt_fatal("vk: bringing %llu bytes the GPU wrote down to guest memory failed (%d)", (unsigned long long)bytes,
(int)r);
}
vkResetFences(c->device, 1, &lz_fence);
vkResetFences(c->device, 1, &lz_slot[s].fence);
__atomic_fetch_or(&lz_slot_free, 1u << s, __ATOMIC_RELEASE);
}
// memory.c's test before it arms [lo, hi) of the pool ctx (frz_lock held there, the pages already inaccessible): the
@@ -3831,12 +3883,20 @@ static int wb_lazy(void) {
VkCommandPoolCreateInfo pi = {.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO,
.queueFamilyIndex = c->gfx_family};
VkFenceCreateInfo fi = {.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO};
on = vkCreateCommandPool(c->device, &pi, NULL, &lz_pool) == VK_SUCCESS;
VkCommandBufferAllocateInfo ai = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO, .commandPool = lz_pool,
.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY, .commandBufferCount = 1};
on = on && vkAllocateCommandBuffers(c->device, &ai, &lz_cb) == VK_SUCCESS &&
vkCreateFence(c->device, &fi, NULL, &lz_fence) == VK_SUCCESS;
on = on && mem_lazy_enable(lazy_fetch, lazy_clean);
int made = 0; // the slots that could be made, in order; one is enough
for (; made < LZ_SLOTS; made++) {
if (vkCreateCommandPool(c->device, &pi, NULL, &lz_slot[made].pool) != VK_SUCCESS) break;
VkCommandBufferAllocateInfo ai = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO,
.commandPool = lz_slot[made].pool,
.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY, .commandBufferCount = 1};
if (vkAllocateCommandBuffers(c->device, &ai, &lz_slot[made].cb) != VK_SUCCESS ||
vkCreateFence(c->device, &fi, NULL, &lz_slot[made].fence) != VK_SUCCESS) {
vkDestroyCommandPool(c->device, lz_slot[made].pool, NULL); // frees its command buffer too
break;
}
}
__atomic_store_n(&lz_slot_free, (1u << made) - 1, __ATOMIC_RELEASE);
on = made > 0 && mem_lazy_enable(lazy_fetch, lazy_clean);
}
rt_log("mem: what shaders write into mirrored pools %s (lazy write-back: %s)",
on ? "goes back to guest memory once and then only when the CPU touches it"
@@ -4135,25 +4195,48 @@ static void hot_check(void) {
if (fr - hot_frame0 < HOT_FRAMES || npend || nwbl) return; // nsnl: the previous submission's, recorded already
double frames = (double)(fr - hot_frame0);
hot_frame0 = fr;
PoolRec *out[8];
VkBuffer dead[8];
VmaAllocation dead_a[8];
int nout = 0;
PoolRec *cand[8];
uint64_t rel[8][3]; // each candidate's range and mirror, for mem_lazy_release
int nc = 0;
plat_mutex_lock(&pool_lock);
for (int i = 0; i < nm;) {
for (int i = 0; i < nm; i++) {
PoolRec *r = mlist[i];
uint64_t n = r->import_hi - r->import_lo, per = (uint64_t)(r->hot_bytes / frames), up = (uint64_t)(r->hot_up / frames);
r->hot_bytes = (uint64_t)per; // kept for the log line below
r->hot_up = 0;
if (nout == 8 || per < HOT_MIN || per * 100 < n * (uint64_t)mirror_hot || up > HOT_UP_MAX || !r->mbuf ||
if (nc == 8 || per < HOT_MIN || per * 100 < n * (uint64_t)mirror_hot || up > HOT_UP_MAX || !r->mbuf ||
(!r->twin && r->gpa_rec)) {
r->hot_bytes = 0, i++; // !twin with a table record: the table's writes are in the mirror alone (mirror_evict)
r->hot_bytes = 0; // !twin with a table record: the table's writes are in the mirror alone (mirror_evict)
continue;
}
mem_lazy_release(r->import_lo, r->import_hi - r->import_lo); // what only the mirror holds comes down first
rel[nc][0] = r->import_lo, rel[nc][1] = r->import_hi, rel[nc][2] = (uint64_t)(uintptr_t)r->mbuf;
cand[nc++] = r;
}
if (!nc) {
plat_mutex_unlock(&pool_lock);
return;
}
// What only the mirrors hold comes down first, without pool_lock, as in mirror_evict (which says why, and why the
// release under the lock below usually finds nothing).
plat_mutex_unlock(&pool_lock);
mem_lazy_release((const uint64_t (*)[3])rel, nc);
struct { uint64_t lo, n, per; uint32_t flags; } note[8]; // for the log line: a record may go once the lock drops
VkBuffer dead[8];
VmaAllocation dead_a[8];
int nout = 0;
plat_mutex_lock(&pool_lock);
for (int k = 0; k < nc; k++) {
PoolRec *r = cand[k];
int on = 0;
for (int i = 0; i < nm && !on; i++) on = mlist[i] == r;
if (!on || r->mbuf != (VkBuffer)(uintptr_t)rel[k][2]) continue; // finalized meanwhile: its mirror is fini's
mem_lazy_release((const uint64_t (*)[3])&rel[k], 1); // what was armed since the release above
note[nout].lo = r->lo, note[nout].n = r->import_hi - r->import_lo, note[nout].per = r->hot_bytes;
note[nout].flags = r->flags;
r->hot_bytes = 0;
// as mirror_evict, under the exclusive lock (vk_pool_fini takes pools off the list under it), but for good:
// MS_REFUSED with no retry (retry_sub 0), so neither vk_mem_use nor the budget retries mirror it again
mirror_forget(r); // takes r off mlist: i now names the next pool
mirror_forget(r);
r->mstate = MS_REFUSED, r->retry_sub = 0;
plat_mutex_lock(&flush_lock); // flushes not copied yet: the pool is read in place from now on
__atomic_store_n(&r->fl_n, 0, __ATOMIC_RELEASE);
@@ -4163,19 +4246,16 @@ static void hot_check(void) {
if (r->gpa_rec) r->gpa_rec = vk_gpa_remap_plain(r->import_lo, r->import_hi, r->gpa_rec, r->buf, r->import_lo);
dead[nout] = r->mbuf, dead_a[nout] = r->malloc_;
r->mbuf = VK_NULL_HANDLE, r->malloc_ = NULL;
out[nout++] = r;
nout++;
}
plat_mutex_unlock(&pool_lock);
for (int k = 0; k < nout; k++) {
PoolRec *r = out[k];
uint64_t n = r->import_hi - r->import_lo, per = r->hot_bytes;
r->hot_bytes = 0;
vk_defer_buffer(dead[k], dead_a[k]); // after the submission being recorded
char f[160];
flag_names(r->flags, f, sizeof f);
flag_names(note[k].flags, f, sizeof f);
if (++n_demoted <= 32)
rt_log("mem: pool %#llx (%.1f MiB, %s) copied %.1f MiB a frame through its mirror: system memory from now on",
(unsigned long long)r->lo, n / 1048576.0, f, per / 1048576.0);
(unsigned long long)note[k].lo, note[k].n / 1048576.0, f, note[k].per / 1048576.0);
}
}
@@ -6282,6 +6362,8 @@ void vk_mem_sizes(PerfSizes *s) {
perf_size(s, "of them physical", phys, 0, PERF_SIZE_NONE);
perf_size(s, "VRAM mirrors", (uint64_t)__atomic_load_n(&m_live_n, __ATOMIC_RELAXED), 0,
(uint64_t)__atomic_load_n(&m_live_bytes, __ATOMIC_RELAXED));
perf_size(s, "lazy fetches that found every slot busy", (uint64_t)__atomic_load_n(&lz_slot_waits, __ATOMIC_RELAXED),
0, PERF_SIZE_NONE);
perf_size(s, "deferred destructions", dead_n, 0, PERF_SIZE_NONE);
perf_size(s, "finalized pools destroyed after the GPU", __atomic_load_n(&n_pool_destroyed, __ATOMIC_RELAXED), 0,
PERF_SIZE_NONE);

321
tests/lazy_fetch_check.c
View File

@@ -1,9 +1,10 @@
// runtime/memory.c's lazy write-back (memory.c "Lazy write-back"), built with the Windows platform layer and its real
// fault handler, the rest of the runtime stubbed. The back end's fetch is a stand-in that waits (the GPU copy and the
// wait for it) and then writes a marker over the first 64 bytes of the copy, opening the page for it, as the GPU's
// copy writes past the pages' protection; the back end's clean test may wait too. Built and run by
// runtime/memory.c's lazy write-back (memory.c "Lazy write-back", "Ranges in flight"), built with the Windows platform
// layer and its real fault handler, the rest of the runtime stubbed. The back end's fetch is a stand-in that waits (the
// GPU copy and the wait for it) and then writes a marker over the first 64 bytes of each copy, opening the page for
// it, as the GPU's copy writes past the pages' protection; the back end's clean test may wait too. Built and run by
// tests/lazy_fetch_check.py, one scenario per process (the argument), so an access violation passed to the crash
// handler shows as the exit code. Prints each failed check and exits 1 on any.
// handler shows as the exit code. CHECK_CHUNK in the environment sets NX2NAVI_WB_LAZY_CHUNK=1, CHECK_LOCKED
// NX2NAVI_WB_LAZY_UNLOCKED=0. Prints each failed check and exits 1 on any.
#include "memory.c"
#include <stdarg.h>
@@ -22,6 +23,7 @@ const char *cfg_value(const char *name) { (void)name; return "0"; }
int cfg_bool(const char *name) {
if (!strcmp(name, "NX2NAVI_WB_LAZY")) return 1;
if (!strcmp(name, "NX2NAVI_WB_LAZY_CHUNK")) return getenv("CHECK_CHUNK") != NULL;
if (!strcmp(name, "NX2NAVI_WB_LAZY_UNLOCKED")) return getenv("CHECK_LOCKED") == NULL;
return 0;
}
void cfg_init(void) {}
@@ -58,22 +60,45 @@ void rt_fatal(const char *fmt, ...) {
// ---- the back end's stand-ins
static uint64_t heap;
static volatile uint64_t fetch_ns = 300000000ull, clean_ns;
static volatile int fetch_calls;
static volatile int in_fetch, fetch_calls, fetch_copies, fetch_now, fetch_most, fetch_batch_most;
static volatile uint8_t marker = 0xA0;
static void fake_fetch(uint64_t h0, uint64_t h1, uint64_t off, uint64_t size) {
(void)h0, (void)size;
static struct { uint64_t a, end; } ends[256]; // when the fetch that copied the range at a ended
static volatile int nends;
static void fake_fetch(const MemLazyCopy *c, int n) {
int now = __atomic_add_fetch(&fetch_now, 1, __ATOMIC_SEQ_CST);
for (int m = fetch_most; now > m && !__atomic_compare_exchange_n(&fetch_most, &m, now, 0, __ATOMIC_SEQ_CST,
__ATOMIC_SEQ_CST);) {}
for (int m = fetch_batch_most; n > m && !__atomic_compare_exchange_n(&fetch_batch_most, &m, n, 0, __ATOMIC_SEQ_CST,
__ATOMIC_SEQ_CST);) {}
__atomic_add_fetch(&fetch_calls, 1, __ATOMIC_SEQ_CST);
__atomic_add_fetch(&fetch_copies, n, __ATOMIC_SEQ_CST);
__atomic_store_n(&in_fetch, 1, __ATOMIC_SEQ_CST);
plat_wait_ns(fetch_ns);
uint64_t a = h1 + off; // the GPU's copy writes past the protection; here the page is opened for it
plat_mem_protect(a & ~(PAGE - 1), PAGE, PLAT_PROT_RW);
for (int j = 0; j < 64; j++) ((volatile uint8_t *)(uintptr_t)a)[j] = marker;
plat_mem_protect(a & ~(PAGE - 1), PAGE, PLAT_PROT_NONE);
uint64_t end = plat_now_ns();
for (int i = 0; i < n; i++) { // the GPU's copy writes past the protection; here the page is opened for it
uint64_t a = c[i].h1 + c[i].off;
plat_mem_protect(a & ~(PAGE - 1), PAGE, PLAT_PROT_RW);
for (int j = 0; j < 64; j++) ((volatile uint8_t *)(uintptr_t)a)[j] = marker;
plat_mem_protect(a & ~(PAGE - 1), PAGE, PLAT_PROT_NONE);
int k = __atomic_fetch_add(&nends, 1, __ATOMIC_SEQ_CST);
if (k < 256) ends[k].a = a, ends[k].end = end;
}
__atomic_sub_fetch(&fetch_now, 1, __ATOMIC_SEQ_CST);
}
static int fake_clean(uint64_t lo, uint64_t hi, void *ctx) {
(void)lo, (void)hi, (void)ctx;
if (clean_ns) plat_wait_ns(clean_ns); // GetWriteWatch while the driver pins memory
return 1;
}
// the end of the fetch that copied the range starting at a, 0 when none did
static uint64_t fetch_end(uint64_t a) {
for (int k = 0; k < nends && k < 256; k++)
if (ends[k].a == a) return ends[k].end;
return 0;
}
static void wait_fetch(void) {
while (!__atomic_load_n(&in_fetch, __ATOMIC_SEQ_CST)) plat_wait_ns(1000000);
}
// the range [a, a + n) of the heap, as written back from buffer h0 into the heap's own (h1 = heap, offset a - heap)
static uint64_t arm_buf(uint64_t a, uint64_t n, uint64_t h0) { return mem_lazy_arm(a, n, h0, heap, a - heap, NULL); }
@@ -82,6 +107,15 @@ static int stale(uint64_t a, uint64_t n, uint64_t h0) {
uint64_t out[16][2];
return mem_lazy_skip(a, n, h0, -(int64_t)heap, out, 16);
}
static int armed(uint64_t a, uint64_t n) { return mem_lazy_armed(a, n); }
static void release(uint64_t lo, uint64_t hi, uint64_t h0) {
uint64_t rg[1][3] = {{lo, hi, h0}};
mem_lazy_release((const uint64_t (*)[3])rg, 1);
}
static int page_rw(uint64_t a) {
uint64_t end;
return plat_mem_state(a & ~(PAGE - 1), (a & ~(PAGE - 1)) + PAGE, &end) == PLAT_PROT_RW;
}
static double ms(uint64_t t0) { return (plat_now_ns() - t0) / 1e6; }
// the number after "what" in the last log line, -1 when it is not there
static double logged(const char *what) {
@@ -105,13 +139,19 @@ static void *touch(Touch *t, uint64_t a, int write) {
t->at = a, t->write = write, t->done = 0;
return plat_thread_new(toucher, t, "toucher");
}
static volatile int spin_stop;
static volatile uint64_t spin_reads;
static void spinner(void *arg) {
volatile uint8_t *p = arg;
while (!__atomic_load_n(&spin_stop, __ATOMIC_SEQ_CST)) (void)*p, spin_reads++;
}
static uint64_t rel_rg[1][3];
static volatile int rel_done;
static void releaser(void *arg) {
(void)arg;
mem_lazy_release((const uint64_t (*)[3])rel_rg, 1);
__atomic_store_n(&rel_done, 1, __ATOMIC_SEQ_CST);
}
int main(int argc, char **argv) {
setvbuf(stdout, NULL, _IONBF, 0);
@@ -120,34 +160,247 @@ int main(int argc, char **argv) {
if (mem_set_heap_size(256ull << 20, &heap) != R_OK || !heap) return printf("SetHeapSize failed\n"), 2;
if (!mem_lazy_enable(fake_fetch, fake_clean)) return printf("the lazy write-back is off\n"), 2;
memset((void *)(uintptr_t)heap, 0x11, 256ull << 20);
const uint64_t K64 = 0x10000;
printf("%s: chunk %d, unlocked %d\n", what, lz_chunk, lz_unlocked);
const uint64_t K64 = 0x10000, M = 1ull << 20;
if (!strcmp(what, "arm_race") || !strcmp(what, "arm_race_slow")) {
// A thread reads a page over and over while the completion thread's arm (vk_memory.c vk_gpu_written_back)
// makes its range inaccessible and a release gives it back: no read may reach the crash handler. The slow one
// has the clean test take 100 ms, as GetWriteWatch can while the driver pins memory, with the pages
// inaccessible all along.
uint64_t t = heap + (4ull << 20), r = t >> FRZ_SHIFT;
uint64_t t = heap + 4 * M, r = t >> FRZ_SHIFT;
void *th = plat_thread_new(spinner, (void *)(uintptr_t)(t + 0x100), "spinner");
plat_wait_ns(20000000);
int slow = !strcmp(what, "arm_race_slow");
if (slow) clean_ns = 100000000;
uint64_t armed = 0;
uint64_t n = 0;
for (int i = 0; i < (slow ? 1 : 20000); i++) {
plat_mutex_lock(&frz_lock); // the hold a touch leaves, so that every turn arms again
lz_hot[r] = 0;
plat_mutex_unlock(&frz_lock);
armed += arm_buf(t, K64, 1) != 0;
mem_lazy_release(t, K64);
n += arm_buf(t, K64, 1) != 0;
release(t, t + K64, 0);
}
spin_stop = 1;
plat_thread_join(th);
CHECK(armed, "nothing was armed");
printf(" %llu arms, %llu reads\n", (unsigned long long)armed, (unsigned long long)spin_reads);
CHECK(n, "nothing was armed");
printf(" %llu arms, %llu reads\n", (unsigned long long)n, (unsigned long long)spin_reads);
} else if (!strcmp(what, "no_stall")) {
// During a 300 ms fetch on another thread nothing below waits for it, and the range counts as one a read
// would fault on.
uint64_t s = heap + 8 * M, other = heap + 12 * M, frz_at = heap + 16 * M, thaw_at = heap + 20 * M;
uint64_t rel_at = heap + 24 * M;
CHECK(arm_buf(s, K64, 1) == K64 && stale(s, K64, 1) == 1, "arm");
CHECK(arm_buf(rel_at, K64, 2) == K64, "arm of the range to release");
Touch a;
void *th = touch(&a, s + 0x10, 0);
wait_fetch();
uint64_t t0 = plat_now_ns();
int f = mem_freeze(frz_at);
double d = ms(t0);
CHECK(d < 50 && f == 1, "mem_freeze of another region took %.1f ms (froze %d)", d, f);
uint8_t buf[64], ok, *dst = buf;
uint64_t src[1][2] = {{other + 0x40, 64}};
t0 = plat_now_ns();
int got = mem_copy_readable(&dst, (const uint64_t (*)[2])src, 1, &ok);
d = ms(t0);
CHECK(d < 50 && got == 1, "mem_copy_readable of another range took %.1f ms", d);
t0 = plat_now_ns();
uint64_t n = arm_buf(other, K64, 1);
d = ms(t0);
CHECK(d < 50 && n == K64, "an arm of another range took %.1f ms (%llu bytes)", d, (unsigned long long)n);
t0 = plat_now_ns();
mem_thaw(thaw_at, 0x1000);
d = ms(t0);
CHECK(d < 50, "mem_thaw of another region took %.1f ms", d);
t0 = plat_now_ns();
release(rel_at, rel_at + K64, 2);
d = ms(t0);
CHECK(d < 50 && !armed(rel_at, K64), "a release of another buffer took %.1f ms", d);
CHECK(mem_lazy_touches(s + 0x20, 8) && !armed(s, K64), "the range in flight does not count as touched");
CHECK(!a.done, "the fetch had ended before the checks");
plat_thread_join(th);
CHECK(a.got == marker, "the toucher read %#x, not the fetched bytes", a.got);
CHECK(!mem_lazy_touches(s + 0x20, 8) && page_rw(s), "the range given back still counts as touched");
} else if (!strcmp(what, "join")) {
// B reads the range A is fetching, and with the fetch outside frz_lock this thread holds that lock all along:
// both finish after one fetch, B after it, and the report line counts B's wait.
uint64_t s = heap + 8 * M;
CHECK(arm_buf(s, K64, 1) == K64 && stale(s, K64, 1) == 1, "arm");
Touch a, b;
void *ta = touch(&a, s + 0x10, 0);
wait_fetch();
if (lz_unlocked) plat_mutex_lock(&frz_lock); // with NX2NAVI_WB_LAZY_UNLOCKED=0 the fetch holds it
void *tb = touch(&b, s + 0x10, 0);
for (int i = 0; i < 100 && !(a.done && b.done); i++) plat_wait_ns(10000000);
int both = a.done && b.done;
if (lz_unlocked) plat_mutex_unlock(&frz_lock);
CHECK(both, "the fetch (%d) or the access that waited for it (%d) did not end", a.done, b.done);
plat_thread_join(ta);
plat_thread_join(tb);
CHECK(fetch_calls == 1, "%d fetches", fetch_calls);
CHECK(b.end_ns >= fetch_end(s), "B's read ended before the fetch did");
CHECK(lz_joins == 1, "%llu joins", (unsigned long long)lz_joins);
CHECK(a.got == marker && b.got == marker, "read %#x and %#x, not the fetched bytes", a.got, b.got);
mem_lazy_report(1, 0);
double w = logged("touches waited for a range another thread was fetching, the longest ");
CHECK(w >= 150 && w < 400, "the report's longest wait for a range in flight read %.2f ms", w);
} else if (!strcmp(what, "loser")) {
// B writes the range A is fetching: one fetch, B waits for it alone, and the region's other range stays armed
// (B went to thaw_fault before, which released every armed range of the region)
fetch_ns = 200000000ull;
uint64_t r = heap + 32 * M, s1 = r, s2 = r + M;
CHECK(arm_buf(s1, K64, 1) == K64 && arm_buf(s2, K64, 1) == K64, "arms");
stale(s1, K64, 1), stale(s2, K64, 1);
Touch a, b;
void *ta = touch(&a, s1 + 0x10, 0);
wait_fetch();
plat_wait_ns(20000000);
void *tb = touch(&b, s1 + 0x20, 1);
plat_thread_join(ta);
plat_thread_join(tb);
CHECK(fetch_calls == 1, "%d fetches", fetch_calls);
CHECK(armed(s2, K64), "the region's other range was released");
CHECK(lz_other == 0, "%llu other releases", (unsigned long long)lz_other);
CHECK(b.took < 250 && b.got == 0x77, "the write waited %.1f ms and read back %#x", b.took, b.got);
} else if (!strcmp(what, "concurrent")) {
// two threads touch two ranges at once: their fetches run at the same time
uint64_t s1 = heap + 40 * M, s2 = heap + 48 * M;
CHECK(arm_buf(s1, K64, 1) == K64 && arm_buf(s2, K64, 1) == K64, "arms");
stale(s1, K64, 1), stale(s2, K64, 1);
Touch a, b;
uint64_t t0 = plat_now_ns();
void *ta = touch(&a, s1, 0), *tb = touch(&b, s2, 1);
plat_thread_join(ta);
plat_thread_join(tb);
double d = ms(t0);
CHECK(fetch_most == 2 && d < 450, "at most %d fetches at once, %.1f ms in all", fetch_most, d);
} else if (!strcmp(what, "thaw")) {
// a thaw (a file read into the buffer) of a page in a region with a range in flight returns once that range's
// pages are back, and fetches nothing again
uint64_t s = heap + 56 * M;
CHECK(arm_buf(s, K64, 1) == K64 && stale(s, K64, 1) == 1, "arm");
Touch a;
void *ta = touch(&a, s + 0x10, 0);
wait_fetch();
mem_thaw(s + 0x1000, 0x1000);
uint64_t now = plat_now_ns(), fe = fetch_end(s);
CHECK(fe && now >= fe, "mem_thaw returned before the fetch ended");
CHECK(page_rw(s + 0x1000), "the page is not read-write after the thaw");
plat_thread_join(ta);
CHECK(fetch_calls == 1, "%d fetches", fetch_calls);
} else if (!strcmp(what, "mapping")) {
// a mapping change (a commit, decommit or protect: thaw_range_locked under frz_lock) in a region with a range
// in flight and another armed: the armed one is fetched under the lock, the flight ends first, both given back
uint64_t r = heap + 60 * M, s1 = r, s2 = r + M;
CHECK(arm_buf(s1, K64, 1) == K64 && arm_buf(s2, K64, 1) == K64, "arms");
stale(s1, K64, 1), stale(s2, K64, 1);
Touch a;
void *ta = touch(&a, s1 + 0x10, 0);
wait_fetch();
plat_mutex_lock(&frz_lock);
thaw_wait_locked(r, FRZ_SIZE);
thaw_range_locked(r, FRZ_SIZE, 3);
uint64_t now = plat_now_ns();
int rw = page_rw(s1) && page_rw(s2), mark = lzr[r >> FRZ_SHIFT];
plat_mutex_unlock(&frz_lock);
CHECK(fetch_end(s1) && now >= fetch_end(s1), "the mapping change went ahead of the range in flight");
CHECK(fetch_end(s2) && rw && !mark && !armed(s2, K64), "the region is not settled (pages %d, mark %d)", rw,
mark);
plat_thread_join(ta);
CHECK(fetch_calls == 2, "%d fetches", fetch_calls);
} else if (!strcmp(what, "release")) {
// A release of one buffer's ranges (a mirror eviction): one fetch of its two armed stale segments, and it
// returns after the fetch another thread runs of its third; the other buffer's range stays armed.
uint64_t base = heap + 64 * M;
uint64_t x1 = base, x2 = base + M, x3 = base + 3 * M, y = base + 2 * M;
CHECK(arm_buf(x1, K64, 1) == K64 && arm_buf(x2, K64, 1) == K64 && arm_buf(x3, K64, 1) == K64, "arms of X");
CHECK(arm_buf(y, K64, 7) == K64, "arm of Y");
stale(x1, K64, 1), stale(x2, K64, 1), stale(x3, K64, 1), stale(y, K64, 7);
Touch a;
void *ta = touch(&a, x3 + 0x10, 0); // in flight on another thread
wait_fetch();
release(base, base + 4 * M, 1);
uint64_t now = plat_now_ns(), fe = fetch_end(x3);
CHECK(fe && now >= fe, "the release returned before the other thread's fetch of the buffer ended");
CHECK(!armed(x1, K64) && !armed(x2, K64) && page_rw(x1) && page_rw(x2), "X still armed");
CHECK(armed(y, K64), "Y was released");
CHECK(fetch_batch_most == 2, "the release's fetch had %d copies", fetch_batch_most);
plat_thread_join(ta);
CHECK(fetch_calls == 2 && fetch_copies == 3, "%d fetches of %d copies", fetch_calls, fetch_copies);
} else if (!strcmp(what, "release_join")) {
// a thread writes a range a release (a mirror eviction) is fetching: it waits for that fetch, its write
// survives, one fetch in all
uint64_t x = heap + 88 * M;
CHECK(arm_buf(x, K64, 3) == K64 && stale(x, K64, 3) == 1, "arm");
rel_rg[0][0] = x, rel_rg[0][1] = x + K64, rel_rg[0][2] = 3;
void *tr = plat_thread_new(releaser, NULL, "releaser");
wait_fetch();
Touch b;
void *tb = touch(&b, x + 0x20, 1);
plat_thread_join(tr);
plat_thread_join(tb);
CHECK(rel_done && fetch_calls == 1, "%d fetches", fetch_calls);
CHECK(b.end_ns >= fetch_end(x) && b.got == 0x77, "the write ended before the fetch, or was lost (%#x)", b.got);
CHECK(lz_joins == 1 && !armed(x, K64), "%llu joins", (unsigned long long)lz_joins);
} else if (!strcmp(what, "arm_busy")) {
// an arm of a range in flight (its region's hold cleared, as when the range was released, not touched) leaves
// it alone: its pages are given back by the flight, and a later arm takes it again
uint64_t s = heap + 92 * M, r = s >> FRZ_SHIFT;
CHECK(arm_buf(s, K64, 1) == K64 && stale(s, K64, 1) == 1, "arm");
rel_rg[0][0] = s, rel_rg[0][1] = s + K64, rel_rg[0][2] = 0;
void *tr = plat_thread_new(releaser, NULL, "releaser");
wait_fetch();
plat_mutex_lock(&frz_lock);
int held = lz_hot[r] != 0;
plat_mutex_unlock(&frz_lock);
uint64_t n = arm_buf(s, K64, 1);
CHECK(!held && n == 0 && lzr[r], "an arm of the range in flight armed %llu bytes (hold %d, mark %d)",
(unsigned long long)n, held, lzr[r]);
plat_thread_join(tr);
CHECK(page_rw(s) && !lzr[r], "the range in flight was not given back");
CHECK(arm_buf(s, K64, 1) == K64 && armed(s, K64), "the range was not armed again after its flight");
} else if (!strcmp(what, "foreign")) {
// a write by the replay to a range armed for another buffer (a mirror gone since) gives it back without a
// fetch: that buffer may be gone
uint64_t s = heap + 100 * M;
CHECK(arm_buf(s, K64, 1) == K64 && stale(s, K64, 1) == 1, "arm");
CHECK(stale(s, K64, 2) == 0, "the range armed for another buffer was marked");
CHECK(fetch_calls == 0 && !armed(s, K64) && page_rw(s) && !lzr[s >> FRZ_SHIFT] && lz_foreign == 1,
"%d fetches, armed %d, read-write %d", fetch_calls, armed(s, K64), page_rw(s));
} else if (!strcmp(what, "chunk")) {
// NX2NAVI_WB_LAZY_CHUNK=1: lz_lock is free during a piece's fetch; the rest of the range stays armed
uint64_t s = heap + 72 * M, other = heap + 76 * M;
CHECK(arm_buf(s, M, 1) == M && stale(s, M, 1) == 1, "arm");
Touch a;
void *ta = touch(&a, s + 0x10, 0);
wait_fetch();
uint64_t t0 = plat_now_ns();
int am = armed(other, 0x1000);
double d = ms(t0);
CHECK(d < 50 && !am, "mem_lazy_armed took %.1f ms", d);
t0 = plat_now_ns();
int k = stale(s + M / 2, 0x1000, 1);
d = ms(t0);
CHECK(d < 50 && k == 1, "mem_lazy_skip in the same range took %.1f ms (%d)", d, k);
plat_thread_join(ta);
CHECK(armed(s + M / 2, 0x1000) && !armed(s, 0x1000) && page_rw(s), "the piece and the rest");
} else if (!strcmp(what, "locked")) {
// NX2NAVI_WB_LAZY_UNLOCKED=0: the fetch holds frz_lock as before
uint64_t s = heap + 8 * M, frz_at = heap + 16 * M;
CHECK(arm_buf(s, K64, 1) == K64 && stale(s, K64, 1) == 1, "arm");
Touch a;
void *th = touch(&a, s + 0x10, 0);
wait_fetch();
uint64_t t0 = plat_now_ns();
mem_freeze(frz_at);
double d = ms(t0);
CHECK(d > 200, "mem_freeze took %.1f ms, not the fetch's time", d);
plat_thread_join(th);
} else if (!strcmp(what, "report")) {
// The report line gives the longest fetch since the last line, and the longest wait of a fault for frz_lock;
// the next line starts from nothing.
uint64_t s1 = heap + (80ull << 20), s2 = heap + (84ull << 20), s3 = heap + (88ull << 20);
uint64_t s1 = heap + 80 * M, s2 = heap + 84 * M, s3 = heap + 104 * M;
CHECK(arm_buf(s1, K64, 1) == K64 && arm_buf(s2, K64, 1) == K64 && arm_buf(s3, K64, 1) == K64, "arms");
CHECK(stale(s1, K64, 1) == 1 && stale(s2, K64, 1) == 1 && stale(s3, K64, 1) == 1, "stale marks");
Touch a;
@@ -168,6 +421,32 @@ int main(int argc, char **argv) {
CHECK(waited >= 40 && waited < 200, "the longest wait for the lock read %.2f ms", waited);
mem_lazy_report(2, 0);
CHECK(logged("the longest ") == 0 && logged("faults waited up to ") == 0, "the next line did not start over");
} else if (!strcmp(what, "flood")) {
// more threads than entries touch a range each at once: those past LZ_FLY wait for room, and all end
enum { N = 80 };
fetch_ns = 200000000ull;
static Touch t[N];
void *th[N];
for (int i = 0; i < N; i++) {
uint64_t s = heap + 96 * M + (uint64_t)i * M;
CHECK(arm_buf(s, K64, 1) == K64 && stale(s, K64, 1) == 1, "arm %d", i);
}
for (int i = 0; i < N; i++) th[i] = touch(&t[i], heap + 96 * M + (uint64_t)i * M, i & 1);
for (int i = 0; i < N; i++) plat_thread_join(th[i]);
int bad = 0;
for (int i = 0; i < N; i++) bad += t[i].got != (i & 1 ? 0x77 : marker);
CHECK(fetch_copies == N && lz_fly_most == LZ_FLY && !bad && !lz_nfly,
"%d copies, %d in flight at most, %d wrong reads, %d left in flight", fetch_copies, lz_fly_most, bad,
(int)lz_nfly);
} else if (!strcmp(what, "freeze")) {
// a region with a range in flight is not frozen
uint64_t s = heap + 8 * M;
CHECK(arm_buf(s, K64, 1) == K64 && stale(s, K64, 1) == 1, "arm");
Touch a;
void *th = touch(&a, s + 0x10, 0);
wait_fetch();
CHECK(mem_freeze(s) == 0, "a region with a range in flight was frozen");
plat_thread_join(th);
} else {
printf("unknown scenario %s\n", what);
return 2;

29
tests/lazy_fetch_check.py
View File

@@ -1,7 +1,9 @@
"""Build and run tests/lazy_fetch_check.c: runtime/memory.c's lazy write-back (memory.c "Lazy write-back") with the
Windows platform layer and its real fault handler, the rest of the runtime stubbed and the back end's fetch a stand-in
that only waits. Each scenario runs as its own process, so an access violation passed to the crash handler shows as
its exit code; any exit code other than 0, or a scenario that does not end within 120 s, fails.
"""Build and run tests/lazy_fetch_check.c: runtime/memory.c's lazy write-back (memory.c "Lazy write-back", "Ranges in
flight") with the Windows platform layer and its real fault handler, the rest of the runtime stubbed and the back end's
fetch a stand-in that only waits. It checks that a fetch holds up nothing but the touches of its range, that those
wait for it, and that arms, thaws, mapping changes, freezes and releases see it; also with NX2NAVI_WB_LAZY_CHUNK=1 and
with NX2NAVI_WB_LAZY_UNLOCKED=0. Each scenario runs as its own process, so an access violation passed to the crash
handler shows as its exit code; any exit code other than 0, or a scenario that does not end within 120 s, fails.
Windows only (the fault handler is): no game files, no GPU. The module layout the runtime includes is a stand-in
written to the test's build folder.
@@ -17,10 +19,29 @@ from pathlib import Path
ROOT = Path(__file__).resolve().parents[1]
MSYS_BIN = Path(r'C:\msys64\ucrt64\bin')
# name: (scenario argument, extra environment)
CHUNK = {'CHECK_CHUNK': '1'} # NX2NAVI_WB_LAZY_CHUNK=1
LOCKED = {'CHECK_LOCKED': '1'} # NX2NAVI_WB_LAZY_UNLOCKED=0
SCENARIOS = {
'arm_race': ('arm_race', {}),
'arm_race_slow': ('arm_race_slow', {}),
'no_stall': ('no_stall', {}),
'join': ('join', {}),
'loser': ('loser', {}),
'concurrent': ('concurrent', {}),
'thaw': ('thaw', {}),
'mapping': ('mapping', {}),
'release': ('release', {}),
'release_join': ('release_join', {}),
'arm_busy': ('arm_busy', {}),
'foreign': ('foreign', {}),
'chunk': ('chunk', CHUNK),
'locked': ('locked', LOCKED),
'locked_join': ('join', LOCKED),
'locked_loser': ('loser', LOCKED),
'locked_arm_race': ('arm_race_slow', LOCKED),
'report': ('report', {}),
'flood': ('flood', {}),
'freeze': ('freeze', {}),
}