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
Files
8e656b85a44c7c0fc59a40e3f4cc114d2eaca0ee
NX2navi/runtime/vk_memory.c
chrg6 8e656b85a4 Hold less video memory in empty space: smaller VMA blocks, spare blocks only where textures grow, a smaller ASTC scratch, and DLSS features freed once unused, with a VRAM breakdown in the perf log. 
VMA's blocks on large heaps are 64 MiB instead of 256 (vma_block_mb): one session held 2.1 GiB free inside the
device blocks beside 8.2 GiB of allocations after a few area changes, and the driver counts a block whole until it
empties. The headroom thread keeps its spare block only for the device-local types textures and mirrors grow in; the
spare blocks it made for the host types and the PCIe BAR type at boot stayed empty for the whole run. The ASTC
decoder's scratch is 32 MiB, so it stays inside a block, and an area load that makes one per upload batch holds half
as much at once. A DLSS feature no frame used for 3600 evaluations is freed with its work images, and a model change
now frees all five sizes.

The perf log gets a "vram by use" line (textures, released textures, mirrors, physical pools, ASTC scratch, headroom
spare, free space inside the blocks, memory outside VMA and DLSS's own share), the free ranges per memory type, the
addresses whose duplicate texture entries hold the most, and a line when memory outside VMA moves by 256 MiB.
astc_bc_probe in nx2navi.ini turns on the ASTC opacity probe.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
(cherry picked from commit 99cfc2fa14ae208b7b23e057ff9b87986cd6a5e0)
2026-11-09 14:56:17 -05:00

3158 lines
166 KiB
C
Raw Blame History

// Vulkan memory of the graphics layer: NVN memory pools as Vulkan
// buffers, the virtual-pool page table, GPU address resolution, page write generations for re-uploads, and the
// upload batch (command buffer, staging buffers, deferred destruction) that vk_texture.c records into.
//
// Contains code Copyright (c) 2026 jhaizhou, Copyright (c) 2026 Peiikai (MIT; see THIRD_PARTY_NOTICES.md).
// Virtual pools have 64 KiB pages. A texture is stale when a covered page carries a generation newer than its upload;
// every write-watch query resets the OS bits, so each query stamps the pages it saw with a new generation shared by all
// users. Virtual pools are a GPU-side page table
// over device-local memory and dirty pages come from GetWriteWatch (mem_written_pages), so neither host VM aliasing
// nor page protection (a fault handler) is needed.
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "cfg.h"
#include "gfx.h"
#include "host.h"
#include "nvn_api.h"
#include "nvn_pool_index.h"
#include "vk.h"
#include "vk_dlaa.h"
#include "vk_exec.h"
#include "vk_gpa.h"
#include "wb_merge_spv.h"
#include "platform/platform.h"
// Timestamps for the tracking cost, in ns; qpc_ms converts.
static uint64_t qpc(void) { return plat_now_ns(); }
static double qpc_ms(uint64_t ticks) { return (double)ticks / 1e6; }
// ---------------------------------------------------------------- pools
enum { POOL_IMPORT = 1, POOL_PHYS, POOL_VIRT };
enum { MS_NONE, MS_ON, MS_REFUSED };
typedef struct { struct PoolRec *phys; uint64_t off; } VPage; // virtual page -> physical pool and offset
#define FLUSH_RANGES 32 // flush ranges a mirrored pool keeps between two syncs (flush_add)
typedef struct { uint64_t lo, hi; } FlushRange;
typedef struct PoolRec {
const pool_t *key;
int kind;
uint32_t flags;
uint64_t lo, size; // GPU address range of the pool
uint64_t import_lo, import_hi; // POOL_IMPORT: the imported host range (aligned outwards)
VkBuffer buf;
VkDeviceMemory mem; // POOL_IMPORT
VmaAllocation alloc; // POOL_PHYS
VPage *pt; // POOL_VIRT: size / 64 KiB entries
uint64_t npages;
volatile int32_t used; // resolved at least once: the GPU may still read it at finalize
// NX2NAVI_TRACE=perf dirty census (vk_mem_dirty_submit), replay thread under the submit lock only:
uint32_t bound_sub; // last submit that bound it as a buffer (dirty_sub), 0 never
uint32_t synced_gen; // page generation its lazy-mirror count covers; DIRTY_UNSYNCED until scanned
uint8_t bufbound; // ever bound as a vertex, index, uniform, storage or indirect buffer
uint8_t untracked; // the last dirty scan found no write watch over it
// VRAM mirror (POOL_IMPORT, "VRAM mirror" below). mstate and mbuf change on the replay thread only.
uint8_t mstate; // MS_NONE until the first buffer use, then MS_ON or MS_REFUSED
uint32_t gpa_rec; // POOL_IMPORT: its record in the device address table (vk_gpa.c), 0 none
struct GpaMap { uint64_t lo, hi; uint32_t rec; } *vmaps; // POOL_VIRT: its mapped ranges in the table
int nvmaps, cvmaps;
VkBuffer mbuf; // MS_ON: device-local copy of [import_lo, import_hi), same offsets as buf
VmaAllocation malloc_;
uint32_t acc_sub; // replay: the last submission that synced the mirror (mirror_sync)
uint32_t wb_sub; // replay: the last submission that wrote mirror bytes back to guest memory
uint32_t retry_sub; // replay: a mirror refused for the VRAM budget is asked again from here, 0 never
uint8_t asked_room; // replay: its refusal already asked the textures for room (once until it gets one)
uint8_t twin; // MS_ON: the table's writes reach guest memory too (twin_wanted); 0: the mirror only
uint32_t vd_wrote_sub, vd_use_sub; // verify mode: the last submission that bound it written / used its mirror
uint64_t vd_wrote_n; // verify mode: written bindings of it so far
uint32_t mgen; // gen_lock: every page write stamped up to this generation is in the mirror
// (or in a copy queued before its next use)
FlushRange fl[FLUSH_RANGES + 1]; // flush_lock: nvnMemoryPoolFlushMappedRange pages not copied yet (flush_add)
int fl_n; // flush_lock: ranges in fl; read without the lock to skip an empty list
uint64_t fl_calls, fl_sum; // flush_lock: the calls and bytes behind fl, for mirror_sync's statistics
// Shadow storage (POOL_IMPORT in shadow mode, "shadow storage" below): buf is a host-visible copy of the guest pages.
VmaAllocation salloc;
uint8_t *smap; // its mapping; offsets as in buf (from import_lo)
uint32_t sgen; // gen_lock: every page write stamped up to this generation is in the shadow
volatile uint8_t stouched; // on the touched list: the GPU may use it in the submission being recorded
// NX2NAVI_TRACE=perf per-pool mirror traffic since the last mirror_pool_log, replay thread (the log reads and clears
// them without a lock: a count lost to the race only blurs a diagnostic)
uint64_t st_pages, st_flush, st_wb, st_merge; // bytes: written pages and flush ranges copied up, GPU writes back
uint32_t st_syncs; // syncs
} PoolRec;
static void mirror_forget(PoolRec *r); // pool_lock exclusive: off the mirrored list
static int shadow_on; // shadow mode: imported pools get shadow storage instead (vk_mem_class_log)
static int shadow_pool(VkCtx *c, PoolRec *r, uint64_t addr, uint64_t size);
static void shadow_touch(PoolRec *r);
static void shadow_forget(PoolRec *r);
static void mirror_retiring(PoolRec *r); // vk_pool_fini: kept until mirror_retire (or vk_mem_shutdown)
static void mirror_retire(void *arg); // completion thread: destroys a finalized mirrored pool's record
static void grow(void **arr, int *cap, size_t elem);
static PlatMutex pool_lock = PLAT_MUTEX_INIT;
// Bumped under pool_lock (exclusive) by every change to the index or a page table: vk_mem_resolve's per-thread hits
// stay valid while it is unchanged.
static uint64_t pool_gen = 1;
// Bumped when a buffer use may get another buffer for the same span (a mirror created, refused or forgotten) and at
// every submission's start (mirror syncs are per submission). vk_mem_gen: both together.
static uint64_t mirror_gen;
static pidx pool_index; // POOL_IMPORT, POOL_PHYS and POOL_VIRT records by GPU address
static PoolRec **phys_list; // physical pools by key, for MapVirtual requests
static int nphys, cap_phys;
static int32_t n_imports, n_import_failures;
static PFN_vkGetMemoryHostPointerPropertiesEXT p_host_props;
static VkBufferUsageFlags pool_usage(const VkCtx *c) {
VkBufferUsageFlags u = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT |
VK_BUFFER_USAGE_UNIFORM_TEXEL_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_TEXEL_BUFFER_BIT |
VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT |
VK_BUFFER_USAGE_INDIRECT_BUFFER_BIT;
if (c->has_transform_feedback)
u |= VK_BUFFER_USAGE_TRANSFORM_FEEDBACK_BUFFER_BIT_EXT | VK_BUFFER_USAGE_TRANSFORM_FEEDBACK_COUNTER_BUFFER_BIT_EXT;
if (c->has_conditional_rendering) u |= VK_BUFFER_USAGE_CONDITIONAL_RENDERING_BIT_EXT;
if (vk_gpa_wanted()) u |= VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT; // pages of the device address table
return u;
}
static int memory_type(const VkCtx *c, uint32_t bits, VkMemoryPropertyFlags want) {
for (uint32_t i = 0; i < c->mem_props.memoryTypeCount; i++)
if ((bits >> i & 1) && (c->mem_props.memoryTypes[i].propertyFlags & want) == want) return (int)i;
for (uint32_t i = 0; i < c->mem_props.memoryTypeCount; i++)
if (bits >> i & 1) return (int)i;
return -1;
}
// CPU-visible storage: the guest pages themselves become the buffer's memory (VK_EXT_external_memory_host).
static int import_pool(VkCtx *c, PoolRec *r, uint64_t addr, uint64_t size) {
if (shadow_on) return shadow_pool(c, r, addr, size);
if (!p_host_props)
p_host_props = (PFN_vkGetMemoryHostPointerPropertiesEXT)vk_dpa(c->device, "vkGetMemoryHostPointerPropertiesEXT");
uint64_t al = c->min_imported_host_alignment ? c->min_imported_host_alignment : 0x1000;
r->import_lo = addr & ~(al - 1);
r->import_hi = (addr + size + al - 1) & ~(al - 1);
uint64_t n = r->import_hi - r->import_lo;
VkMemoryHostPointerPropertiesEXT hp = {.sType = VK_STRUCTURE_TYPE_MEMORY_HOST_POINTER_PROPERTIES_EXT};
VkResult res = p_host_props ? p_host_props(c->device, VK_EXTERNAL_MEMORY_HANDLE_TYPE_HOST_ALLOCATION_BIT_EXT,
(void *)(uintptr_t)r->import_lo, &hp)
: VK_ERROR_EXTENSION_NOT_PRESENT;
if (res != VK_SUCCESS || !hp.memoryTypeBits) {
rt_log("vk: pool %#llx+%#llx: host pointer not importable (%d)", (unsigned long long)addr, (unsigned long long)size, (int)res);
return 0;
}
VkExternalMemoryBufferCreateInfo eb = {.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO,
.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_HOST_ALLOCATION_BIT_EXT};
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .pNext = &eb, .size = n,
.usage = pool_usage(c), .sharingMode = VK_SHARING_MODE_EXCLUSIVE};
if ((res = vkCreateBuffer(c->device, &bi, NULL, &r->buf)) != VK_SUCCESS) {
rt_log("vk: pool %#llx: vkCreateBuffer for the import failed (%d)", (unsigned long long)addr, (int)res);
return 0;
}
VkMemoryRequirements req;
vkGetBufferMemoryRequirements(c->device, r->buf, &req);
int type = memory_type(c, req.memoryTypeBits & hp.memoryTypeBits,
VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT);
if (type < 0 || req.size > n) {
rt_log("vk: pool %#llx: no memory type for the import (buffer %#x, host %#x, size %#llx of %#llx)",
(unsigned long long)addr, req.memoryTypeBits, hp.memoryTypeBits, (unsigned long long)req.size,
(unsigned long long)n);
vkDestroyBuffer(c->device, r->buf, NULL);
r->buf = VK_NULL_HANDLE;
return 0;
}
VkImportMemoryHostPointerInfoEXT im = {.sType = VK_STRUCTURE_TYPE_IMPORT_MEMORY_HOST_POINTER_INFO_EXT,
.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_HOST_ALLOCATION_BIT_EXT,
.pHostPointer = (void *)(uintptr_t)r->import_lo};
// A buffer with device address usage needs memory allocated for device addresses (vk_gpa.c).
VkMemoryAllocateFlagsInfo fi = {.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_FLAGS_INFO, .pNext = &im,
.flags = VK_MEMORY_ALLOCATE_DEVICE_ADDRESS_BIT};
VkMemoryAllocateInfo ai = {.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO, .pNext = vk_gpa_wanted() ? (void *)&fi : (void *)&im,
.allocationSize = n, .memoryTypeIndex = (uint32_t)type};
if ((res = vkAllocateMemory(c->device, &ai, NULL, &r->mem)) != VK_SUCCESS ||
(res = vkBindBufferMemory(c->device, r->buf, r->mem, 0)) != VK_SUCCESS) {
rt_log("vk: pool %#llx+%#llx: host memory import failed (%d)", (unsigned long long)addr, (unsigned long long)size, (int)res);
if (r->mem) vkFreeMemory(c->device, r->mem, NULL);
vkDestroyBuffer(c->device, r->buf, NULL);
r->buf = VK_NULL_HANDLE;
r->mem = VK_NULL_HANDLE;
return 0;
}
plat_inc32(&n_imports);
return 1;
}
// Physical pool: device-local memory of the pool size that both buffers and optimal-tiling images can be bound to.
static int alloc_physical(VkCtx *c, PoolRec *r) {
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .size = r->size, .usage = pool_usage(c),
.sharingMode = VK_SHARING_MODE_EXCLUSIVE};
VkResult res = vkCreateBuffer(c->device, &bi, NULL, &r->buf);
if (res != VK_SUCCESS) { rt_log("vk: physical pool: vkCreateBuffer failed (%d)", (int)res); return 0; }
VkMemoryRequirements rb;
vkGetBufferMemoryRequirements(c->device, r->buf, &rb);
VkImageCreateInfo ii = {.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO, .imageType = VK_IMAGE_TYPE_2D,
.format = VK_FORMAT_R8G8B8A8_UNORM, .extent = {256, 256, 1}, .mipLevels = 1, .arrayLayers = 1,
.samples = VK_SAMPLE_COUNT_1_BIT, .tiling = VK_IMAGE_TILING_OPTIMAL,
.usage = VK_IMAGE_USAGE_SAMPLED_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT | VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT};
VkDeviceImageMemoryRequirements di = {.sType = VK_STRUCTURE_TYPE_DEVICE_IMAGE_MEMORY_REQUIREMENTS, .pCreateInfo = &ii};
VkMemoryRequirements2 ri = {.sType = VK_STRUCTURE_TYPE_MEMORY_REQUIREMENTS_2};
vkGetDeviceImageMemoryRequirements(c->device, &di, &ri);
uint32_t bits = rb.memoryTypeBits & ri.memoryRequirements.memoryTypeBits;
if (!bits) bits = rb.memoryTypeBits;
VkDeviceSize align = NVN_VPOOL_PAGE;
if (rb.alignment > align) align = rb.alignment;
if (ri.memoryRequirements.alignment > align) align = ri.memoryRequirements.alignment;
VkMemoryRequirements req = {.size = (rb.size + align - 1) & ~(align - 1), .alignment = align, .memoryTypeBits = bits};
VmaAllocationCreateInfo aci = {.flags = VMA_ALLOCATION_CREATE_DEDICATED_MEMORY_BIT,
.requiredFlags = VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT};
if ((res = vmaAllocateMemory(c->vma, &req, &aci, &r->alloc, NULL)) != VK_SUCCESS ||
(res = vmaBindBufferMemory(c->vma, r->alloc, r->buf)) != VK_SUCCESS) {
rt_log("vk: physical pool of %#llx bytes: device-local allocation failed (%d)", (unsigned long long)r->size, (int)res);
if (r->alloc) vmaFreeMemory(c->vma, r->alloc);
vkDestroyBuffer(c->device, r->buf, NULL);
r->buf = VK_NULL_HANDLE;
r->alloc = NULL;
return 0;
}
if (vk_fix_zero()) { // Switch pool memory starts zeroed; once per pool, not per image placed in it
VkCommandBuffer cb = vk_upload_begin();
vkCmdFillBuffer(cb, r->buf, 0, VK_WHOLE_SIZE, 0);
VkMemoryBarrier2 mb = {.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER_2, .srcStageMask = VK_PIPELINE_STAGE_2_ALL_TRANSFER_BIT,
.srcAccessMask = VK_ACCESS_2_TRANSFER_WRITE_BIT, .dstStageMask = VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
.dstAccessMask = VK_ACCESS_2_MEMORY_READ_BIT | VK_ACCESS_2_MEMORY_WRITE_BIT};
VkDependencyInfo di = {.sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO, .memoryBarrierCount = 1, .pMemoryBarriers = &mb};
vkCmdPipelineBarrier2(cb, &di);
vk_upload_end();
}
return 1;
}
int vk_fix_zero(void) {
static int on = -1;
if (on < 0) {
on = cfg_bool("NX2NAVI_FIX_ZERO");
}
return on;
}
static void index_put(PoolRec *r) {
__atomic_add_fetch(&pool_gen, 1, __ATOMIC_RELEASE);
pidx_ent e = {.lo = r->lo, .hi = r->lo + r->size, .base = r->lo, .buf = r, .gpu = r->lo, .key = r->key};
pidx_remove(&pool_index, r->key, NULL); // an object initialized again without Finalize
if (!pidx_insert(&pool_index, &e)) rt_fatal("vk: out of memory for the pool index");
}
static PoolRec *phys_find(const pool_t *key) {
for (int i = 0; i < nphys; i++) if (phys_list[i]->key == key) return phys_list[i];
return NULL;
}
static void destroy_rec(VkCtx *c, PoolRec *r) {
if (r->mbuf) vmaDestroyBuffer(c->vma, r->mbuf, r->malloc_);
if (r->salloc) vmaDestroyBuffer(c->vma, r->buf, r->salloc), r->buf = VK_NULL_HANDLE;
if (r->buf) vkDestroyBuffer(c->device, r->buf, NULL);
if (r->mem) vkFreeMemory(c->device, r->mem, NULL);
if (r->alloc) vmaFreeMemory(c->vma, r->alloc);
free(r->pt);
free(r);
}
int vk_pool_init(pool_t *p) {
VkCtx *c = vk_ctx();
if (!c) { rt_log("vk: MemoryPoolInitialize before vk_init"); return 0; }
// The GPU never sees a GPU_NO_ACCESS pool, except a physical one (0x189; 1.4.3 also 0x18c, CPU-visible) through the
// virtual pools that map it. Either gets device-local memory: 1.4.3's CPU never writes its 0x18c pool's storage
// (unchanged over 3600 frames of the intro), so the guest storage and the GPU copy need no sync.
if ((p->flags & NVN_MEMORY_POOL_FLAGS_GPU_NO_ACCESS) && !(p->flags & NVN_MEMORY_POOL_FLAGS_PHYSICAL)) return 1;
// A pool of zero bytes (NX2navi makes some now and then) holds no address the GPU can use: it gets no record. A record
// would import 0 bytes (invalid usage; the driver refuses the memory) and put an empty range into the index at the
// pool's address, where a lookup of the pool that starts there found the empty one first. vk_pool_fini finds no
// record for it and does nothing.
if (!p->size) return 1;
PoolRec *r = calloc(1, sizeof *r);
if (!r) rt_fatal("vk: out of memory for a pool record");
r->key = p;
r->flags = p->flags;
r->lo = (uintptr_t)p->mem;
r->size = p->size;
int ok;
if (p->flags & NVN_MEMORY_POOL_FLAGS_VIRTUAL) {
r->kind = POOL_VIRT;
r->npages = p->size / NVN_VPOOL_PAGE;
r->pt = calloc(r->npages ? r->npages : 1, sizeof *r->pt);
ok = r->pt != NULL;
} else if (p->flags & NVN_MEMORY_POOL_FLAGS_PHYSICAL) {
r->kind = POOL_PHYS;
ok = alloc_physical(c, r);
} else {
r->kind = POOL_IMPORT;
mem_thaw((uintptr_t)p->mem, p->size); // the import may need write access to the pages
ok = import_pool(c, r, (uintptr_t)p->mem, p->size);
if (!ok) {
// Not fatal for the game: the pool works for the CPU, and vk_mem_resolve reports its addresses as unbacked.
long n = plat_inc32(&n_import_failures);
if (n <= 8) rt_log("vk: pool %p (flags %#x, %#llx bytes) has no GPU buffer (%ld so far)", (void *)p->mem,
p->flags, (unsigned long long)p->size, n);
r->buf = VK_NULL_HANDLE;
ok = 1;
}
}
if (!ok) { free(r->pt); free(r); return 0; }
plat_mutex_lock(&pool_lock);
index_put(r);
if (r->kind == POOL_PHYS) {
if (nphys == cap_phys) {
cap_phys = cap_phys ? cap_phys * 2 : 16;
phys_list = realloc(phys_list, (size_t)cap_phys * sizeof *phys_list);
if (!phys_list) rt_fatal("vk: out of memory for the physical pool list");
}
phys_list[nphys++] = r;
}
plat_mutex_unlock(&pool_lock);
// The imported pages go into the device address table: shaders that read a pointer from memory reach them there.
if (r->kind == POOL_IMPORT && r->mem && r->buf && r->import_hi > r->import_lo)
r->gpa_rec = vk_gpa_map(r->import_lo, r->import_hi, r->buf, r->import_lo);
return 1;
}
// vk_gpa_unmap: the record of another live imported pool that holds the 4 KiB page at a whole (overlapping imports
// of the same guest memory), 0 for none.
static uint32_t gpa_cover(uint64_t a) {
uint32_t rec = 0;
plat_mutex_lock_shared(&pool_lock);
for (int i = pidx_upper(&pool_index, a) - 1; i >= 0 && !rec; i--) {
const PoolRec *o = pool_index.e[i].buf;
if (o->kind == POOL_IMPORT && o->gpa_rec && o->import_lo <= a && o->import_hi >= a + 4096) rec = o->gpa_rec;
if (i && pool_index.maxhi[i - 1] <= a) break;
}
plat_mutex_unlock_shared(&pool_lock);
return rec;
}
static void gpa_retire(void *arg) { destroy_rec(vk_ctx(), arg); }
void vk_pool_fini(pool_t *p) {
VkCtx *c = vk_ctx();
if (!c) return;
pidx_ent e;
plat_mutex_lock(&pool_lock);
int had = pidx_remove(&pool_index, p, &e);
__atomic_add_fetch(&pool_gen, 1, __ATOMIC_RELEASE);
PoolRec *r = had ? e.buf : NULL;
if (r && r->mstate == MS_ON) mirror_forget(r);
if (r && r->kind == POOL_PHYS) {
for (int i = 0; i < nphys; i++)
if (phys_list[i] == r) { phys_list[i] = phys_list[--nphys]; break; }
for (int i = 0; i < pool_index.n; i++) { // virtual pages that still map it become unmapped
PoolRec *v = pool_index.e[i].buf;
if (v->kind != POOL_VIRT) continue;
for (uint64_t k = 0; k < v->npages; k++)
if (v->pt[k].phys == r) {
v->pt[k].phys = NULL;
vk_gpa_clear(v->lo + k * NVN_VPOOL_PAGE, v->lo + (k + 1) * NVN_VPOOL_PAGE); // the buffer goes away
}
}
}
plat_mutex_unlock(&pool_lock);
if (!r) return;
if (r->smap) shadow_forget(r);
// Off the device address table first; the buffer outlives the submissions that may still read the old pages.
if (r->gpa_rec) vk_gpa_unmap(r->import_lo, r->import_hi, r->gpa_rec, gpa_cover);
for (int i = 0; i < r->nvmaps; i++) vk_gpa_unmap(r->vmaps[i].lo, r->vmaps[i].hi, r->vmaps[i].rec, NULL);
free(r->vmaps), r->vmaps = NULL, r->nvmaps = 0;
if (r->kind == POOL_IMPORT) mem_thaw(r->lo, r->size); // the storage goes back to the guest writable
// Images cut from this range go first; a used buffer or memory waits until the GPU is done with it, since the
// guest gets the storage back when this returns.
vk_texture_range_gone(r->lo, r->size, r->alloc);
if (r->used && r->kind != POOL_VIRT) {
vk_upload_flush(1);
vk_queue_lock();
vkQueueWaitIdle(c->gfx_queue);
vk_queue_unlock();
}
vk_mem_collect();
// A mirrored pool's record and buffers outlive the submission being recorded, whose sync copies and write-backs may
// name them (the game frees a pool only once the GPU is done with it, but its CPU writes may still be queued).
if (r->mbuf) mirror_retiring(r), vk_completion_add(vk_sync_issued() + 1, mirror_retire, r);
else if (r->gpa_rec && !r->used) vk_completion_add(vk_sync_issued() + 1, gpa_retire, r); // the queue was not drained
// an unused physical pool may still be read through the virtual pages that mapped it (device addresses, which do not
// mark it used): its buffer outlives the submission being recorded too
else if (r->kind == POOL_PHYS && !r->used) vk_completion_add(vk_sync_issued() + 1, gpa_retire, r);
else destroy_rec(c, r);
}
// pool_lock held: the device address table follows a MapVirtual request on [lo, hi) of virtual pool v. Ranges it
// covers whole leave the table (their records retire); the pages then name the physical pool's buffer at poff, or none.
static void gpa_map_virtual(PoolRec *v, uint64_t lo, uint64_t hi, PoolRec *ph, uint64_t poff) {
if (!vk_gpa_wanted()) return;
for (int i = 0; i < v->nvmaps;)
if (v->vmaps[i].lo >= lo && v->vmaps[i].hi <= hi) {
vk_gpa_unmap(v->vmaps[i].lo, v->vmaps[i].hi, v->vmaps[i].rec, NULL);
v->vmaps[i] = v->vmaps[--v->nvmaps];
} else {
i++;
}
vk_gpa_clear(lo, hi);
if (!ph || !ph->buf) return;
uint32_t rec = vk_gpa_map(lo, hi, ph->buf, lo - poff);
if (!rec) return;
if (v->nvmaps == v->cvmaps) grow((void **)&v->vmaps, &v->cvmaps, sizeof *v->vmaps);
v->vmaps[v->nvmaps++] = (struct GpaMap){lo, hi, rec};
}
// Requests were checked by the front end (nvn_api_core.c); apply them in order.
int vk_pool_map_virtual(pool_t *v, int n, const void *requests) {
const nvn_mapping_t *rq = requests;
plat_mutex_lock(&pool_lock);
const pidx_ent *e = pidx_find(&pool_index, (uintptr_t)v->mem, 0);
PoolRec *r = e && e->key == v ? e->buf : NULL;
int ok = r && r->kind == POOL_VIRT;
__atomic_add_fetch(&pool_gen, 1, __ATOMIC_RELEASE);
for (int i = 0; ok && i < n; i++) {
PoolRec *ph = rq[i].physical_pool ? phys_find((const pool_t *)(const void *)rq[i].physical_pool) : NULL;
if (rq[i].physical_pool && !ph) {
rt_log("vk: MapVirtual request %d names physical pool %p that has no Vulkan memory", i,
(const void *)rq[i].physical_pool);
ok = 0;
break;
}
uint64_t first = (uint64_t)rq[i].virtual_offset / NVN_VPOOL_PAGE, count = rq[i].size / NVN_VPOOL_PAGE;
if (first + count > r->npages) { ok = 0; break; }
for (uint64_t k = 0; k < count; k++)
r->pt[first + k] = (VPage){ph, ph ? (uint64_t)rq[i].physical_offset + k * NVN_VPOOL_PAGE : 0};
gpa_map_virtual(r, r->lo + first * NVN_VPOOL_PAGE, r->lo + (first + count) * NVN_VPOOL_PAGE, ph,
(uint64_t)rq[i].physical_offset);
}
plat_mutex_unlock(&pool_lock);
if (!ok) rt_log("vk: MapVirtual on pool %p failed", (void *)v->mem);
return ok;
}
// Per-thread copies of resolved import and physical pools that overlap no other index entry (so pidx_find picks them for
// any range inside), valid while pool_gen is unchanged: the replay resolves a few dozen pools thousands of times per
// frame (16 vertex streams and every buffer binding of each draw). Direct-mapped by the 4 KiB page of the address, so
// a hit is one compare however many pools a frame touches. NX2NAVI_RESOLVE_CHECK=1 compares every hit with the locked
// lookup and stops on a difference.
static int rec_class(const PoolRec *r) {
if (r->kind != POOL_IMPORT) return VK_MCLS_DEVICE;
return r->flags & NVN_MEMORY_POOL_FLAGS_CPU_NO_ACCESS ? VK_MCLS_HOST_NOACC
: r->flags & NVN_MEMORY_POOL_FLAGS_CPU_CACHED ? VK_MCLS_HOST_CACHED : VK_MCLS_HOST_UNC;
}
typedef struct { uint64_t gen, lo, hi, base, end; VkBuffer buf; VmaAllocation alloc; int cls; void *pool; } ResolveHit;
#define RHIT_SLOTS 512
static _Thread_local ResolveHit rhit[RHIT_SLOTS];
static ResolveHit *rhit_slot(uint64_t addr) { return &rhit[(addr >> 12 ^ addr >> 21) & (RHIT_SLOTS - 1)]; }
static int resolve_check_on(void) {
static int on = -1;
if (on < 0) {
const char *e = cfg_get("NX2NAVI_RESOLVE_CHECK");
on = e && *e && strcmp(e, "0");
}
return on;
}
static int resolve_locked(uint64_t addr, uint64_t size, VkMemSpan *out);
int vk_mem_resolve(uint64_t addr, uint64_t size, VkMemSpan *out) {
memset(out, 0, sizeof *out);
uint64_t g = __atomic_load_n(&pool_gen, __ATOMIC_ACQUIRE), end = addr + size;
const ResolveHit *h = rhit_slot(addr);
if (end >= addr && h->gen == g && addr >= h->lo && end <= h->hi) {
out->buffer = h->buf;
out->offset = addr - h->base;
out->avail = h->end - addr;
out->alloc = h->alloc;
out->cls = h->cls;
out->pool = h->pool;
if (h->alloc) out->alloc_offset = out->offset;
if (shadow_on && h->pool && !((PoolRec *)h->pool)->stouched) shadow_touch(h->pool);
if (resolve_check_on()) {
VkMemSpan s;
if (!resolve_locked(addr, size, &s) || memcmp(&s, out, sizeof s))
rt_fatal("vk: resolve cache of 0x%llx differs from the pool index", (unsigned long long)addr);
}
return 1;
}
return resolve_locked(addr, size, out);
}
// The pool index lookup under the lock; a non-overlapped import or physical pool goes into the address's cache slot.
static int resolve_locked(uint64_t addr, uint64_t size, VkMemSpan *out) {
memset(out, 0, sizeof *out);
plat_mutex_lock_shared(&pool_lock);
const pidx_ent *e = pidx_find(&pool_index, (uintptr_t)addr, size);
PoolRec *r = e ? e->buf : NULL;
int ok = 0;
if (r && (r->kind == POOL_IMPORT ? r->buf != VK_NULL_HANDLE : r->kind == POOL_PHYS)) {
int i = (int)(e - pool_index.e);
if ((i == 0 || pool_index.maxhi[i - 1] <= e->lo) && (i + 1 == pool_index.n || pool_index.e[i + 1].lo >= e->hi)) {
ResolveHit *h = rhit_slot(addr);
h->gen = __atomic_load_n(&pool_gen, __ATOMIC_RELAXED); // stable: changed only under the exclusive lock
h->lo = e->lo, h->hi = e->hi, h->buf = r->buf, h->cls = rec_class(r);
if (r->kind == POOL_IMPORT) h->base = r->import_lo, h->end = r->import_hi, h->alloc = NULL, h->pool = r;
else h->base = r->lo, h->end = r->lo + r->size, h->alloc = r->alloc, h->pool = NULL;
}
}
if (r && r->kind == POOL_IMPORT && r->buf) {
out->cls = rec_class(r);
out->pool = r;
out->buffer = r->buf;
out->offset = addr - r->import_lo;
out->avail = r->import_hi - addr;
if (r->smap && !r->stouched) shadow_touch(r);
ok = 1;
} else if (r && r->kind == POOL_PHYS) {
out->buffer = r->buf;
out->offset = out->alloc_offset = addr - r->lo;
out->avail = r->size - out->offset;
out->alloc = r->alloc;
out->cls = VK_MCLS_DEVICE;
ok = 1;
} else if (r && r->kind == POOL_VIRT) {
uint64_t rel = addr - r->lo, k = rel / NVN_VPOOL_PAGE;
if (k < r->npages && r->pt[k].phys) {
PoolRec *ph = r->pt[k].phys;
uint64_t off = r->pt[k].off + rel % NVN_VPOOL_PAGE, end = r->pt[k].off + NVN_VPOOL_PAGE;
for (uint64_t j = k + 1; j < r->npages && r->pt[j].phys == ph && r->pt[j].off == end; j++) end += NVN_VPOOL_PAGE;
out->buffer = ph->buf;
out->offset = out->alloc_offset = off;
out->avail = end - off;
out->alloc = ph->alloc;
out->cls = VK_MCLS_DEVICE;
ok = out->avail >= size;
if (ok && !ph->used) plat_xchg32(&ph->used, 1); // read first: the flag is only ever set
}
}
if (ok && !r->used) plat_xchg32(&r->used, 1);
plat_mutex_unlock_shared(&pool_lock);
if (!ok) memset(out, 0, sizeof *out);
return ok;
}
uint64_t vk_mem_gen(void) {
return __atomic_load_n(&pool_gen, __ATOMIC_ACQUIRE) + __atomic_load_n(&mirror_gen, __ATOMIC_ACQUIRE);
}
int vk_mem_class_of(uint64_t addr) {
VkMemSpan m;
return vk_mem_resolve(addr, 1, &m) ? m.cls : VK_MCLS_NONE;
}
// ---------------------------------------------------------------- census
static volatile int64_t traffic[VK_MCLS_N][VK_MUSE_N];
// Every caller is the replay (vk_exec*, under vk_sync_submit's lock), about ten adds per draw: one writer at a time, so
// a plain add, not a locked one. An add racing the census's reset may be lost.
void vk_mem_count(int cls, int use, uint64_t bytes) {
if ((unsigned)cls < VK_MCLS_N && (unsigned)use < VK_MUSE_N)
__atomic_store_n(&traffic[cls][use], __atomic_load_n(&traffic[cls][use], __ATOMIC_RELAXED) + (int64_t)bytes,
__ATOMIC_RELAXED);
}
static void flag_names(uint32_t f, char *buf, size_t n) {
static const char *const names[] = {"CPU_NO_ACCESS", "CPU_UNCACHED", "CPU_CACHED", "GPU_NO_ACCESS", "GPU_UNCACHED",
"GPU_CACHED", "SHADER_CODE", "COMPRESSIBLE", "PHYSICAL", "VIRTUAL"};
size_t k = 0;
buf[0] = 0;
for (int i = 0; i < 10; i++)
if (f >> i & 1) k += (size_t)snprintf(buf + k, k < n ? n - k : 0, "%s%s", k ? "|" : "", names[i]);
if (!k) snprintf(buf, n, "none");
}
// One line per (kind, flags) group of the live pools, then the bytes the replay bound or copied per frame by pool
// class since the previous census (the first one: since frame 1200). Logged at frame 1800 (in the world from the
// run-game saves), at frame 600, and with NX2NAVI_TRACE=perf every 1800 frames.
static void dirty_frame(uint64_t frame);
static void mirror_log(uint64_t frame);
static void mirror_pool_log(uint64_t frame);
static void shadow_log(uint64_t frame);
static int dead_pending(void);
static void vram_log(uint64_t frame);
static void outside_watch(uint64_t frame);
static volatile uint64_t mirror_frame; // the frame vk_mem_census last saw, for mirror_sync's log lines
void vk_mem_census(uint64_t frame) {
mirror_frame = frame;
static uint64_t window_start;
if ((rt_trace & RT_TRACE_PERF) && frame % 600 == 0) vram_log(frame), mirror_pool_log(frame);
if ((rt_trace & RT_TRACE_PERF) && frame % 60 == 0) outside_watch(frame);
if ((rt_trace & RT_TRACE_PERF) && !vk_mirror_on) dirty_frame(frame);
if ((rt_trace & RT_TRACE_PERF) ? frame % 60 == 0 : frame == 600 || frame == 1800)
mirror_log(frame), shadow_log(frame); // mirror_log also logs the write-watch line
if (frame == 1200) { // loading is over: the first window starts here
for (int c = 0; c < VK_MCLS_N; c++)
for (int u = 0; u < VK_MUSE_N; u++) plat_xchg64(&traffic[c][u], 0);
window_start = frame;
return;
}
// Frame 600 as well: runs on slow machines end before 1800.
if (frame != 600 && (frame % 1800 || !(frame == 1800 || (rt_trace & RT_TRACE_PERF)))) return;
typedef struct { int kind; uint32_t flags; uint64_t n, bytes; } Group;
Group g[64];
int ng = 0;
uint64_t total[4] = {0}, count[4] = {0};
plat_mutex_lock_shared(&pool_lock);
for (int i = 0; i < pool_index.n; i++) {
const PoolRec *r = pool_index.e[i].buf;
int j = 0;
while (j < ng && (g[j].kind != r->kind || g[j].flags != r->flags)) j++;
if (j == ng && ng < 64) g[ng++] = (Group){r->kind, r->flags, 0, 0};
if (j < ng) g[j].n++, g[j].bytes += r->size;
total[r->kind] += r->size, count[r->kind]++;
}
plat_mutex_unlock_shared(&pool_lock);
rt_log("mem: census at frame %llu: import %llu pools %.1f MiB (%d without a GPU buffer), physical %llu pools %.1f MiB, "
"virtual %llu pools %.1f MiB of address space",
(unsigned long long)frame, (unsigned long long)count[POOL_IMPORT], total[POOL_IMPORT] / 1048576.0,
n_import_failures, (unsigned long long)count[POOL_PHYS], total[POOL_PHYS] / 1048576.0,
(unsigned long long)count[POOL_VIRT], total[POOL_VIRT] / 1048576.0);
static const char *const kinds[] = {"", "import", "physical", "virtual"};
for (int j = 0; j < ng; j++) {
char f[160];
flag_names(g[j].flags, f, sizeof f);
rt_log("mem: %s %#x %s: %llu pools, %.1f MiB", kinds[g[j].kind], g[j].flags, f, (unsigned long long)g[j].n,
g[j].bytes / 1048576.0);
}
uint64_t frames = window_start && frame > window_start ? frame - window_start : frame;
static const char *const cls[] = {"", "host CPU_NO_ACCESS", "host CPU_UNCACHED", "host CPU_CACHED", "device", "scratch"};
rt_log("mem: traffic per frame over %llu frames, MiB (bound ranges, not bytes read): class: vertex index uniform "
"storage indirect copy-src gpu-write cpu-read", (unsigned long long)frames);
for (int c = 1; c < VK_MCLS_N; c++) {
double v[VK_MUSE_N];
for (int u = 0; u < VK_MUSE_N; u++) v[u] = plat_xchg64(&traffic[c][u], 0) / 1048576.0 / (double)(frames ? frames : 1);
rt_log("mem: %s: %.2f %.2f %.2f %.2f %.3f %.2f %.2f %.3f", cls[c], v[0], v[1], v[2], v[3], v[4], v[5], v[6], v[7]);
}
window_start = frame;
}
// NX2NAVI_VRAM_BUDGET_MB (test; unset or 0: the device's): the device-local budget every decision here and in
// vk_texture.c compares with is at most this many MiB, so a large card behaves as a small one does under pressure
// (mirrors refused or evicted, textures dropped). The heap lines still log what the driver reports.
static uint64_t budget_cap(void) {
static int64_t mb = -1;
if (mb < 0) {
const char *s = cfg_get("NX2NAVI_VRAM_BUDGET_MB");
int64_t v = s && *s ? atoll(s) : 0;
if (v > 0) rt_log("mem: the VRAM budget is taken as at most %lld MiB (NX2NAVI_VRAM_BUDGET_MB, a test)", (long long)v);
mb = v > 0 ? v : 0;
}
return (uint64_t)mb << 20;
}
static int vram_query(uint64_t *used, uint64_t *budget, int *exact) {
VkCtx *c = vk_ctx();
*used = *budget = 0, *exact = 0;
if (!c) return 0;
const VkPhysicalDeviceMemoryProperties *mp = &c->mem_props;
if (c->has_memory_budget) {
VkPhysicalDeviceMemoryBudgetPropertiesEXT b = {.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_BUDGET_PROPERTIES_EXT};
VkPhysicalDeviceMemoryProperties2 p2 = {.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_PROPERTIES_2, .pNext = &b};
vkGetPhysicalDeviceMemoryProperties2(c->phys, &p2);
for (uint32_t h = 0; h < p2.memoryProperties.memoryHeapCount && h < VK_MAX_MEMORY_HEAPS; h++)
if (p2.memoryProperties.memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT)
*used += b.heapUsage[h], *budget += b.heapBudget[h];
*exact = 1;
return *budget != 0;
}
VmaBudget vb[VK_MAX_MEMORY_HEAPS];
vmaGetHeapBudgets(c->vma, vb);
for (uint32_t h = 0; h < mp->memoryHeapCount && h < VK_MAX_MEMORY_HEAPS; h++)
if (mp->memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT)
*used += vb[h].statistics.blockBytes, *budget += mp->memoryHeaps[h].size;
return *budget != 0;
}
int vk_mem_vram(uint64_t *used, uint64_t *budget, int *exact) {
int ok = vram_query(used, budget, exact);
uint64_t cap = budget_cap();
if (ok && cap && *budget > cap) *budget = cap;
return ok;
}
// vk_mem_vram's use less the free space inside VMA's device-local blocks: the driver counts a block whole for as long
// as one allocation lives in it, and the next allocation of its type takes a hole first. The evictions measure with
// this: a texture or mirror freed in the middle of a block left the driver's figure where it was, so an eviction that
// waited for that figure to fall went on until nearly every idle texture and mirror was gone (16.2 GiB in use, 9.2 GiB
// of it free inside blocks after the evictions).
int vk_mem_vram_live(uint64_t *used, uint64_t *budget) {
int exact;
if (!vk_mem_vram(used, budget, &exact)) return 0;
VkCtx *c = vk_ctx();
VmaBudget vb[VK_MAX_MEMORY_HEAPS];
vmaGetHeapBudgets(c->vma, vb);
uint64_t slack = 0;
for (uint32_t h = 0; h < c->mem_props.memoryHeapCount && h < VK_MAX_MEMORY_HEAPS; h++)
if ((c->mem_props.memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT) &&
vb[h].statistics.blockBytes > vb[h].statistics.allocationBytes)
slack += vb[h].statistics.blockBytes - vb[h].statistics.allocationBytes;
*used = *used > slack ? *used - slack : 0;
return 1;
}
// ---------------------------------------------------------------- page write generations
// Two-level table over the guest address space (below 2^39): 256 MiB chunks of 65,536 page generations, allocated
// when a page of the chunk is first stamped. 0 = never stamped.
#define GEN_CHUNK_SHIFT 28
#define GEN_PAGE_SHIFT 12
#define GEN_CHUNKS (RT_ADDRESS_LIMIT >> GEN_CHUNK_SHIFT)
#define GEN_PER_CHUNK (1u << (GEN_CHUNK_SHIFT - GEN_PAGE_SHIFT))
static uint32_t *gen_chunk[GEN_CHUNKS];
static uint32_t gen_now; // last generation handed out
static PlatMutex gen_lock = PLAT_MUTEX_INIT;
static uint64_t gen_pages[1024];
// The newest generation stamped on any page of each 2 MiB region (the write-watch snapshot regions below): the mirror
// sync skips regions with nothing newer than its pool's mirror.
#define REG_SHIFT 21
static uint32_t reg_max[RT_ADDRESS_LIMIT >> REG_SHIFT];
// The same two tables without the stamps of a mirror's own write-back (vk_gpu_written_back): what the CPU and the host
// paths wrote. The mirror sync reads these, so the bytes a mirror wrote back into guest memory are not copied up into it
// again at its next sync: the storage buffers the game's compute shaders write went down and came back up every frame,
// 20 MiB a frame over PCIe in Kokiri Forest (2026-11-09).
static uint32_t *cgen_chunk[GEN_CHUNKS];
static uint32_t creg_max[RT_ADDRESS_LIMIT >> REG_SHIFT];
// Every write-watch query of the generation table (gen_lock held exclusively): calls, bytes and time since the last
// mirror summary, all consumers together; the queries that found written pages (they also flush TLBs) apart.
static uint64_t ww_calls, ww_bytes, ww_ticks, ww_dirty_calls, ww_dirty_ticks, ww_pages;
static uint64_t ww_query(uint64_t addr, uint64_t size) {
uint64_t t0 = qpc();
uint64_t n = mem_written_pages(addr, size, gen_pages, sizeof gen_pages / sizeof *gen_pages);
uint64_t dt = qpc() - t0;
ww_ticks += dt, ww_calls++, ww_bytes += size;
if (n && n != UINT64_MAX) ww_dirty_calls++, ww_dirty_ticks += dt, ww_pages += n;
return n;
}
// Frozen regions: a 2 MiB region whose pages nothing wrote for freeze_idle submissions is made read-only
// (memory.c mem_freeze), and write-watch queries skip it until a write or a file read thaws it. Most of the pool memory
// the mirror and the texture checks cover is static (meshes, textures), and a query costs about 1 us per MiB (4-5 on an
// 8th-gen i7), so this keeps the queries to the regions the CPU writes. The region need not lie wholly in pools: NVN
// pools are small and not 2 MiB aligned, so most regions also hold other guest memory, and a rule that refused those
// left 340 of 436 queried MiB a frame unfrozen in town although nothing had written them for 300 submissions. The
// other memory is safe to freeze for the same reason pools are: the only kernel writes into guest memory are file
// reads, which thaw their buffer first (fs_host.c, romfs.c), and every user-mode write faults and thaws.
#define FREEZE_PER_SUB 32
static int freeze_idle; // 0: off
static uint32_t ww_ep = 1; // gen_lock: the submission being recorded (vk_mem_mirror_begin)
static uint32_t reg_last[RT_ADDRESS_LIMIT >> REG_SHIFT]; // gen_lock: ww_ep of the region's last stamp
static uint32_t reg_tried[RT_ADDRESS_LIMIT >> REG_SHIFT]; // gen_lock: ww_ep of its last freeze candidacy
static uint32_t *fcand; // gen_lock: regions to try at the next freeze pass
static int nfcand, cfcand;
static uint64_t n_frz_refused; // gen_lock: candidates mem_freeze refused
static uint32_t *slot_in(uint32_t **tab, uint64_t page_addr, int create) {
uint64_t ci = page_addr >> GEN_CHUNK_SHIFT;
if (ci >= GEN_CHUNKS) return NULL;
if (!tab[ci]) {
if (!create) return NULL;
tab[ci] = calloc(GEN_PER_CHUNK, sizeof(uint32_t));
if (!tab[ci]) rt_fatal("vk: out of memory for page generations");
}
return &tab[ci][(page_addr >> GEN_PAGE_SHIFT) & (GEN_PER_CHUNK - 1)];
}
static uint32_t *gen_slot(uint64_t page_addr, int create) { return slot_in(gen_chunk, page_addr, create); }
// gen_lock held exclusively; g is the newest generation handed out. back: a mirror's write-back, which leaves the
// CPU-side tables alone.
static void stamp_page_as(uint64_t a, uint32_t g, int back) {
uint32_t *s = gen_slot(a, 1);
if (!s) return;
*s = g;
reg_max[a >> REG_SHIFT] = g;
reg_last[a >> REG_SHIFT] = ww_ep;
if (back) return;
uint32_t *c = slot_in(cgen_chunk, a, 1);
if (c) *c = g, creg_max[a >> REG_SHIFT] = g;
}
static void stamp_page(uint64_t a, uint32_t g) { stamp_page_as(a, g, 0); }
static void stamp_range(uint64_t lo, uint64_t hi, uint32_t g) {
for (uint64_t a = lo; a < hi; a += 1ull << GEN_PAGE_SHIFT) stamp_page(a, g);
}
void vk_guest_written(uint64_t addr, uint64_t size) {
if (!size || nvn_gpu_only(addr)) return;
uint64_t lo = addr & ~0xfffull, hi = (addr + size + 0xfff) & ~0xfffull;
plat_mutex_lock(&gen_lock);
stamp_range(lo, hi, ++gen_now);
plat_mutex_unlock(&gen_lock);
}
int vk_mem_changed(uint64_t addr, uint64_t size, uint32_t since, uint32_t *now) {
if (!size) { *now = since; return 0; }
uint64_t lo = addr & ~0xfffull, hi = (addr + size + 0xfff) & ~0xfffull;
plat_mutex_lock(&gen_lock);
uint32_t g = ++gen_now;
int tracked = 1, frozen = 1;
for (uint64_t a = lo & ~((1ull << REG_SHIFT) - 1); a < hi && frozen; a += 1ull << REG_SHIFT) frozen = mem_frozen(a);
while (!frozen) {
uint64_t n = ww_query(lo, hi - lo);
if (n == UINT64_MAX) { tracked = 0; break; }
for (uint64_t k = 0; k < n; k++) stamp_page(gen_pages[k], g);
if (n < sizeof gen_pages / sizeof *gen_pages) break; // a full answer may have left pages for the next call
}
int changed = 0;
for (uint64_t a = lo; tracked && a < hi && !changed; a += 1ull << GEN_PAGE_SHIFT) {
uint32_t *s = gen_slot(a, 0);
if (s && *s > since) changed = 1;
}
plat_mutex_unlock(&gen_lock);
*now = g;
return tracked ? changed : -1;
}
// Write-watch snapshots per submission: the guest address space in 2 MiB regions, each queried at most once per epoch
// (a guest submission), at the first check that covers it. A query costs about 1.3 us plus 0.55 us per MiB, so one per
// region instead of one per texture. The snapshot only adds observations to the generation table: a page written
// before it carries a generation newer than any check before it, and a page written after it keeps its OS bit until
// the next query over it (at the latest the next epoch's snapshot).
#define SNAP_SHIFT 21
#define SNAP_REGIONS (RT_ADDRESS_LIMIT >> SNAP_SHIFT)
static uint32_t snap_at[SNAP_REGIONS]; // epoch of the region's last snapshot; 0 never
static uint8_t snap_untracked[SNAP_REGIONS];
static uint64_t n_snap_checks, n_snap_calls, n_snap_untracked;
// gen_lock held exclusively. Queries the regions r0..r1 that were not snapped in epoch, one query per run of them;
// returns the number of queries.
static uint64_t snap_regions(uint64_t r0, uint64_t r1, uint32_t epoch) {
uint32_t g = 0;
uint64_t calls = 0;
for (uint64_t r = r0; r <= r1;) {
if (snap_at[r] == epoch) { r++; continue; }
if (mem_frozen(r << SNAP_SHIFT)) { // nothing in it changed since its freeze
snap_at[r] = epoch, snap_untracked[r] = 0;
r++;
continue;
}
uint64_t k = r; // one query over a run of regions not seen in this epoch
while (k + 1 <= r1 && snap_at[k + 1] != epoch && !mem_frozen((k + 1) << SNAP_SHIFT)) k++;
uint64_t a = r << SNAP_SHIFT, b = (k + 1) << SNAP_SHIFT;
int tracked = 1;
if (!g) g = ++gen_now;
for (;;) {
uint64_t n = ww_query(a, b - a);
if (n == UINT64_MAX) { tracked = 0; break; }
for (uint64_t j = 0; j < n; j++) stamp_page(gen_pages[j], g);
if (n < sizeof gen_pages / sizeof *gen_pages) break;
}
calls++;
for (uint64_t j = r; j <= k; j++) {
snap_at[j] = epoch, snap_untracked[j] = (uint8_t)!tracked;
if (tracked && freeze_idle && ww_ep - reg_last[j] >= (uint32_t)freeze_idle &&
ww_ep - reg_tried[j] >= (uint32_t)freeze_idle) {
reg_tried[j] = ww_ep;
if (nfcand == cfcand) grow((void **)&fcand, &cfcand, sizeof *fcand);
fcand[nfcand++] = (uint32_t)j;
}
}
r = k + 1;
}
return calls;
}
int vk_mem_changed_snap(uint64_t addr, uint64_t size, uint32_t since, uint32_t epoch, uint32_t *now) {
if (!size) { *now = since; return 0; }
uint64_t lo = addr & ~0xfffull, hi = (addr + size + 0xfff) & ~0xfffull;
uint64_t r0 = lo >> SNAP_SHIFT, r1 = (hi - 1) >> SNAP_SHIFT;
if (!epoch || r1 >= SNAP_REGIONS) return vk_mem_changed(addr, size, since, now);
int excl = 0;
plat_mutex_lock_shared(&gen_lock);
for (uint64_t r = r0; r <= r1; r++)
if (snap_at[r] != epoch) { excl = 1; break; }
if (excl) {
plat_mutex_unlock_shared(&gen_lock);
plat_mutex_lock(&gen_lock);
n_snap_calls += snap_regions(r0, r1, epoch);
}
int untracked = 0, changed = 0;
for (uint64_t r = r0; r <= r1; r++) untracked |= snap_untracked[r];
// Pages only in regions stamped after since (reg_max; REG_SHIFT is SNAP_SHIFT): an unchanged texture of 4 MiB was
// 1,024 page reads, at about 1,100 texture checks a frame.
for (uint64_t r = r0; !untracked && r <= r1 && !changed; r++) {
if (reg_max[r] <= since) continue;
uint64_t a = r << REG_SHIFT > lo ? r << REG_SHIFT : lo, b = (r + 1) << REG_SHIFT < hi ? (r + 1) << REG_SHIFT : hi;
for (; a < b && !changed; a += 1ull << GEN_PAGE_SHIFT) {
uint32_t *s = gen_slot(a, 0);
if (s && *s > since) changed = 1;
}
}
*now = gen_now; // every observation so far is in the table
__atomic_add_fetch(&n_snap_checks, 1, __ATOMIC_RELAXED);
if (untracked) __atomic_add_fetch(&n_snap_untracked, 1, __ATOMIC_RELAXED);
if (excl) plat_mutex_unlock(&gen_lock);
else plat_mutex_unlock_shared(&gen_lock);
return untracked ? -1 : changed;
}
void vk_mem_snap_stats(uint64_t *checks, uint64_t *calls, uint64_t *untracked) {
*checks = __atomic_exchange_n(&n_snap_checks, 0, __ATOMIC_RELAXED);
*calls = __atomic_exchange_n(&n_snap_calls, 0, __ATOMIC_RELAXED);
*untracked = __atomic_exchange_n(&n_snap_untracked, 0, __ATOMIC_RELAXED);
}
// ---------------------------------------------------------------- dirty census (NX2NAVI_TRACE=perf)
// What a device-local mirror of the buffer-bound pools would upload. The replay notes every
// imported pool it binds as a vertex, index, uniform, storage or indirect buffer (vk_mem_bound_note). At the start of
// each submit one write-watch query per run of such pools stamps the pages written since with a new generation, in the
// shared generation table (no second consumer of the OS bits: texture checks see the stamps), and the pages of those
// pools are counted:
// - eager: pages stamped after the previous submit's scan, in every buffer-bound pool (upload at every submit);
// - lazy: pages stamped after the pool's previous count, in pools the previous submit bound (upload when bound);
// - first use: the size of a pool the first time it is bound (a mirror copies the whole pool then).
// Classes: CPU_CACHED, CPU_UNCACHED, other (CPU_NO_ACCESS). Logged as `mem: dirty` every 60 frames.
#define DIRTY_UNSYNCED UINT32_MAX
#define DIRTY_GAP 0x100000ull // pools closer than this share one query (the gap's pages are stamped too)
#define DIRTY_WIN 60
enum { DC_CACHED, DC_UNC, DC_OTHER, DC_N };
typedef struct {
uint64_t eager[DC_N], lazy[DC_N], first[DC_N];
uint64_t submits, queries, scan_ns, sub_max, frames;
} DirtyAcc;
static uint32_t dirty_sub = 1, dirty_gen; // submit lock: the current submit number, generation of the last scan
static uint64_t first_pending[DC_N], ever_n[DC_N], ever_bytes[DC_N]; // submit lock
static PoolRec **dlist; // submit lock: the buffer-bound pools of a scan
static int dcap;
static PlatMutex dirty_lock = PLAT_MUTEX_INIT; // everything below
static DirtyAcc dfr, dwin; // this frame, this window
static uint32_t dwin_kib[DIRTY_WIN][DC_N]; // eager KiB of each frame of the window
static uint64_t live_n[DC_N], live_bytes[DC_N], live_untracked, ever_snap_n[DC_N], ever_snap_bytes[DC_N];
static int dc_class(const PoolRec *r) {
int c = rec_class(r);
return c == VK_MCLS_HOST_CACHED ? DC_CACHED : c == VK_MCLS_HOST_UNC ? DC_UNC : DC_OTHER;
}
void vk_mem_bound_note(void *pool) {
PoolRec *r = pool;
if (r->bound_sub == dirty_sub) return;
r->bound_sub = dirty_sub;
if (r->bufbound) return;
r->bufbound = 1, r->synced_gen = DIRTY_UNSYNCED;
int k = dc_class(r);
ever_n[k]++, ever_bytes[k] += r->size, first_pending[k] += r->size;
}
static uint64_t page_lo(const PoolRec *r) { return r->lo & ~0xfffull; }
static uint64_t page_hi(const PoolRec *r) { return (r->lo + r->size + 0xfff) & ~0xfffull; }
// gen_lock held exclusively. 0: the range is not write-watched.
static int query_stamp(uint64_t lo, uint64_t hi, uint32_t g) {
for (;;) {
uint64_t n = ww_query(lo, hi - lo);
if (n == UINT64_MAX) return 0;
for (uint64_t k = 0; k < n; k++) stamp_page(gen_pages[k], g);
if (n < sizeof gen_pages / sizeof *gen_pages) return 1;
}
}
// ---------------------------------------------------------------- shadow storage
// Where importing guest memory costs too much (Linux: re-protecting pages that amdgpu imported as a userptr
// invalidates the buffer, about 1.3 ms per 16 MiB at the next submit, and the import's write pinning clears the
// write watch), a CPU-visible pool gets a host-visible buffer of its own instead, at the offsets the import would have:
// - CPU -> shadow: right before each submit (vk_mem_shadow_sync), the pools the replay resolved since the last one are
// queried for written pages; pages stamped after a pool's sgen are copied in by the CPU.
// - GPU -> guest: the replay notes each GPU write into CPU-visible memory at record time (vk_mem_gpu_pending) and
// reports it at the fence (vk_gpu_written), which copies those bytes back into guest memory. Until then the CPU
// copies skip them: a page the CPU and the GPU both write keeps the GPU's bytes.
// NX2NAVI_SHADOW=1/0 forces the mode on or off; the default is on for Linux and Wine.
typedef struct { uint64_t lo, hi; } SRange;
static PlatMutex shadow_lock = PLAT_MUTEX_INIT; // after pool_lock; before gen_lock
static PoolRec **stouch; // shadow_lock: pools resolved since the last sync
static int nstouch, cstouch;
static SRange *spend; // shadow_lock: GPU writes recorded and not copied back yet
static int nspend, cspend;
static int64_t s_live_n, s_live_bytes; // shadow_lock
static uint64_t s_syncs, s_copied, s_queries, s_untracked, s_wb_bytes, s_ticks; // shadow_lock: since the last log
static int shadow_pool(VkCtx *c, PoolRec *r, uint64_t addr, uint64_t size) {
r->import_lo = addr & ~0xfffull;
r->import_hi = (addr + size + 0xfff) & ~0xfffull;
uint64_t n = r->import_hi - r->import_lo;
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .size = n, .usage = pool_usage(c),
.sharingMode = VK_SHARING_MODE_EXCLUSIVE};
// Cached where the device has it: the fence copies read it back on the CPU.
VmaAllocationCreateInfo ai = {.flags = VMA_ALLOCATION_CREATE_MAPPED_BIT,
.requiredFlags = VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
.preferredFlags = VK_MEMORY_PROPERTY_HOST_CACHED_BIT};
VmaAllocationInfo info;
VkResult res = vmaCreateBuffer(c->vma, &bi, &ai, &r->buf, &r->salloc, &info);
if (res != VK_SUCCESS) {
rt_log("vk: pool %#llx+%#llx: shadow allocation failed (%d)", (unsigned long long)addr, (unsigned long long)size,
(int)res);
r->buf = VK_NULL_HANDLE, r->salloc = NULL;
return 0;
}
r->smap = info.pMappedData;
plat_mutex_lock(&gen_lock);
r->sgen = gen_now; // the copy below holds every write stamped so far; later stamps are copied again
plat_mutex_unlock(&gen_lock);
memcpy(r->smap, (const void *)(uintptr_t)r->import_lo, n);
plat_mutex_lock(&shadow_lock);
s_live_n++, s_live_bytes += (int64_t)n;
plat_mutex_unlock(&shadow_lock);
plat_inc32(&n_imports);
return 1;
}
static void shadow_touch(PoolRec *r) {
plat_mutex_lock(&shadow_lock);
if (!r->stouched) {
if (nstouch == cstouch) grow((void **)&stouch, &cstouch, sizeof *stouch);
stouch[nstouch++] = r;
r->stouched = 1;
}
plat_mutex_unlock(&shadow_lock);
}
// vk_pool_fini, after the pool left the index: off the touched list before its record goes.
static void shadow_forget(PoolRec *r) {
plat_mutex_lock(&shadow_lock);
for (int i = 0; i < nstouch; i++)
if (stouch[i] == r) { stouch[i] = stouch[--nstouch]; break; }
s_live_n--, s_live_bytes -= (int64_t)(r->import_hi - r->import_lo);
plat_mutex_unlock(&shadow_lock);
}
// shadow_lock held: copies guest [a, b) into the shadow, leaving out the bytes of GPU writes not copied back yet
// (spend from index i on).
static void shadow_copy(PoolRec *r, uint64_t a, uint64_t b, int i) {
for (; i < nspend; i++)
if (spend[i].lo < b && spend[i].hi > a) {
if (spend[i].lo > a) shadow_copy(r, a, spend[i].lo, i + 1);
if (spend[i].hi < b) shadow_copy(r, spend[i].hi, b, i + 1);
return;
}
memcpy(r->smap + (a - r->import_lo), (const void *)(uintptr_t)a, b - a);
s_copied += b - a;
}
// shadow_lock and gen_lock held: stamps the pool's written pages and copies those newer than its shadow.
static void shadow_sync_pool(PoolRec *r) {
uint64_t lo = r->import_lo, hi = r->import_hi;
uint32_t g = ++gen_now;
int tracked = 1;
for (uint64_t a = lo; a < hi;) { // one query per run of regions that are not frozen
uint64_t b = (a | ((1ull << REG_SHIFT) - 1)) + 1;
if (b > hi) b = hi;
if (mem_frozen(a)) { a = b; continue; }
while (b < hi && !mem_frozen(b)) b = b + (1ull << REG_SHIFT) < hi ? b + (1ull << REG_SHIFT) : hi;
s_queries++;
if (!query_stamp(a, b, g)) tracked = 0;
a = b;
}
uint32_t since = r->sgen;
r->sgen = gen_now;
if (!tracked) { // nothing tells which pages changed: all of them
s_untracked++;
shadow_copy(r, lo, hi, 0);
return;
}
uint64_t run = 0;
for (uint64_t k = lo >> REG_SHIFT; k <= (hi - 1) >> REG_SHIFT; k++) {
uint64_t a = k << REG_SHIFT, b = a + (1ull << REG_SHIFT);
if (a < lo) a = lo;
if (b > hi) b = hi;
const uint32_t *ch = gen_chunk[a >> GEN_CHUNK_SHIFT]; // a 2 MiB region lies in one chunk
if (reg_max[k] <= since || !ch) {
if (run) shadow_copy(r, run, a, 0), run = 0;
continue;
}
for (uint64_t x = a; x < b; x += 1ull << GEN_PAGE_SHIFT) {
int dirty = ch[(x >> GEN_PAGE_SHIFT) & (GEN_PER_CHUNK - 1)] > since;
if (dirty && !run) run = x;
else if (!dirty && run) shadow_copy(r, run, x, 0), run = 0;
}
}
if (run) shadow_copy(r, run, hi, 0);
}
void vk_mem_shadow_sync(void) {
if (!shadow_on) return;
uint64_t t0 = qpc();
plat_mutex_lock(&shadow_lock);
if (nstouch) {
plat_mutex_lock(&gen_lock);
for (int i = 0; i < nstouch; i++) shadow_sync_pool(stouch[i]), stouch[i]->stouched = 0;
plat_mutex_unlock(&gen_lock);
s_syncs += (uint64_t)nstouch;
nstouch = 0;
}
s_ticks += qpc() - t0;
plat_mutex_unlock(&shadow_lock);
}
void vk_mem_gpu_pending(uint64_t addr, uint64_t size) {
if (!shadow_on || !size) return;
plat_mutex_lock(&shadow_lock);
if (nspend == cspend) grow((void **)&spend, &cspend, sizeof *spend);
spend[nspend++] = (SRange){addr, addr + size};
plat_mutex_unlock(&shadow_lock);
}
// NX2NAVI_WB_NO_REUP (default 1; 0 = off): the pages a VRAM mirror wrote back into guest memory are not copied up into
// it again at its next sync. vk_gpu_written stamps the pages the GPU wrote, at the fence, so the texture cache re-reads
// them; the same stamps made mirror_sync copy them back up into the mirror the write-back had read them from: the
// storage buffers the game's compute shaders write went down and came back up every frame, about 57 of the 60 MiB copied
// up a frame outdoors (RTX 4080 SUPER, 2026-11-08) and 20 MiB a frame in Kokiri Forest (RTX 4090, PCIe 3.0, 2026-11-09).
// A write into a mirror (vk_mem_mirrored, noted at record time) stamps only the texture tables (vk_gpu_written_back);
// mirror_sync reads the CPU-side ones (stamp_page_as), which a CPU write or a host path still stamps.
static int wb_reup_off = -1;
static int wb_no_reup(void) {
if (wb_reup_off < 0) {
int on = rt_knob("WB_NO_REUP", "1") == 1;
rt_log("mem: pages a VRAM mirror wrote back %s (NX2NAVI_WB_NO_REUP)",
on ? "are not copied up into it again" : "are copied up into it again");
__atomic_store_n(&wb_reup_off, on, __ATOMIC_RELEASE);
}
return wb_reup_off;
}
static void gpu_written(uint64_t addr, uint64_t size, int back) {
if (shadow_on && size) {
plat_mutex_lock_shared(&pool_lock);
const pidx_ent *e = pidx_find(&pool_index, addr, size);
PoolRec *r = e ? e->buf : NULL;
plat_mutex_lock(&shadow_lock);
if (r && r->smap && addr >= r->import_lo && addr + size <= r->import_hi) {
mem_thaw(addr, size);
memcpy((void *)(uintptr_t)addr, r->smap + (addr - r->import_lo), size);
s_wb_bytes += size;
}
for (int i = 0; i < nspend; i++)
if (spend[i].lo == addr && spend[i].hi == addr + size) { spend[i] = spend[--nspend]; break; }
plat_mutex_unlock(&shadow_lock);
plat_mutex_unlock_shared(&pool_lock);
}
if (!size || nvn_gpu_only(addr)) return;
uint64_t lo = addr & ~0xfffull, hi = (addr + size + 0xfff) & ~0xfffull;
back = back && wb_no_reup(); // read before gen_lock: its first call logs
plat_mutex_lock(&gen_lock);
uint32_t g = ++gen_now;
for (uint64_t a = lo; a < hi; a += 1ull << GEN_PAGE_SHIFT) stamp_page_as(a, g, back);
plat_mutex_unlock(&gen_lock);
}
void vk_gpu_written(uint64_t addr, uint64_t size) { gpu_written(addr, size, 0); }
void vk_gpu_written_back(uint64_t addr, uint64_t size) { gpu_written(addr, size, 1); }
int vk_mem_mirrored(uint64_t addr, uint64_t size) {
VkMemSpan sp;
if (!vk_mirror_on || !size || !vk_mem_resolve(addr, size, &sp) || !sp.pool) return 0;
const PoolRec *r = sp.pool;
return r->mstate == MS_ON && addr >= r->import_lo && addr + size <= r->import_hi;
}
static void shadow_log(uint64_t frame) {
if (!shadow_on) return;
plat_mutex_lock(&shadow_lock);
rt_log("mem: shadow at frame %llu: %lld pools %.1f MiB; since the last: %llu pool syncs, %llu queries, %.1f MiB "
"copied in (%llu untracked), %.1f MiB written back, %.2f ms, %d GPU writes pending",
(unsigned long long)frame, (long long)s_live_n, s_live_bytes / 1048576.0, (unsigned long long)s_syncs,
(unsigned long long)s_queries, s_copied / 1048576.0, (unsigned long long)s_untracked, s_wb_bytes / 1048576.0,
qpc_ms(s_ticks), nspend);
s_syncs = s_queries = s_copied = s_untracked = s_wb_bytes = s_ticks = 0;
plat_mutex_unlock(&shadow_lock);
}
void vk_mem_dirty_submit(void) {
if (!(rt_trace & RT_TRACE_PERF) || vk_mirror_on) return; // the mirror's own line replaces the census
uint64_t t0 = kern_now_ns();
uint32_t prev = dirty_sub; // the submit whose replay noted the bindings
dirty_sub = prev + 1 ? prev + 1 : 1;
uint64_t eager[DC_N] = {0}, lazy[DC_N] = {0}, ln[DC_N] = {0}, lb[DC_N] = {0}, untracked = 0, queries = 0;
plat_mutex_lock_shared(&pool_lock);
if (dcap < pool_index.n) {
dcap = pool_index.n * 2;
dlist = realloc(dlist, (size_t)dcap * sizeof *dlist);
if (!dlist) rt_fatal("vk: out of memory for the dirty census");
}
int n = 0;
for (int i = 0; i < pool_index.n; i++) {
PoolRec *r = pool_index.e[i].buf;
if (r->kind == POOL_IMPORT && r->bufbound) dlist[n++] = r;
}
plat_mutex_lock(&gen_lock);
uint32_t g = ++gen_now, since = dirty_gen;
dirty_gen = g;
for (int i = 0; i < n;) { // one query per run of pools (sorted by address)
uint64_t lo = page_lo(dlist[i]), hi = page_hi(dlist[i]);
int j = i + 1;
for (; j < n && page_lo(dlist[j]) <= hi + DIRTY_GAP; j++)
if (page_hi(dlist[j]) > hi) hi = page_hi(dlist[j]);
queries++;
int ok = query_stamp(lo, hi, g);
for (int k = i; k < j; k++) {
dlist[k]->untracked = 0;
if (!ok) queries++, dlist[k]->untracked = (uint8_t)!query_stamp(page_lo(dlist[k]), page_hi(dlist[k]), g);
}
i = j;
}
uint64_t covered = 0;
for (int i = 0; i < n; i++) {
PoolRec *r = dlist[i];
int k = dc_class(r);
ln[k]++, lb[k] += r->size;
int fresh = r->synced_gen == DIRTY_UNSYNCED, lz = !fresh && r->bound_sub == prev;
uint32_t ls = r->synced_gen;
if (fresh || r->bound_sub == prev) r->synced_gen = g;
if (r->untracked) { untracked += r->size; continue; }
if (fresh) continue; // counted as first use
uint64_t lo = page_lo(r), hi = page_hi(r), ne = 0, nl = 0;
if (lo < covered) lo = covered; // overlapping pools: each page once
if (lo >= hi) continue;
covered = hi;
for (uint64_t a = lo; a < hi;) {
uint64_t ci = a >> GEN_CHUNK_SHIFT, ce = (ci + 1) << GEN_CHUNK_SHIFT;
if (ce > hi) ce = hi;
const uint32_t *ch = ci < GEN_CHUNKS ? gen_chunk[ci] : NULL;
if (ch)
for (uint64_t x = a; x < ce; x += 1ull << GEN_PAGE_SHIFT) {
uint32_t v = ch[(x >> GEN_PAGE_SHIFT) & (GEN_PER_CHUNK - 1)];
ne += v > since;
nl += lz && v > ls;
}
a = ce;
}
eager[k] += ne << GEN_PAGE_SHIFT, lazy[k] += nl << GEN_PAGE_SHIFT;
}
plat_mutex_unlock(&gen_lock);
plat_mutex_unlock_shared(&pool_lock);
uint64_t tot = 0;
for (int k = 0; k < DC_N; k++) tot += eager[k];
plat_mutex_lock(&dirty_lock);
for (int k = 0; k < DC_N; k++) {
dfr.eager[k] += eager[k], dfr.lazy[k] += lazy[k], dfr.first[k] += first_pending[k];
first_pending[k] = 0;
live_n[k] = ln[k], live_bytes[k] = lb[k], ever_snap_n[k] = ever_n[k], ever_snap_bytes[k] = ever_bytes[k];
}
live_untracked = untracked;
dfr.submits++, dfr.queries += queries;
if (tot > dfr.sub_max) dfr.sub_max = tot;
dfr.scan_ns += kern_now_ns() - t0;
plat_mutex_unlock(&dirty_lock);
}
static void dirty_frame(uint64_t frame) {
plat_mutex_lock(&dirty_lock);
if (dwin.frames < DIRTY_WIN)
for (int k = 0; k < DC_N; k++) dwin_kib[dwin.frames][k] = (uint32_t)(dfr.eager[k] >> 10);
for (int k = 0; k < DC_N; k++) dwin.eager[k] += dfr.eager[k], dwin.lazy[k] += dfr.lazy[k], dwin.first[k] += dfr.first[k];
dwin.submits += dfr.submits, dwin.queries += dfr.queries, dwin.scan_ns += dfr.scan_ns, dwin.frames++;
if (dfr.sub_max > dwin.sub_max) dwin.sub_max = dfr.sub_max;
memset(&dfr, 0, sizeof dfr);
if (frame % DIRTY_WIN) { plat_mutex_unlock(&dirty_lock); return; }
DirtyAcc w = dwin;
uint32_t kib[DIRTY_WIN][DC_N];
memcpy(kib, dwin_kib, sizeof kib);
uint64_t lv[DC_N], lvb[DC_N], ev[DC_N], evb[DC_N], unt = live_untracked;
for (int k = 0; k < DC_N; k++) lv[k] = live_n[k], lvb[k] = live_bytes[k], ev[k] = ever_snap_n[k], evb[k] = ever_snap_bytes[k];
memset(&dwin, 0, sizeof dwin);
plat_mutex_unlock(&dirty_lock);
if (!w.frames) return;
uint64_t fl[2][3][2];
nvn_flush_stats(fl);
double f = (double)w.frames, M = 1048576.0;
uint64_t sub_tot = 0;
for (int k = 0; k < DC_N; k++) sub_tot += w.eager[k];
#define C3(a, d) (a)[0] / (d), (a)[1] / (d), (a)[2] / (d)
rt_log("mem: dirty frame %llu: frames %llu submits_per_frame %.2f; per frame, cached/uncached/other: eager_mib "
"%.3f/%.3f/%.3f lazy_mib %.3f/%.3f/%.3f first_mib %.3f/%.3f/%.3f flush_pool_calls %.1f/%.1f/%.1f "
"flush_pool_mib %.3f/%.3f/%.3f flush_buf_calls %.1f/%.1f/%.1f flush_buf_mib %.3f/%.3f/%.3f; per submit: "
"avg_mib %.3f max_mib %.3f; mirror: live_pools %llu/%llu/%llu live_mib %.1f/%.1f/%.1f ever_pools %llu/%llu/%llu "
"ever_mib %.1f/%.1f/%.1f untracked_mib %.1f; cost: queries_per_frame %.1f scan_ms_per_frame %.3f",
(unsigned long long)frame, (unsigned long long)w.frames, w.submits / f,
C3(w.eager, f * M), C3(w.lazy, f * M), C3(w.first, f * M),
fl[0][0][0] / f, fl[0][1][0] / f, fl[0][2][0] / f, fl[0][0][1] / f / M, fl[0][1][1] / f / M, fl[0][2][1] / f / M,
fl[1][0][0] / f, fl[1][1][0] / f, fl[1][2][0] / f, fl[1][0][1] / f / M, fl[1][1][1] / f / M, fl[1][2][1] / f / M,
w.submits ? sub_tot / (double)w.submits / M : 0.0, w.sub_max / M,
(unsigned long long)lv[0], (unsigned long long)lv[1], (unsigned long long)lv[2], C3(lvb, M),
(unsigned long long)ev[0], (unsigned long long)ev[1], (unsigned long long)ev[2], C3(evb, M), unt / M,
w.queries / f, w.scan_ns / 1e6 / f);
#undef C3
char line[DIRTY_WIN * 36 + 64];
size_t k = (size_t)snprintf(line, sizeof line, "mem: dirty frames KiB cached/uncached/other:");
for (uint64_t i = 0; i < w.frames && i < DIRTY_WIN; i++)
k += (size_t)snprintf(line + k, k < sizeof line ? sizeof line - k : 0, " %u/%u/%u", kib[i][0], kib[i][1], kib[i][2]);
rt_log("%s", line);
}
// ---------------------------------------------------------------- VRAM mirror of buffer-bound pools
// An imported pool keeps vertices, indices, constants
// and storage buffers in system memory, so a discrete GPU reads them over PCIe on every draw (GTX 1050 Ti and RTX
// 3090 Ti: about 110 ms of GPU time per frame in the field). In mirror mode (device class discrete, NX2NAVI_MEMORY) an
// imported pool gets a device-local buffer over its imported range the first time the replay uses it as a vertex,
// index, uniform, storage or indirect buffer (vk_mem_use). From then on those bindings, and the buffer copies, fills
// and uniform updates that touch the pool (vk_mem_pick), use the mirror at the same offsets. Guest memory stays what
// the CPU reads and writes; textures and every host path that reads or writes pool bytes keep using it.
// - CPU -> mirror: the first access of a submission to a mirror syncs it (mirror_sync): the 2 MiB write-watch
// snapshots of its regions for this epoch (snap_regions, shared with the texture checks: the OS bits keep one
// consumer, the shared page generations) and every page the CPU-side tables (cgen_chunk) stamped after the mirror's
// mgen goes into a copy. Where a region of the pool is not write-watched, so do the nvnMemoryPoolFlushMappedRange
// ranges, each on its own: one span from the lowest to the highest flush of a 252 MiB pool copied up to the whole
// pool for a few MiB flushed, 30-50 ms of GPU time over PCIe 3.0 x8 in crowded towns (NX2NAVI_FLUSH_SPAN=1 restores
// it). Where the write watch covers the pool, a flushed page the CPU wrote is a written page already, and the game's
// flushes of whole pools it had not written cost 47-57 MiB of copies a frame in Kokiri Forest (2026-11-09);
// NX2NAVI_FLUSH_UPLOAD=1 copies them anyway. The first use copies the whole range. The copies run in a prologue command
// buffer that vk_sync.c submits right before the part that accessed the mirror first, between full barriers.
// vk_mem_mirror_begin snaps the regions of the pools the previous submission accessed in a few merged queries.
// - GPU -> guest: what the replay writes into a mirror (storage buffers its shaders write, buffer copies, fills,
// uniform updates, texture-to-buffer copies: vk_mem_gpu_wrote) goes back into the imported buffer at the end of the
// submission, exact byte ranges, before the fence signals (vk_mem_mirror_end). At the fence those pages get a new
// generation for the texture checks (vk_gpu_written_back) but not in the CPU-side tables: the mirror holds the bytes.
// - NX2NAVI_MEMORY=verify: every submission also copies up to VERIFY_BYTES of mirror pages, round robin, into host memory;
// at its fence the pages the CPU has not written since their sync are compared with guest memory.
// A pool stays imported when its mirror would not fit the VRAM budget, its range is not write-watched, it overlaps
// another pool or the allocation fails.
#define VERIFY_BYTES (4u << 20)
#define VERIFY_SAMPLES 64
#define VERIFY_LOG 200
enum { MR_BUDGET, MR_UNTRACKED, MR_OVERLAP, MR_ALLOC, MR_N };
#define RETRY_SUBS 600 // a pool refused for the budget asks again this many submissions later (eviction)
int vk_mirror_on;
static int mirror_verify;
static int mirror_gap = 2; // pre-snap: hot pools at most this many 2 MiB regions apart share one query
static PoolRec **mlist; // pool_lock: mirrored pools by import_lo
static int nm, mcap;
static uint32_t msub; // replay: number of the submission being recorded
static uint32_t mepoch; // replay: its write-watch snapshot epoch (vk_texture_epoch)
typedef struct { PoolRec *r; uint64_t off, size; } MRange;
static MRange *pend, *wbl; // replay: copies for the next prologue; GPU writes into mirrors this submission
static int npend, cpend, nwbl, cwbl;
static VkBufferCopy *mregions; // replay: vkCmdCopyBuffer regions
static int cregions;
static PlatMutex flush_lock = PLAT_MUTEX_INIT;
static int flush_cap = FLUSH_RANGES; // ranges a pool keeps; 1 with NX2NAVI_FLUSH_SPAN=1
static int flush_upload; // NX2NAVI_FLUSH_UPLOAD=1: flush ranges are copied where the write watch sees them too
typedef struct {
uint64_t subs, syncs, upload_bytes, upload_copies, first_n, first_bytes, wb_bytes, wb_copies, wb_ranges;
uint64_t presnap_queries, sync_queries, presnap_t, sync_t, prologues, presnap_bytes;
uint64_t snap_bytes, merge_bytes, merge_dispatches, snap_allocs; // write-back merge
uint64_t page_bytes, flush_n, flush_bytes, flush_calls, flush_sum; // what mirror_sync queues, by source
uint64_t flush_skipped; // flush ranges not copied: the write watch covers them
} MirrorAcc;
static MirrorAcc mrep; // replay: this submission
static MirrorAcc macc; // mirror_stat_lock: since the last summary
static PlatMutex mirror_stat_lock = PLAT_MUTEX_INIT;
static volatile int64_t m_live_n, m_live_bytes, m_refused_n[MR_N], m_refused_bytes[MR_N];
static volatile int64_t m_flush_calls, m_flush_bytes, v_pages, v_skipped, v_mismatch;
static void mirror_mode(int discrete) {
const char *e = cfg_value("NX2NAVI_MEMORY");
int mode = discrete ? 1 : 0; // 0 import, 1 mirror, 2 verify
const char *why = discrete ? "auto, discrete gpu" : "auto, uma";
if (strcmp(e, "auto")) {
if (!strcmp(e, "import")) mode = 0, why = "NX2NAVI_MEMORY=import";
else if (!strcmp(e, "mirror")) mode = 1, why = "NX2NAVI_MEMORY=mirror";
else if (!strcmp(e, "verify")) mode = 2, why = "NX2NAVI_MEMORY=verify";
else rt_log("mem: NX2NAVI_MEMORY=%s is not import, mirror, verify or auto: auto", e);
}
vk_mirror_on = mode > 0;
mirror_verify = mode == 2;
const char *g = cfg_get("NX2NAVI_MIRROR_GAP");
mirror_gap = g && *g ? atoi(g) : 2;
if (mirror_gap < 0) mirror_gap = 0;
// Submissions a region stays unwritten before it is frozen. 60 rather than 300: in a fight 100 of the 257 MiB
// queried a frame lay in regions idle for 30-299 submissions, and a thaw by a write costs about 60 us.
if (rt_knob("FLUSH_SPAN", "0")) {
flush_cap = 1;
rt_log("mem: a pool's flushes copy one span from the lowest to the highest (NX2NAVI_FLUSH_SPAN=1)");
}
flush_upload = rt_knob("FLUSH_UPLOAD", "0") == 1;
if (flush_upload) rt_log("mem: flushed ranges are copied into the mirrors as well as written pages (NX2NAVI_FLUSH_UPLOAD=1)");
const char *fi = cfg_get("NX2NAVI_FREEZE_IDLE");
freeze_idle = fi && *fi ? atoi(fi) : 60;
if (freeze_idle < 0) freeze_idle = 0;
const char *sh = cfg_get("NX2NAVI_SHADOW");
#ifdef _WIN32
shadow_on = sh && *sh ? strcmp(sh, "0") != 0 : plat_is_wine();
#else
shadow_on = !sh || !*sh || strcmp(sh, "0");
#endif
if (shadow_on) rt_log("mem: shadow storage: CPU-visible pools are host-visible copies of guest memory, not imports%s",
sh && *sh ? " (NX2NAVI_SHADOW)" : "");
rt_log("mem: mode %s (%s): %s", mode == 2 ? "verify" : mode ? "mirror" : "import", why,
mode ? "pools used as buffers get a device-local mirror" : "every CPU-visible pool is imported host memory");
// Import is the integrated-GPU mode. On a discrete card every draw reads the pools over PCIe (the comment above
// this section has the measurements), which is the difference between playable and 5 fps when it is not intended.
if (!mode && discrete)
rt_log("mem: warning: import on a discrete GPU: the GPU reads the game's pools over PCIe on every draw, far "
"slower than mirror (the auto default on a discrete GPU); use memory = auto unless testing");
}
static void range_add(MRange **a, int *n, int *cap, PoolRec *r, uint64_t off, uint64_t size) {
if (*n && (*a)[*n - 1].r == r && (*a)[*n - 1].off + (*a)[*n - 1].size == off) {
(*a)[*n - 1].size += size;
return;
}
if (*n == *cap) grow((void **)a, cap, sizeof **a);
(*a)[(*n)++] = (MRange){r, off, size};
}
static int range_cmp(const void *x, const void *y) {
const MRange *a = x, *b = y;
if (a->r != b->r) return (uintptr_t)a->r < (uintptr_t)b->r ? -1 : 1;
return a->off < b->off ? -1 : a->off > b->off;
}
// Records the ranges as one vkCmdCopyBuffer per pool, overlapping and adjacent ranges merged; to_mirror: buf -> mbuf,
// else mbuf -> buf. Returns the commands recorded; adds what they move to *bytes and *copies.
static int copy_ranges(VkCommandBuffer cb, MRange *a, int n, int to_mirror, uint64_t *bytes, uint64_t *copies) {
qsort(a, (size_t)n, sizeof *a, range_cmp);
int rec = 0;
for (int i = 0; i < n;) {
PoolRec *r = a[i].r;
uint64_t lim = r->import_hi - r->import_lo;
int k = 0;
for (; i < n && a[i].r == r; i++) {
uint64_t lo = a[i].off, hi = a[i].off + a[i].size;
if (hi > lim) hi = lim;
if (lo >= hi) continue;
if (k && lo <= mregions[k - 1].srcOffset + mregions[k - 1].size) {
uint64_t end = mregions[k - 1].srcOffset + mregions[k - 1].size;
if (hi > end) mregions[k - 1].size += hi - end;
continue;
}
if (k == cregions) grow((void **)&mregions, &cregions, sizeof *mregions);
mregions[k++] = (VkBufferCopy){lo, lo, hi - lo};
}
if (!k) continue;
if (to_mirror) vkCmdCopyBuffer(cb, r->buf, r->mbuf, (uint32_t)k, mregions);
else vkCmdCopyBuffer(cb, r->mbuf, r->buf, (uint32_t)k, mregions);
rec++;
for (int j = 0; j < k; j++) *bytes += mregions[j].size;
*copies += (uint64_t)k;
}
return rec;
}
static void mbarrier(VkCommandBuffer cb, VkPipelineStageFlags2 ss, VkAccessFlags2 sa, VkPipelineStageFlags2 ds,
VkAccessFlags2 da) {
VkMemoryBarrier2 mb = {.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER_2, .srcStageMask = ss, .srcAccessMask = sa,
.dstStageMask = ds, .dstAccessMask = da};
VkDependencyInfo di = {.sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO, .memoryBarrierCount = 1, .pMemoryBarriers = &mb};
vkCmdPipelineBarrier2(cb, &di);
}
static void refuse(PoolRec *r, int why, const char *what) {
r->mstate = MS_REFUSED;
r->retry_sub = why == MR_BUDGET ? (msub + RETRY_SUBS ? msub + RETRY_SUBS : 1) : 0;
uint64_t n = r->import_hi - r->import_lo;
static volatile int32_t logged; // budget refusals are asked again (vk_mem_use): a few lines in all
if (plat_inc64(&m_refused_n[why]) <= 4 && plat_inc32(&logged) <= 12) {
char f[160];
flag_names(r->flags, f, sizeof f);
rt_log("mem: pool %#llx+%#llx (%s) stays imported: %s", (unsigned long long)r->lo, (unsigned long long)r->size,
f, what);
}
plat_add64(&m_refused_bytes[why], (int64_t)n);
}
// ---- eviction (NX2NAVI_MIRROR_EVICT, default on; 0 = a pool over the budget stays imported, as before)
// A pool whose mirror would not fit the VRAM budget stayed imported for good, however hot, while mirrors nothing had
// used for minutes kept their VRAM; draws that read an imported pool over PCIe bring a discrete GPU down to a few
// frames per second. Now a new mirror that does not fit takes the room of mirrors no submission synced for MIRROR_IDLE
// submissions, least recently synced first (mirror_evict), and only when they free enough; when they cannot, idle
// textures are asked to make the rest (vk_texture_want_room) and the pool asks again RETRY_SOON submissions later, once
// per refusal series; after that, and for any other pool refused for the budget, every RETRY_SUBS submissions
// (vk_mem_use). Asked every RETRY_SOON submissions without end, 113 refused pools made 1,795 retries in 6,000 frames. Mirrors never go to make room for textures: evicted
// under VRAM pressure during a menu, 395 of them were refused again when the world came back and the game ran at 4 fps.
// The budget counts the free space inside VMA's blocks as room (vk_mem_vram_live).
// A victim was neither synced, used nor written back in the submission being recorded (it has nothing in the copy,
// write-back or snapshot lists), and every GPU write to it went back to guest memory at the end of the submission that
// made it (vk_mem_mirror_end), so the imported buffer holds everything. The device address table moves its pages back
// to the imported buffer (vk_gpa_remap_plain); a shader of a submission in flight that reads the old record still finds
// the mirror, destroyed only once the submission being recorded is done (vk_defer_buffer). A mirror without a twin
// (twin_wanted) is never a victim: what shaders wrote into it through the table is in the mirror alone. The pool is
// mirrored again, from a whole copy, on its next buffer use. Its freed bytes count as room at once (evict_room), until
// the GPU is past the submission and the destroy shows in VMA's figures.
#define MIRROR_IDLE 120
#define RETRY_SOON 16
static int mirror_evict_on = -1;
static int64_t m_evicted_n, m_evicted_bytes, m_retries;
static uint64_t room_bytes, room_fence; // replay: bytes evicted that the budget may still count, until room_fence
static int evict_enabled(void) {
if (mirror_evict_on < 0) {
mirror_evict_on = rt_knob("MIRROR_EVICT", "1");
rt_log("mem: idle mirrors %s (NX2NAVI_MIRROR_EVICT=%d)", mirror_evict_on ? "make room for new ones and for textures under VRAM pressure"
: "stay; a pool over the VRAM budget stays imported", mirror_evict_on);
}
return mirror_evict_on;
}
static uint64_t evict_room(void) {
if (room_bytes && vk_sync_fence_done() >= room_fence) room_bytes = 0;
return room_bytes;
}
// Replay thread: evicts mirrors idle for at least min_idle submissions, least recently synced first, until need bytes
// are freed; none when the candidates cannot free that much. Returns the bytes freed.
static uint64_t mirror_evict(uint64_t need, uint32_t min_idle, const char *why) {
enum { MAXV = 64 };
PoolRec *v[MAXV];
int nv = 0;
plat_mutex_lock(&pool_lock);
for (int i = 0; i < nm; i++) { // the MAXV least recently synced candidates, oldest first
PoolRec *r = mlist[i];
uint32_t age = msub - r->acc_sub;
if (age < min_idle || r->wb_sub == msub || msub - r->wb_sub < 2) continue;
if (!r->twin && r->gpa_rec) continue; // the table's writes are in the mirror alone
int at = nv;
while (at > 0 && msub - v[at - 1]->acc_sub < age) at--;
if (at == MAXV) continue;
if (nv < MAXV) nv++;
memmove(v + at + 1, v + at, sizeof *v * (size_t)(nv - 1 - at));
v[at] = r;
}
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;
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];
idle_for[n] = msub - r->acc_sub;
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_;
r->mbuf = VK_NULL_HANDLE, r->malloc_ = NULL;
plat_mutex_lock(&flush_lock); // flushes not copied yet: the next mirror starts from a whole copy
__atomic_store_n(&r->fl_n, 0, __ATOMIC_RELEASE);
r->fl_calls = r->fl_sum = 0;
plat_mutex_unlock(&flush_lock);
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);
freed += r->import_hi - r->import_lo;
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);
}
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
if (n) {
plat_add64(&m_evicted_n, n);
plat_add64(&m_evicted_bytes, (int64_t)freed);
room_bytes += freed, room_fence = vk_sync_issued() + 2;
}
return freed;
}
// ---- twins (NX2NAVI_GPA_TWIN_MAX_KB, default 0: every mirror has one)
// Shaders write through the device address table into both a mirror and its pool's imported buffer, guest memory
// (vk_gpa.c, "The twin"): such writes are not recorded as GPU writes, so with the mirror alone the CPU never saw
// them and a sync of a page the CPU wrote copied the stale guest page over them. On a discrete GPU the second write
// crosses PCIe: on PCIe 3.0 the twins cost about 3 ms of GPU time a frame in Kokiri Forest (NX2NAVI_GPA_TWIN=0 against
// on). The writes that went wrong without them were the foliage counts and records in pools of 20 and 44 KiB.
// GPA_TWIN_MAX_KB=n gives a twin only to mirrors of pools up to n KiB; a larger mirror keeps the table's writes to
// itself (its record names the mirror twice, which sends the second write to the table's scratch page) and its
// buffer fills stay in the mirror (vk_mem_guest_twin). Such a mirror is never evicted (mirror_evict), since guest
// memory lacks what the shaders wrote into it. For measuring which pools need a twin: a pool that does and is left
// without draws stretched or flickering foliage.
static int64_t twin_max = -1; // bytes; 0 no limit
static int64_t m_untwinned_n, m_untwinned_bytes;
static int twin_wanted(const PoolRec *r) {
if (twin_max < 0) {
int kb = cfg_int("NX2NAVI_GPA_TWIN_MAX_KB");
twin_max = kb > 0 ? (int64_t)kb << 10 : 0;
if (twin_max)
rt_log("mem: only mirrors of pools up to %d KiB send the device address table's writes to guest memory "
"too (NX2NAVI_GPA_TWIN_MAX_KB); larger ones keep them and are not evicted", kb);
}
return !twin_max || (int64_t)(r->import_hi - r->import_lo) <= twin_max;
}
// Replay thread: the pool's mirror on its first buffer use.
static void mirror_create(PoolRec *r) {
VkCtx *c = vk_ctx();
uint64_t n = r->import_hi - r->import_lo;
r->mstate = MS_REFUSED;
__atomic_add_fetch(&mirror_gen, 1, __ATOMIC_RELEASE);
if (r->kind != POOL_IMPORT || !r->buf || !n) return;
plat_mutex_lock_shared(&pool_lock);
int i = pidx_upper(&pool_index, r->lo) - 1;
while (i >= 0 && pool_index.e[i].buf != r) i--;
int alone = i >= 0 && (i == 0 || pool_index.maxhi[i - 1] <= pool_index.e[i].lo) &&
(i + 1 == pool_index.n || pool_index.e[i + 1].lo >= pool_index.e[i].hi);
plat_mutex_unlock_shared(&pool_lock);
if (i < 0) return; // finalized meanwhile
if (!alone) return refuse(r, MR_OVERLAP, "it overlaps another pool");
uint64_t r0 = r->import_lo >> SNAP_SHIFT, r1 = (r->import_hi - 1) >> SNAP_SHIFT;
if (r1 >= SNAP_REGIONS) return refuse(r, MR_UNTRACKED, "outside the write-watched address space");
uint64_t used, budget;
int exact;
int evict = evict_enabled();
if (evict ? vk_mem_vram_live(&used, &budget) : vk_mem_vram(&used, &budget, &exact)) {
if (evict) {
uint64_t room = evict_room();
used = used > room ? used - room : 0;
}
if (used + n + budget / 8 > budget && evict) { // room from mirrors nothing synced for a while
char why[96];
snprintf(why, sizeof why, "room for a %.1f MiB mirror", n / 1048576.0);
uint64_t f = mirror_evict(used + n + budget / 8 - budget, MIRROR_IDLE, why);
used = used > f ? used - f : 0;
}
if (used + n + budget / 8 > budget) {
char why[128];
snprintf(why, sizeof why, "its %.1f MiB mirror does not fit the VRAM budget (%.0f of %.0f MiB in use)",
n / 1048576.0, used / 1048576.0, budget / 1048576.0);
refuse(r, MR_BUDGET, why);
if (evict && !r->asked_room) { // idle textures make the rest; asked again once they had the time
vk_texture_want_room(used + n + budget / 8 - budget);
r->retry_sub = msub + RETRY_SOON ? msub + RETRY_SOON : 1;
r->asked_room = 1; // later refusals wait RETRY_SUBS: the textures had nothing to give
}
return;
}
}
plat_mutex_lock(&gen_lock);
mrep.sync_queries += snap_regions(r0, r1, mepoch);
int untracked = 0;
for (uint64_t k = r0; k <= r1; k++) untracked |= snap_untracked[k];
r->mgen = gen_now; // the whole-range copy below holds every write observed so far
plat_mutex_unlock(&gen_lock);
if (untracked) return refuse(r, MR_UNTRACKED, "its range is not write-watched");
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .size = n, .usage = pool_usage(c),
.sharingMode = VK_SHARING_MODE_EXCLUSIVE};
VmaAllocationCreateInfo ai = {.usage = VMA_MEMORY_USAGE_AUTO_PREFER_DEVICE,
.requiredFlags = VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT};
VkResult res = vmaCreateBuffer(c->vma, &bi, &ai, &r->mbuf, &r->malloc_, NULL);
if (res != VK_SUCCESS) {
r->mbuf = VK_NULL_HANDLE, r->malloc_ = NULL;
char why[96];
snprintf(why, sizeof why, "device-local allocation of %.1f MiB failed (%d)", n / 1048576.0, (int)res);
return refuse(r, MR_ALLOC, why);
}
plat_mutex_lock(&pool_lock);
if (nm == mcap) grow((void **)&mlist, &mcap, sizeof *mlist);
int at = nm;
while (at > 0 && mlist[at - 1]->import_lo > r->import_lo) at--;
memmove(mlist + at + 1, mlist + at, (size_t)(nm - at) * sizeof *mlist);
mlist[at] = r;
nm++;
plat_mutex_unlock(&pool_lock);
r->acc_sub = msub;
r->mstate = MS_ON;
r->asked_room = 0;
// The device address table names the mirror from now on: the GPU's writes go there, and the sync copies the CPU's.
// With a twin they also go to the imported buffer; without one (twin_wanted) the record names the mirror twice.
r->twin = (uint8_t)twin_wanted(r);
if (!r->twin) {
plat_inc64(&m_untwinned_n);
plat_add64(&m_untwinned_bytes, (int64_t)n);
}
if (r->gpa_rec)
r->gpa_rec = r->twin ? vk_gpa_remap(r->import_lo, r->import_hi, r->gpa_rec, r->mbuf, r->import_lo)
: vk_gpa_remap_plain(r->import_lo, r->import_hi, r->gpa_rec, r->mbuf, r->import_lo);
range_add(&pend, &npend, &cpend, r, 0, n);
mrep.first_n++, mrep.first_bytes += n;
plat_inc64(&m_live_n);
plat_add64(&m_live_bytes, (int64_t)n);
}
static void mirror_forget(PoolRec *r) {
__atomic_add_fetch(&mirror_gen, 1, __ATOMIC_RELEASE);
for (int i = 0; i < nm; i++)
if (mlist[i] == r) {
memmove(mlist + i, mlist + i + 1, (size_t)(nm - i - 1) * sizeof *mlist);
nm--;
plat_dec64(&m_live_n);
plat_add64(&m_live_bytes, -(int64_t)(r->import_hi - r->import_lo));
break;
}
}
// Finalized mirrored pools waiting for mirror_retire; vk_mem_shutdown destroys those whose callback never ran.
static PoolRec **retiring;
static int nretiring, cretiring;
static PlatMutex retire_lock = PLAT_MUTEX_INIT;
static void mirror_retiring(PoolRec *r) {
plat_mutex_lock(&retire_lock);
if (nretiring == cretiring) grow((void **)&retiring, &cretiring, sizeof *retiring);
retiring[nretiring++] = r;
plat_mutex_unlock(&retire_lock);
}
static void mirror_retire(void *arg) {
plat_mutex_lock(&retire_lock);
for (int i = 0; i < nretiring; i++)
if (retiring[i] == arg) { retiring[i] = retiring[--nretiring]; break; }
plat_mutex_unlock(&retire_lock);
destroy_rec(vk_ctx(), arg);
}
// Replay thread, the first access of this submission: queues the copies of the pages written since the last sync.
static void mirror_sync(PoolRec *r) {
uint64_t t0 = qpc();
r->acc_sub = msub;
uint64_t lo = r->import_lo, hi = r->import_hi;
uint64_t r0 = lo >> SNAP_SHIFT, r1 = (hi - 1) >> SNAP_SHIFT;
int excl = 0;
plat_mutex_lock_shared(&gen_lock);
for (uint64_t k = r0; k <= r1; k++)
if (snap_at[k] != mepoch) { excl = 1; break; }
if (excl) {
plat_mutex_unlock_shared(&gen_lock);
plat_mutex_lock(&gen_lock);
mrep.sync_queries += snap_regions(r0, r1, mepoch);
}
uint32_t since = r->mgen;
int untracked = 0;
for (uint64_t k = r0; k <= r1; k++) {
untracked |= snap_untracked[k];
if (creg_max[k] <= since) continue;
uint64_t a = k << SNAP_SHIFT, b = a + (1ull << SNAP_SHIFT);
if (a < lo) a = lo;
if (b > hi) b = hi;
const uint32_t *ch = cgen_chunk[a >> GEN_CHUNK_SHIFT]; // a 2 MiB region lies in one chunk
if (!ch) continue;
for (uint64_t x = a; x < b; x += 1ull << GEN_PAGE_SHIFT)
if (ch[(x >> GEN_PAGE_SHIFT) & (GEN_PER_CHUNK - 1)] > since) {
range_add(&pend, &npend, &cpend, r, x - lo, 1ull << GEN_PAGE_SHIFT), mrep.page_bytes += 1ull << GEN_PAGE_SHIFT;
r->st_pages += 1ull << GEN_PAGE_SHIFT;
}
}
r->mgen = gen_now;
if (excl) plat_mutex_unlock(&gen_lock);
else plat_mutex_unlock_shared(&gen_lock);
if (__atomic_load_n(&r->fl_n, __ATOMIC_ACQUIRE)) {
FlushRange fl[FLUSH_RANGES + 1];
plat_mutex_lock(&flush_lock);
int n = r->fl_n;
memcpy(fl, r->fl, (size_t)n * sizeof *fl);
uint64_t calls = r->fl_calls, sum = r->fl_sum;
r->fl_n = 0, r->fl_calls = r->fl_sum = 0;
plat_mutex_unlock(&flush_lock);
mrep.flush_calls += calls, mrep.flush_sum += sum;
for (int i = 0; i < n; i++) {
uint64_t a = fl[i].lo > lo ? fl[i].lo : lo, b = fl[i].hi < hi ? fl[i].hi : hi;
if (a >= b) continue;
if (!untracked && !flush_upload) { // the write watch saw every write in them: the pages above
mrep.flush_skipped += b - a;
continue;
}
range_add(&pend, &npend, &cpend, r, a - lo, b - a);
mrep.flush_n++, mrep.flush_bytes += b - a, r->st_flush += b - a;
if ((rt_trace & RT_TRACE_PERF) && b - a >= 16u << 20)
rt_log("mem: mirror frame %llu: flush range %.1f MiB at +%.1f MiB of a %.1f MiB pool (%d ranges from %llu "
"flushes of %.3f MiB)", (unsigned long long)mirror_frame, (b - a) / 1048576.0, (a - lo) / 1048576.0,
(hi - lo) / 1048576.0, n, (unsigned long long)calls, sum / 1048576.0);
}
}
mrep.syncs++, r->st_syncs++;
mrep.sync_t += qpc() - t0;
}
void vk_mem_use(VkMemSpan *sp) {
PoolRec *r = sp->pool;
if (r->mstate == MS_REFUSED && r->retry_sub && (int32_t)(msub - r->retry_sub) >= 0 && evict_enabled()) {
plat_dec64(&m_refused_n[MR_BUDGET]); // counted again if refused again
plat_add64(&m_refused_bytes[MR_BUDGET], -(int64_t)(r->import_hi - r->import_lo));
plat_inc64(&m_retries);
r->mstate = MS_NONE, r->retry_sub = 0; // refuse sets it again
}
if (r->mstate == MS_NONE) mirror_create(r);
if (r->mstate != MS_ON) return;
if (r->acc_sub != msub) mirror_sync(r);
r->vd_use_sub = msub;
sp->buffer = r->mbuf;
}
void vk_mem_pick(VkMemSpan *sp) {
PoolRec *r = sp->pool;
if (!r || r->mstate != MS_ON) return;
if (r->acc_sub != msub) mirror_sync(r);
sp->buffer = r->mbuf;
}
VkBuffer vk_mem_guest_twin(const VkMemSpan *sp) {
const PoolRec *r = sp->pool;
return r && r->mstate == MS_ON && r->twin && sp->buffer == r->mbuf ? r->buf : VK_NULL_HANDLE;
}
int vk_mem_mirror_probe(uint64_t addr) {
VkMemSpan sp;
if (!vk_mirror_on || !vk_mem_resolve(addr, 1, &sp) || !sp.pool) return 0;
mepoch = vk_texture_epoch();
vk_mem_use(&sp);
return ((PoolRec *)sp.pool)->mstate == MS_ON;
}
// ---- write-back merge
// A storage buffer a shader writes whole records of does not
// go back to guest memory as a plain copy. The CPU may write fields of the same records while the GPU runs: NX2navi's
// particle emitters write a new particle's whole record into a free slot of a buffer that the simulation shaders update
// (the waterfall splashes and mist, 2026-11-02), and a plain copy puts the mirror's stale copy of those words back over
// them. Instead:
// - record time (gpu_wrote from vk_mem_gpu_wrote_shader): each distinct range of the submission gets room in a snapshot
// chunk (snap_add);
// - each part's prologue, after the CPU -> mirror copies: the mirror's words of the ranges the part added are copied
// into their snapshots (snap_copies), the state before the part's GPU work. Words of an earlier part's writes that a
// later range overlaps are covered by the earlier range's own snapshot;
// - vk_mem_mirror_end: shaders/wb_merge.comp writes each 32-bit word that differs between the mirror and the snapshot
// into the pool's imported buffer (wb_merge_record), one dispatch per pool and group of ranges.
// Words the GPU left unchanged keep the CPU's value; a word both changed takes the GPU's, as a plain copy would.
// NX2NAVI_WB_MERGE=0 restores the plain copies; so does a pipeline or allocation failure.
#define SNAP_CHUNK (1u << 20)
#define MERGE_SPAN (64ull << 20) // the bytes of one pool a merge dispatch binds at most
#define MERGE_WORDS (65535u * 64u) // the words one merge dispatch covers at most
// Larger written ranges stay plain copies (NX2NAVI_WB_MERGE_MAX, MiB, default 32). A plain copy writes the whole range
// over PCIe at the end of each submission even when the shader changed a few words; a merge copies it within video
// memory and writes only the words that changed. With a 4 MiB limit the field ran about 46 MiB of plain write-back a
// frame in ranges of about 5.5 MiB, just above it; at 16 the same outdoor route went from a median 47.5 to 67.3 fps and
// from 11.3 to 7.2 ms of GPU time a frame (RTX 4080 SUPER, 120 fps limit, 2026-11-08). Kokiri Forest writes seven
// storage buffer ranges of 4-15 MiB a frame; at 32, merged in MERGE_PIECE parts so that no dispatch covers more than
// MERGE_WORDS, the plain copies of about 4 ms of GPU time a frame on PCIe 3.0 went (RTX 4090, 2026-11-09).
#define MERGE_PIECE (4ull << 20)
#define MERGE_MAX_LIMIT (256ull << 20) // the snapshot pool's largest size (snap_pool_cap)
static uint64_t merge_max(void) {
static uint64_t m;
if (!m) {
const char *e = cfg_get("NX2NAVI_WB_MERGE_MAX");
uint64_t mib = e && *e ? (uint64_t)strtoull(e, NULL, 10) : 32;
m = mib << 20;
if (!m) m = 1;
if (m > MERGE_MAX_LIMIT) m = MERGE_MAX_LIMIT;
rt_log("mem: write-back merge covers ranges up to %.1f MiB (NX2NAVI_WB_MERGE_MAX)", m / 1048576.0);
}
return m;
}
typedef struct { PoolRec *r; uint64_t off, size, soff; int chunk; } SnapRange;
typedef struct { VkBuffer buf; VmaAllocation alloc; uint64_t size, used; int host; } SnapChunk;
typedef struct { int n; SnapChunk c[]; } SnapFree; // a submission's chunks and range lists, pooled at its fence
static SnapRange *snl; // replay: this submission's merged ranges; [0, snl_copied) have their snapshot copy
static int nsnl, csnl, snl_copied;
static SnapChunk *chunks; // replay: this submission's snapshot chunks and merge range lists
static int nchunks, cchunks;
// Chunks a finished submission used (snap_free) wait here for chunk_new, so a frame with merges creates no buffer once
// the pool holds the frames in flight: a VMA allocation that needs a new memory block costs 8-11 ms on the replay
// thread. Chunk sizes are powers of two so that one fits the next request; what the pool cannot hold is destroyed. The
// pool holds about three frames of Kokiri Forest's snapshots (65 MiB a frame); at 16 MiB it held less than one frame's
// and two buffers a frame were created and destroyed again (2026-11-09).
#define SNAP_LIST_MIN (64u << 10)
#define SNAP_POOL_N 128
// The pool holds up to 1/64 of the VRAM budget, between 64 and 256 MiB: room for merges of multi-MiB ranges across the
// submissions in flight. Its chunks prefer device memory and are not evictable, so a small card keeps less of them.
static uint64_t snap_pool_cap(void) {
static uint64_t cap;
if (!cap) {
uint64_t used = 0, budget = 0;
int exact = 0;
uint64_t c = vk_mem_vram(&used, &budget, &exact) && budget ? budget / 64 : (64ull << 20);
if (c < (64ull << 20)) c = 64ull << 20;
if (c > (256ull << 20)) c = 256ull << 20;
cap = c;
rt_log("mem: write-back snapshot pool keeps up to %llu MiB", (unsigned long long)(cap >> 20));
}
return cap;
}
static PlatMutex snap_pool_lock = PLAT_MUTEX_INIT;
static SnapChunk snap_pool[SNAP_POOL_N];
static int nsnap_pool;
static uint64_t snap_pool_bytes;
static int merge_on = -1;
static struct {
int tried, ok;
PFN_vkCmdPushDescriptorSetKHR push;
VkDescriptorSetLayout dsl;
VkPipelineLayout layout;
VkPipeline pipe;
} WM;
static int merge_init(VkCtx *c) {
WM.tried = 1;
WM.push = (PFN_vkCmdPushDescriptorSetKHR)vk_dpa(c->device, "vkCmdPushDescriptorSetKHR");
VkDescriptorSetLayoutBinding b[4];
for (int i = 0; i < 4; i++)
b[i] = (VkDescriptorSetLayoutBinding){.binding = (uint32_t)i, .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
.descriptorCount = 1, .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT};
VkDescriptorSetLayoutCreateInfo dci = {.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
.flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
.bindingCount = 4, .pBindings = b};
VkPushConstantRange pr = {VK_SHADER_STAGE_COMPUTE_BIT, 0, 8};
VkShaderModuleCreateInfo mi = {.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO, .codeSize = sizeof wb_merge_spv,
.pCode = wb_merge_spv};
VkShaderModule m = VK_NULL_HANDLE;
if (!WM.push || vkCreateDescriptorSetLayout(c->device, &dci, NULL, &WM.dsl) != VK_SUCCESS) goto fail;
VkPipelineLayoutCreateInfo li = {.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, .setLayoutCount = 1,
.pSetLayouts = &WM.dsl, .pushConstantRangeCount = 1, .pPushConstantRanges = &pr};
if (vkCreatePipelineLayout(c->device, &li, NULL, &WM.layout) != VK_SUCCESS ||
vkCreateShaderModule(c->device, &mi, NULL, &m) != VK_SUCCESS)
goto fail;
VkComputePipelineCreateInfo ci = {.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
.stage = {.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
.stage = VK_SHADER_STAGE_COMPUTE_BIT, .module = m, .pName = "main"},
.layout = WM.layout};
VkResult res = vkCreateComputePipelines(c->device, VK_NULL_HANDLE, 1, &ci, NULL, &WM.pipe);
vkDestroyShaderModule(c->device, m, NULL);
if (res != VK_SUCCESS) goto fail;
return WM.ok = 1;
fail:
rt_log("mem: write-back merge off: its pipeline could not be created");
return 0;
}
static int merging(void) {
if (merge_on < 0) {
merge_on = cfg_bool("NX2NAVI_WB_MERGE");
rt_log("mem: write-back merge %s (NX2NAVI_WB_MERGE)", merge_on ? "on" : "off");
}
if (!merge_on) return 0;
if (!WM.tried) merge_init(vk_ctx());
return WM.ok;
}
// A chunk of at least size bytes in chunks[]: the smallest fitting one of the pool, else a new buffer; -1 when the
// allocation fails.
static int chunk_new(uint64_t size, int host) {
uint64_t want = host ? SNAP_LIST_MIN : SNAP_CHUNK;
while (want < size) want <<= 1;
if (nchunks == cchunks) grow((void **)&chunks, &cchunks, sizeof *chunks);
SnapChunk *ch = &chunks[nchunks];
int best = -1;
plat_mutex_lock(&snap_pool_lock);
for (int i = 0; i < nsnap_pool; i++)
if (snap_pool[i].host == host && snap_pool[i].size >= want && (best < 0 || snap_pool[i].size < snap_pool[best].size))
best = i;
if (best >= 0) {
*ch = snap_pool[best];
snap_pool[best] = snap_pool[--nsnap_pool];
snap_pool_bytes -= ch->size;
}
plat_mutex_unlock(&snap_pool_lock);
if (best < 0) {
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .size = want,
.usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT,
.sharingMode = VK_SHARING_MODE_EXCLUSIVE};
VmaAllocationCreateInfo ai = {.usage = host ? VMA_MEMORY_USAGE_AUTO_PREFER_HOST : VMA_MEMORY_USAGE_AUTO_PREFER_DEVICE,
.flags = host ? VMA_ALLOCATION_CREATE_HOST_ACCESS_SEQUENTIAL_WRITE_BIT |
VMA_ALLOCATION_CREATE_MAPPED_BIT : 0};
if (vmaCreateBuffer(vk_ctx()->vma, &bi, &ai, &ch->buf, &ch->alloc, NULL) != VK_SUCCESS) return -1;
ch->size = want, ch->host = host;
mrep.snap_allocs++;
}
ch->used = 0;
return nchunks++;
}
// Record time: room in a snapshot chunk for [off, off+size) of r's mirror, once per submission. 0: plain copy instead.
static int snap_add(PoolRec *r, uint64_t off, uint64_t size) {
for (int i = nsnl; i-- > 0;)
if (snl[i].r == r && snl[i].off == off && snl[i].size == size) return 1;
int k = -1;
for (int i = nchunks; i-- > 0;) // the newest snapshot chunk, if the range fits
if (chunks[i].used != UINT64_MAX && chunks[i].size - chunks[i].used >= size) { k = i; break; }
if (k < 0 && (k = chunk_new(size > SNAP_CHUNK ? size : SNAP_CHUNK, 0)) < 0) return 0;
if (nsnl == csnl) grow((void **)&snl, &csnl, sizeof *snl);
snl[nsnl++] = (SnapRange){r, off, size, chunks[k].used, k};
chunks[k].used += (size + 15) & ~15ull;
return 1;
}
// Prologue: the snapshot copies of the ranges added since the last prologue.
static void snap_copies(VkCommandBuffer cb) {
for (int i = snl_copied; i < nsnl; i++) {
VkBufferCopy cp = {snl[i].off, snl[i].soff, snl[i].size};
vkCmdCopyBuffer(cb, snl[i].r->mbuf, chunks[snl[i].chunk].buf, 1, &cp);
mrep.snap_bytes += snl[i].size;
}
snl_copied = nsnl;
}
static int snap_cmp(const void *x, const void *y) {
const SnapRange *a = x, *b = y;
if (a->r != b->r) return (uintptr_t)a->r < (uintptr_t)b->r ? -1 : 1;
if (a->chunk != b->chunk) return a->chunk < b->chunk ? -1 : 1;
return a->off < b->off ? -1 : a->off > b->off;
}
// One dispatch: ranges s[0..n) share a pool and a chunk and lie in [base, base + span) of the pool.
static void merge_dispatch(VkCommandBuffer cb, const SnapRange *s, int n, uint64_t base, uint64_t span, int list,
uint64_t list_off, uint32_t total) {
VkDescriptorBufferInfo bi[4] = {{s->r->mbuf, base, span}, {chunks[s->chunk].buf, 0, chunks[s->chunk].size},
{s->r->buf, base, span}, {chunks[list].buf, list_off, (uint64_t)n * 16}};
VkWriteDescriptorSet w[4];
for (int i = 0; i < 4; i++)
w[i] = (VkWriteDescriptorSet){.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, .dstBinding = (uint32_t)i,
.descriptorCount = 1, .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
.pBufferInfo = &bi[i]};
rec_push(WM.push, cb, VK_PIPELINE_BIND_POINT_COMPUTE, WM.layout, 0, 4, w); // in call order (vk_rec.h)
uint32_t pc[2] = {(uint32_t)n, total};
vkCmdPushConstants(cb, WM.layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof pc, pc);
vkCmdDispatch(cb, (total + 63) / 64, 1, 1);
}
// vk_mem_mirror_end: the merge dispatches of this submission's ranges. Returns the commands recorded; 0 with nsnl
// left as it was when the range list cannot be allocated (the caller copies them plainly).
static int wb_merge_record(VkCommandBuffer cb) {
VkCtx *c = vk_ctx();
uint64_t align = c->props.limits.minStorageBufferOffsetAlignment;
if (align < 16) align = 16;
SnapRange *s = malloc(sizeof *s * (size_t)nsnl);
if (!s) rt_fatal("vk: out of memory");
memcpy(s, snl, sizeof *s * (size_t)nsnl);
qsort(s, (size_t)nsnl, sizeof *s, snap_cmp);
// the range lists: at most one 16-byte entry per range, each dispatch's list at an aligned offset
int list = chunk_new((uint64_t)nsnl * 16 + (uint64_t)nsnl * align, 1);
if (list < 0) { free(s); return 0; }
chunks[list].used = UINT64_MAX; // not a snapshot chunk
VmaAllocationInfo li;
vmaGetAllocationInfo(c->vma, chunks[list].alloc, &li);
uint32_t *lm = li.pMappedData;
uint64_t lo = 0;
int rec = 0;
vkCmdBindPipeline(cb, VK_PIPELINE_BIND_POINT_COMPUTE, WM.pipe);
for (int i = 0; i < nsnl;) {
uint64_t base = s[i].off & ~(align - 1), end = s[i].off + s[i].size;
uint32_t total = 0;
int j = i;
lo = (lo + align - 1) & ~(align - 1);
uint64_t l0 = lo;
for (; j < nsnl && s[j].r == s[i].r && s[j].chunk == s[i].chunk; j++) {
uint64_t e = s[j].off + s[j].size > end ? s[j].off + s[j].size : end;
if (j > i && (e - base > MERGE_SPAN || total + s[j].size / 4 > MERGE_WORDS)) break;
end = e;
uint32_t *q = lm + lo / 4;
q[0] = (uint32_t)((s[j].off - base) / 4), q[1] = (uint32_t)(s[j].soff / 4), q[2] = (uint32_t)(s[j].size / 4);
q[3] = total;
total += (uint32_t)(s[j].size / 4);
lo += 16;
mrep.merge_bytes += s[j].size;
}
merge_dispatch(cb, s + i, j - i, base, end - base, list, l0, total);
mrep.merge_dispatches++;
rec += 3;
i = j;
}
vmaFlushAllocation(c->vma, chunks[list].alloc, 0, VK_WHOLE_SIZE);
for (int i = 0; i < nsnl; i++) snl[i].r->wb_sub = msub;
free(s);
return rec;
}
static void snap_free(void *arg) { // completion thread: the submission is done with its chunks
SnapFree *f = arg;
plat_mutex_lock(&snap_pool_lock);
for (int i = 0; i < f->n; i++) {
SnapChunk *ch = &f->c[i];
if (nsnap_pool == SNAP_POOL_N || snap_pool_bytes + ch->size > snap_pool_cap()) continue;
snap_pool[nsnap_pool++] = *ch;
snap_pool_bytes += ch->size;
ch->buf = VK_NULL_HANDLE;
}
plat_mutex_unlock(&snap_pool_lock);
for (int i = 0; i < f->n; i++)
if (f->c[i].buf) vmaDestroyBuffer(vk_ctx()->vma, f->c[i].buf, f->c[i].alloc);
free(f);
}
static void gpu_wrote(const VkMemSpan *sp, uint64_t size, int shader) {
PoolRec *r = sp->pool;
if (!r || r->mstate != MS_ON || sp->buffer != r->mbuf || !size) return;
r->vd_wrote_sub = msub, r->vd_wrote_n++;
// A binding to the end of its pool (no size) can span hundreds of MiB: those stay copies.
uint64_t o = 0;
if (shader && !(sp->offset & 3) && !(size & 3) && size <= merge_max() && merging())
for (uint64_t n; o < size; o += n)
if (!snap_add(r, sp->offset + o, n = size - o < MERGE_PIECE ? size - o : MERGE_PIECE)) break;
r->st_merge += o;
if (o == size) return;
r->st_wb += size - o;
range_add(&wbl, &nwbl, &cwbl, r, sp->offset + o, size - o); // what found no snapshot room
}
void vk_mem_gpu_wrote(const VkMemSpan *sp, uint64_t size) { gpu_wrote(sp, size, 0); }
void vk_mem_gpu_wrote_shader(const VkMemSpan *sp, uint64_t size) { gpu_wrote(sp, size, 1); }
// Adds [a, b) to l, n sorted ranges that do not touch; returns the new count. Above cap ranges (l holds cap + 1) the two
// nearest merge, so a full list copies the fewest bytes that were not flushed; cap 1 keeps one span over all flushes.
static int flush_add(FlushRange *l, int n, uint64_t a, uint64_t b, int cap) {
int i = 0;
while (i < n && l[i].hi < a) i++;
int j = i;
for (; j < n && l[j].lo <= b; j++) {
if (l[j].lo < a) a = l[j].lo;
if (l[j].hi > b) b = l[j].hi;
}
memmove(l + i + 1, l + j, (size_t)(n - j) * sizeof *l);
l[i] = (FlushRange){a, b};
n += 1 - (j - i);
if (n > cap) {
int k = 0;
for (int x = 1; x + 1 < n; x++)
if (l[x + 1].lo - l[x].hi < l[k + 1].lo - l[k].hi) k = x;
l[k].hi = l[k + 1].hi;
memmove(l + k + 1, l + k + 2, (size_t)(n - k - 2) * sizeof *l);
n--;
}
return n;
}
int vk_mem_flush_add_probe(void) {
FlushRange l[5];
int n = 0;
n = flush_add(l, n, 0x8000, 0x9000, 4);
n = flush_add(l, n, 0x1000, 0x2000, 4);
n = flush_add(l, n, 0x4000, 0x5000, 4);
if (n != 3 || l[0].lo != 0x1000 || l[1].lo != 0x4000 || l[2].hi != 0x9000) return 1; // kept apart, sorted
n = flush_add(l, n, 0x2000, 0x3000, 4);
if (n != 3 || l[0].hi != 0x3000) return 2; // touching ranges join
n = flush_add(l, n, 0x800, 0x8800, 4);
if (n != 1 || l[0].lo != 0x800 || l[0].hi != 0x9000) return 3; // one covers several
n = 0;
n = flush_add(l, n, 0x10000, 0x11000, 2);
n = flush_add(l, n, 0x40000, 0x41000, 2);
n = flush_add(l, n, 0x12000, 0x13000, 2); // over the cap: nearest merge
if (n != 2 || l[0].lo != 0x10000 || l[0].hi != 0x13000 || l[1].lo != 0x40000) return 4;
n = 0;
n = flush_add(l, n, 0x10000, 0x11000, 1);
n = flush_add(l, n, 0x40000, 0x41000, 1);
if (n != 1 || l[0].lo != 0x10000 || l[0].hi != 0x41000) return 5; // NX2NAVI_FLUSH_SPAN=1
return 0;
}
uint64_t vk_mem_mirror_uploaded(void) {
plat_mutex_lock(&mirror_stat_lock);
uint64_t n = macc.upload_bytes;
plat_mutex_unlock(&mirror_stat_lock);
return n;
}
void nvn_pool_flushed(const pool_t *p, uint64_t off, uint64_t size) {
if (!vk_mirror_on || !size || !p->mem) return;
plat_inc64(&m_flush_calls);
plat_add64(&m_flush_bytes, (int64_t)size);
plat_mutex_lock_shared(&pool_lock);
const pidx_ent *e = pidx_find(&pool_index, (uintptr_t)p->mem, 0);
PoolRec *r = e && e->key == p ? e->buf : NULL;
if (r && r->mstate == MS_ON) {
uint64_t a = ((uintptr_t)p->mem + off) & ~0xfffull, b = ((uintptr_t)p->mem + off + size + 0xfff) & ~0xfffull;
plat_mutex_lock(&flush_lock);
__atomic_store_n(&r->fl_n, flush_add(r->fl, r->fl_n, a, b, flush_cap), __ATOMIC_RELEASE);
r->fl_calls++, r->fl_sum += size;
plat_mutex_unlock(&flush_lock);
}
plat_mutex_unlock_shared(&pool_lock);
}
// gen_lock exclusive: freezes up to FREEZE_PER_SUB candidate regions still idle. The query after the
// protection collects the writes made before it took effect.
static void freeze_pass(void) {
for (int done = 0; nfcand && done < FREEZE_PER_SUB;) {
uint64_t j = fcand[--nfcand], a = j << REG_SHIFT;
if (ww_ep - reg_last[j] < (uint32_t)freeze_idle || mem_frozen(a)) continue;
done++;
if (!mem_freeze(a)) { n_frz_refused++; continue; }
query_stamp(a, a + (1ull << REG_SHIFT), ++gen_now);
}
}
void vk_mem_mirror_begin(void) {
if (!vk_mirror_on) { // import mode: the texture checks' snapshots still skip frozen regions
if (!freeze_idle) return;
plat_mutex_lock(&gen_lock);
ww_ep++;
freeze_pass();
plat_mutex_unlock(&gen_lock);
return;
}
uint64_t t0 = qpc();
nsnl = snl_copied = nchunks = 0; // the previous submission handed its chunks to snap_free
uint32_t prev = msub;
msub = msub + 1 ? msub + 1 : 1;
__atomic_add_fetch(&mirror_gen, 1, __ATOMIC_RELEASE);
mepoch = vk_texture_epoch();
uint64_t q = 0, run0 = 0, run1 = 0;
int have = 0;
plat_mutex_lock_shared(&pool_lock);
plat_mutex_lock(&gen_lock);
ww_ep++;
uint64_t b0 = ww_bytes;
for (int i = 0; i < nm; i++) {
const PoolRec *r = mlist[i];
if (r->acc_sub != prev) continue;
uint64_t a = r->import_lo >> SNAP_SHIFT, b = (r->import_hi - 1) >> SNAP_SHIFT;
if (have && a <= run1 + 1 + (uint64_t)mirror_gap) {
if (b > run1) run1 = b;
continue;
}
if (have) q += snap_regions(run0, run1, mepoch);
run0 = a, run1 = b, have = 1;
}
if (have) q += snap_regions(run0, run1, mepoch);
mrep.presnap_bytes += ww_bytes - b0;
if (freeze_idle) freeze_pass();
plat_mutex_unlock(&gen_lock);
plat_mutex_unlock_shared(&pool_lock);
mrep.subs++;
mrep.presnap_queries += q;
mrep.presnap_t += qpc() - t0;
}
int vk_mem_mirror_pending(void) { return npend || snl_copied < nsnl; }
void vk_mem_mirror_prologue(VkCommandBuffer cb) {
if (!npend && snl_copied == nsnl) return;
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);
copy_ranges(cb, pend, npend, 1, &mrep.upload_bytes, &mrep.upload_copies);
if (snl_copied < nsnl) { // the snapshots read what the uploads above wrote
if (npend) mbarrier(cb, VK_PIPELINE_STAGE_2_COPY_BIT, VK_ACCESS_2_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_2_COPY_BIT,
VK_ACCESS_2_TRANSFER_READ_BIT);
snap_copies(cb);
}
mbarrier(cb, VK_PIPELINE_STAGE_2_COPY_BIT, VK_ACCESS_2_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
VK_ACCESS_2_MEMORY_READ_BIT | VK_ACCESS_2_MEMORY_WRITE_BIT);
npend = 0;
mrep.prologues++;
}
// ---- verify mode
typedef struct { PoolRec *r; uint64_t off, len, dst; uint32_t mgen; } VSample;
typedef struct { uint32_t sub; int n; VkBuffer buf; VmaAllocation alloc; const uint8_t *map; VSample s[VERIFY_SAMPLES]; } VJob;
static int ver_i;
static uint64_t ver_off;
// The CPU wrote the page since the mirror's sync that mgen records: stamped later, or its write-watch bit is set.
static int page_dirty(uint64_t a, uint32_t mgen) {
uint32_t *s = gen_slot(a, 0);
if (s && __atomic_load_n(s, __ATOMIC_RELAXED) > mgen) return 1;
uint64_t p[1];
return mem_written_peek(a, 1ull << GEN_PAGE_SHIFT, p, 1) != 0; // UINT64_MAX (not watched) counts as written
}
static void verify_cb(void *arg) { // completion thread
VJob *j = arg;
static volatile int32_t logged;
uint64_t pages = 0, skipped = 0, bad = 0;
for (int i = 0; i < j->n; i++) {
const VSample *s = &j->s[i];
PoolRec *r = s->r;
uint32_t wb = __atomic_load_n(&r->wb_sub, __ATOMIC_RELAXED);
if (wb != j->sub && wb - j->sub < 0x80000000u) { // a later submission wrote back into it
skipped += s->len >> GEN_PAGE_SHIFT;
continue;
}
for (uint64_t o = 0; o < s->len; o += 1ull << GEN_PAGE_SHIFT) {
uint64_t a = r->import_lo + s->off + o;
const uint8_t *m = j->map + s->dst + o, *g = (const uint8_t *)(uintptr_t)a;
pages++;
if (!memcmp(m, g, 1ull << GEN_PAGE_SHIFT)) continue;
// gen_lock: a write-watch query clears the page's bit and stamps it in one hold (vk_mem_changed, mirror_sync),
// so the check sees one or the other; without it a page between the two counted as a mismatch.
plat_mutex_lock_shared(&gen_lock);
int dirty = page_dirty(a, s->mgen);
plat_mutex_unlock_shared(&gen_lock);
if (dirty) { skipped++; continue; }
bad++;
if (plat_inc32(&logged) <= VERIFY_LOG) {
int k = 0, e = (1 << GEN_PAGE_SHIFT) - 1;
while (k < e && m[k] == g[k]) k++;
while (e > k && m[e] == g[e]) e--;
char f[160];
flag_names(r->flags, f, sizeof f);
rt_log("mem: verify mismatch: pool %#llx+%#llx (%s) offset %#llx: bytes %#x..%#x of the page differ, "
"mirror %02x guest %02x at the first (submission %u)", (unsigned long long)r->lo,
(unsigned long long)r->size, f, (unsigned long long)(a - r->lo), k, e, m[k], g[k], j->sub);
const uint32_t *gs = gen_slot(a & ~((1ull << GEN_PAGE_SHIFT) - 1), 0);
char mh[100], gh[100];
for (int b = 0; b < 16; b++) snprintf(mh + 3 * b, 4, "%02x ", m[b]), snprintf(gh + 3 * b, 4, "%02x ", g[b]);
rt_log("mem: verify mismatch detail: written binds %llu, last %u; used %u; written back %u; gpa record %u; "
"page cpu stamp %u (synced to %u, now %u); mirror %s| guest %s",
(unsigned long long)r->vd_wrote_n, r->vd_wrote_sub, r->vd_use_sub, wb, r->gpa_rec, gs ? *gs : 0,
s->mgen, gen_now, mh, gh);
}
}
}
plat_add64(&v_pages, (int64_t)pages);
plat_add64(&v_skipped, (int64_t)skipped);
plat_add64(&v_mismatch, (int64_t)bad);
vmaDestroyBuffer(vk_ctx()->vma, j->buf, j->alloc);
free(j);
}
static int verify_record(VkCommandBuffer cb, uint64_t fence) {
VJob *j = calloc(1, sizeof *j);
if (!j) rt_fatal("vk: out of memory");
uint64_t left = VERIFY_BYTES, dst = 0;
plat_mutex_lock_shared(&pool_lock);
for (int tries = 0; nm && left && j->n < VERIFY_SAMPLES && tries <= nm; tries++) {
if (ver_i >= nm) ver_i = 0, ver_off = 0;
PoolRec *r = mlist[ver_i];
uint64_t size = r->import_hi - r->import_lo;
if (ver_off >= size) { ver_i++, ver_off = 0; continue; }
uint64_t len = size - ver_off < left ? size - ver_off : left;
j->s[j->n++] = (VSample){r, ver_off, len, dst, r->mgen};
dst += len, left -= len, ver_off += len;
if (ver_off >= size) ver_i++, ver_off = 0;
}
plat_mutex_unlock_shared(&pool_lock);
if (!j->n) { free(j); return 0; }
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .size = dst,
.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT};
VmaAllocationCreateInfo ai = {.usage = VMA_MEMORY_USAGE_AUTO_PREFER_HOST,
.flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_RANDOM_BIT | VMA_ALLOCATION_CREATE_MAPPED_BIT};
VmaAllocationInfo info;
if (vmaCreateBuffer(vk_ctx()->vma, &bi, &ai, &j->buf, &j->alloc, &info) != VK_SUCCESS) { free(j); return 0; }
j->map = info.pMappedData;
j->sub = msub;
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);
for (int i = 0; i < j->n; i++) {
VkBufferCopy c = {j->s[i].off, j->s[i].dst, j->s[i].len};
vkCmdCopyBuffer(cb, j->s[i].r->mbuf, j->buf, 1, &c);
}
mbarrier(cb, VK_PIPELINE_STAGE_2_COPY_BIT, VK_ACCESS_2_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_2_HOST_BIT,
VK_ACCESS_2_HOST_READ_BIT);
vk_completion_add(fence, verify_cb, j);
return j->n + 2;
}
int vk_mem_mirror_end(VkCommandBuffer cb, uint64_t fence) {
if (!vk_mirror_on) return 0;
int rec = 0;
if (nsnl) {
// the merge reads the mirror the submission's shaders wrote and writes the imported buffers
mbarrier(cb, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT, VK_ACCESS_2_MEMORY_WRITE_BIT,
VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, VK_ACCESS_2_SHADER_STORAGE_READ_BIT);
int m = wb_merge_record(cb);
if (!m) { // no range list: copy them plainly
for (int i = 0; i < nsnl; i++) range_add(&wbl, &nwbl, &cwbl, snl[i].r, snl[i].off, snl[i].size);
}
rec += m + 1;
mbarrier(cb, VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, VK_ACCESS_2_SHADER_STORAGE_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);
}
if (nchunks) { // the last part's prologue still records snapshot copies into them; back to the pool at the fence
SnapFree *f = malloc(sizeof *f + sizeof(SnapChunk) * (size_t)nchunks);
if (!f) rt_fatal("vk: out of memory");
f->n = nchunks;
memcpy(f->c, chunks, sizeof(SnapChunk) * (size_t)nchunks);
vk_completion_add(fence, snap_free, f);
}
if (nwbl) {
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);
for (int i = 0; i < nwbl; i++) wbl[i].r->wb_sub = msub;
mrep.wb_ranges += (uint64_t)nwbl;
rec += copy_ranges(cb, wbl, nwbl, 0, &mrep.wb_bytes, &mrep.wb_copies) + 2;
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);
nwbl = 0;
}
if (mirror_verify) rec += verify_record(cb, fence);
plat_mutex_lock(&mirror_stat_lock);
uint64_t *d = (uint64_t *)&macc, *s = (uint64_t *)&mrep;
for (size_t i = 0; i < sizeof macc / sizeof(uint64_t); i++) d[i] += s[i];
plat_mutex_unlock(&mirror_stat_lock);
memset(&mrep, 0, sizeof mrep);
return rec;
}
// NX2NAVI_TRACE=perf, every 600 frames: the mirrored pools that moved the most bytes since the last list, per frame: what
// their CPU writes and flushes copied up, how often a submission synced them, and what the GPU's writes sent back
// (merged or copied). A pool whose upload is close to its size every frame is rewritten whole each frame.
#define POOL_LOG_TOP 12
static uint64_t pool_score(const PoolRec *r) { return r->st_pages + r->st_flush + r->st_wb + r->st_merge; }
static void mirror_pool_log(uint64_t frame) {
static uint64_t last;
double fr = frame > last ? (double)(frame - last) : 1.0;
last = frame;
PoolRec *top[POOL_LOG_TOP];
int nt = 0;
uint64_t all_up = 0;
plat_mutex_lock_shared(&pool_lock);
for (int i = 0; i < nm; i++) {
PoolRec *r = mlist[i];
all_up += r->st_pages + r->st_flush;
uint64_t sc = pool_score(r);
if (!sc || (nt == POOL_LOG_TOP && sc <= pool_score(top[nt - 1]))) continue;
int k = nt < POOL_LOG_TOP ? nt++ : nt - 1; // insertion into top[], largest first
for (; k > 0 && pool_score(top[k - 1]) < sc; k--) top[k] = top[k - 1];
top[k] = r;
}
rt_log("mem: mirror pools frame %llu: %d mirrored, %.1f MiB a frame copied up in all; the top %d (MiB a frame: written pages, "
"flush ranges (overlap with the pages counted twice), syncs a frame, GPU writes back merged / copied):",
(unsigned long long)frame, nm, all_up / fr / 1048576.0, nt);
for (int i = 0; i < nt; i++) {
const PoolRec *r = top[i];
char f[160];
flag_names(r->flags, f, sizeof f);
rt_log("mem: pool %#llx %.1f MiB %s: up %.2f + %.2f, %.1f syncs, back %.2f / %.2f", (unsigned long long)r->lo,
r->size / 1048576.0, f, r->st_pages / fr / 1048576.0, r->st_flush / fr / 1048576.0, r->st_syncs / fr,
r->st_merge / fr / 1048576.0, r->st_wb / fr / 1048576.0);
}
for (int i = 0; i < nm; i++) {
PoolRec *r = mlist[i];
r->st_pages = r->st_flush = r->st_wb = r->st_merge = 0, r->st_syncs = 0;
}
plat_mutex_unlock_shared(&pool_lock);
}
// One summary line: with NX2NAVI_TRACE=perf every 60 frames, else at frames 600 and 1800.
static void mirror_log(uint64_t frame) {
static uint64_t last, ww_last;
plat_mutex_lock(&gen_lock);
uint64_t wc = ww_calls, wb = ww_bytes, wt = ww_ticks, dc = ww_dirty_calls, dt = ww_dirty_ticks, wp = ww_pages;
ww_pages = 0;
uint64_t nr = n_frz_refused;
ww_calls = ww_bytes = ww_ticks = ww_dirty_calls = ww_dirty_ticks = n_frz_refused = 0;
plat_mutex_unlock(&gen_lock);
double wf = frame > ww_last ? (double)(frame - ww_last) : 1.0;
ww_last = frame;
uint64_t fz, ff, ft, fo;
mem_freeze_stats(&fz, &ff, &ft, &fo);
if (rt_trace & RT_TRACE_PERF)
rt_log("mem: write-watch frame %llu: per frame, every consumer: %.1f queries over %.1f MiB, %.3f ms (%.1f found "
"%.0f written pages, %.3f ms); frozen %.0f MiB; since the last line: freezes %llu, thaws by a write %llu, by a file read "
"or mapping change %llu, refused %llu",
(unsigned long long)frame, wc / wf, wb / wf / 1048576.0, qpc_ms(wt) / wf, dc / wf, wp / wf, qpc_ms(dt) / wf,
fz * 2.0, (unsigned long long)ff, (unsigned long long)ft, (unsigned long long)fo, (unsigned long long)nr);
if (!vk_mirror_on) return;
plat_mutex_lock(&mirror_stat_lock);
MirrorAcc a = macc;
memset(&macc, 0, sizeof macc);
plat_mutex_unlock(&mirror_stat_lock);
double f = frame > last ? (double)(frame - last) : 1.0, M = 1048576.0;
last = frame;
uint64_t fc = (uint64_t)plat_xchg64(&m_flush_calls, 0), fb = (uint64_t)plat_xchg64(&m_flush_bytes, 0);
char ver[160] = "";
if (mirror_verify)
snprintf(ver, sizeof ver, "; verify: pages compared %lld, skipped %lld (written since), mismatches %lld",
(long long)plat_xchg64(&v_pages, 0), (long long)plat_xchg64(&v_skipped, 0),
(long long)plat_xchg64(&v_mismatch, 0));
vk_gpa_report();
rt_log("mem: mirror frame %llu: pools %lld, %.1f MiB device-local; stay imported: budget %lld (%.1f MiB), untracked "
"%lld, overlapping %lld, failed %lld; per frame: submits %.2f, syncs %.1f, upload %.3f MiB in %.1f ranges, "
"first use %.2f pools %.3f MiB, write-back %.3f MiB in %.1f ranges, merge %.3f MiB in %.1f dispatches "
"(snapshots %.3f MiB, new buffers %.2f), flush calls %.1f (%.3f MiB), queries "
"%.1f pre-snap (%.1f MiB) + %.1f sync, prologues %.2f, tracking cpu %.3f ms (pre-snap %.3f, syncs %.3f); so far "
"evicted %lld (%.1f MiB), budget refusals asked again %lld, made without a twin %lld (%.1f MiB)%s",
(unsigned long long)frame, (long long)m_live_n, m_live_bytes / M, (long long)m_refused_n[MR_BUDGET],
m_refused_bytes[MR_BUDGET] / M, (long long)m_refused_n[MR_UNTRACKED], (long long)m_refused_n[MR_OVERLAP],
(long long)m_refused_n[MR_ALLOC], a.subs / f, a.syncs / f, a.upload_bytes / f / M, a.upload_copies / f,
a.first_n / f, a.first_bytes / f / M, a.wb_bytes / f / M, a.wb_copies / f, a.merge_bytes / f / M,
a.merge_dispatches / f, a.snap_bytes / f / M, a.snap_allocs / f, fc / f, fb / f / M,
a.presnap_queries / f, a.presnap_bytes / f / M, a.sync_queries / f, a.prologues / f, qpc_ms(a.presnap_t + a.sync_t) / f,
qpc_ms(a.presnap_t) / f, qpc_ms(a.sync_t) / f, (long long)m_evicted_n, m_evicted_bytes / M, (long long)m_retries,
(long long)m_untwinned_n, m_untwinned_bytes / M, ver);
rt_log("mem: mirror frame %llu: upload sources per frame, overlaps counted twice: written pages %.3f MiB; flush ranges "
"%.2f, %.3f MiB, from %.1f flushes of %.3f MiB; flushed and not copied (write-watched) %.3f MiB",
(unsigned long long)frame, a.page_bytes / f / M, a.flush_n / f,
a.flush_bytes / f / M, a.flush_calls / f, a.flush_sum / f / M, a.flush_skipped / f / M);
}
// After vram_log: every heap's usage as the driver reports it (VK_EXT_memory_budget) beside VMA's share of it, VMA by
// memory type, and the process's private commit with its write-combined part (host-visible memory the driver maps
// uncached), to tell which memory a long session grows. When the write-combined total (VMA's blocks where that is not
// known) grew by 2 GiB since the last look, VMA's detailed map goes to vma-frame<N>.json beside the log, 4 at most.
// The process's part walks the whole address space (plat_mem_census), 100-240 ms in a long session, so it runs on a
// thread of its own (census_main) and logs a line of its own; a look while the last walk still runs goes without it.
typedef struct { VkCtx *c; uint64_t frame; } CensusJob;
static volatile int32_t census_busy;
static void census_main(void *arg) {
CensusJob j = *(CensusJob *)arg;
free(arg);
const double M = 1048576.0;
uint64_t priv = 0, wc = 0, wcn = 0, now;
int have = plat_mem_census(&priv, &wc, &wcn);
if (have) {
rt_log("mem: process frame %llu: private %.0f MiB, write-combined %.0f MiB in %llu allocations",
(unsigned long long)j.frame, priv / M, wc / M, (unsigned long long)wcn);
now = wc;
} else {
VmaBudget vb[VK_MAX_MEMORY_HEAPS];
vmaGetHeapBudgets(j.c->vma, vb);
now = 0;
for (uint32_t h = 0; h < j.c->mem_props.memoryHeapCount && h < VK_MAX_MEMORY_HEAPS; h++) now += vb[h].statistics.blockBytes;
}
static uint64_t base;
static int dumps = -1; // -1: no baseline yet
if (dumps < 0) base = now, dumps = 0;
else if (dumps < 4 && now >= base + (2ull << 30) && rt_log_path()) {
base = now, dumps++;
const char *lp = rt_log_path(), *sl = strrchr(lp, '/'), *bs = strrchr(lp, '\\');
if (bs && (!sl || bs > sl)) sl = bs;
char path[1024];
snprintf(path, sizeof path, "%.*svma-frame%llu.json", sl ? (int)(sl - lp) + 1 : 0, lp, (unsigned long long)j.frame);
char *json = NULL;
vmaBuildStatsString(j.c->vma, &json, VK_TRUE);
FILE *f = json ? fopen(path, "wb") : NULL;
int ok = f && fputs(json, f) >= 0;
if (f && fclose(f)) ok = 0;
rt_log("mem: vma detailed map %s %s", ok ? "written to" : "not written:", path);
if (json) vmaFreeStatsString(j.c->vma, json);
}
plat_xchg32(&census_busy, 0);
}
static void heap_log(VkCtx *c, uint64_t frame, const VmaBudget *vb) {
const double M = 1048576.0;
VkPhysicalDeviceMemoryBudgetPropertiesEXT bp = {.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_BUDGET_PROPERTIES_EXT};
if (c->has_memory_budget) {
VkPhysicalDeviceMemoryProperties2 p2 = {.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_PROPERTIES_2, .pNext = &bp};
vkGetPhysicalDeviceMemoryProperties2(c->phys, &p2);
}
char line[1536];
int o = snprintf(line, sizeof line, "mem: heaps frame %llu:", (unsigned long long)frame);
for (uint32_t h = 0; h < c->mem_props.memoryHeapCount && h < VK_MAX_MEMORY_HEAPS && o < (int)sizeof line - 160; h++)
o += snprintf(line + o, sizeof line - (size_t)o, " heap %u (%s) used %.0f of %.0f MiB, vma %u blocks %.0f MiB, %u allocations %.0f MiB;",
h, c->mem_props.memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT ? "device" : "host",
(c->has_memory_budget ? bp.heapUsage[h] : vb[h].usage) / M,
(c->has_memory_budget ? bp.heapBudget[h] : vb[h].budget) / M, vb[h].statistics.blockCount,
vb[h].statistics.blockBytes / M, vb[h].statistics.allocationCount, vb[h].statistics.allocationBytes / M);
rt_log("%s", line);
VmaTotalStatistics ts;
vmaCalculateStatistics(c->vma, &ts);
o = snprintf(line, sizeof line, "mem: vma by type frame %llu:", (unsigned long long)frame);
for (uint32_t t = 0; t < c->mem_props.memoryTypeCount && t < VK_MAX_MEMORY_TYPES && o < (int)sizeof line - 120; t++) {
const VmaDetailedStatistics *d = &ts.memoryType[t];
if (!d->statistics.blockCount) continue;
o += snprintf(line + o, sizeof line - (size_t)o,
" type %u: %u blocks %.0f MiB, %u allocations %.0f MiB, largest %.1f MiB, %u free ranges (largest %.1f MiB);",
t, d->statistics.blockCount, d->statistics.blockBytes / M, d->statistics.allocationCount,
d->statistics.allocationBytes / M, d->allocationSizeMax / M, d->unusedRangeCount,
d->unusedRangeCount ? d->unusedRangeSizeMax / M : 0.0);
}
rt_log("%s", line);
// The address-space walk of the census (plat_mem_census, 100-240 ms of VirtualQuery) runs on its own thread, but
// it holds the process's address-space lock, which the game's threads take at every page fault and every
// VirtualAlloc: the present after each walk arrived 65-120 ms late. NX2NAVI_MEM_CENSUS=1 asks for it.
static int census_on = -1;
if (census_on < 0) census_on = cfg_bool("NX2NAVI_MEM_CENSUS");
if (!census_on || plat_xchg32(&census_busy, 1)) return;
CensusJob *j = malloc(sizeof *j);
if (j) *j = (CensusJob){c, frame};
if (!j || !plat_thread_spawn(census_main, j, "mem census", 0)) free(j), plat_xchg32(&census_busy, 0);
}
// After the vram line: the imported pools the GPU reached through a resolve (a binding, copy or texture placement) and
// those it never did, which hold host memory the driver counts in the host heap all the same (the 4.3 GiB imported
// at the title screen exceed the 3.9 GiB host heap of an integrated GPU with 16 GB of RAM); then the VRAM mirrors no
// submission synced for MIRROR_IDLE submissions, which a scene that needs the room could take back.
static void pool_census_log(uint64_t frame) {
uint64_t un = 0, ub = 0, nn = 0, nb = 0, mn = 0, mb = 0, in = 0, ib = 0;
uint32_t sub = __atomic_load_n(&msub, __ATOMIC_RELAXED); // the replay's; a log line may read it a submission late
plat_mutex_lock_shared(&pool_lock);
for (int i = 0; i < pool_index.n; i++) {
const PoolRec *r = pool_index.e[i].buf;
if (r->kind != POOL_IMPORT || !r->buf) continue;
uint64_t n = r->import_hi - r->import_lo;
if (__atomic_load_n(&r->used, __ATOMIC_RELAXED)) un++, ub += n;
else nn++, nb += n;
if (r->mstate != MS_ON) continue;
mn++, mb += n;
if (sub - r->acc_sub >= MIRROR_IDLE) in++, ib += n;
}
plat_mutex_unlock_shared(&pool_lock);
const double M = 1048576.0;
rt_log("mem: pools frame %llu: imported %llu %.0f MiB resolved by the GPU, %llu %.0f MiB never; mirrors %llu %.0f MiB, "
"idle for %d submissions %llu %.0f MiB", (unsigned long long)frame, (unsigned long long)un, ub / M,
(unsigned long long)nn, nb / M, (unsigned long long)mn, mb / M, MIRROR_IDLE, (unsigned long long)in, ib / M);
}
// After the heap lines: the driver's figure for the device-local heaps split by what holds it. VMA's blocks hold the
// textures with their own memory (live and released), the VRAM mirrors, the physical pools, the ASTC decoder's scratch,
// the headroom's spare blocks (their anchors), everything else VMA placed, and the free space inside the blocks, which
// the driver counts as used until a block empties. The rest of the driver's figure was never allocated through VMA:
// the swapchain, DLSS and frame generation (Streamline), compiled pipelines and the driver's own memory; DLSS's share
// as its statistics report it, when known.
static void vram_use_log(VkCtx *c, uint64_t frame, const VmaBudget *vb, const VkTexCensus *t) {
uint64_t usage = 0, blocks = 0, allocs = 0, phys = 0;
VkPhysicalDeviceMemoryBudgetPropertiesEXT bp = {.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_BUDGET_PROPERTIES_EXT};
if (c->has_memory_budget) {
VkPhysicalDeviceMemoryProperties2 p2 = {.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_PROPERTIES_2, .pNext = &bp};
vkGetPhysicalDeviceMemoryProperties2(c->phys, &p2);
}
for (uint32_t h = 0; h < c->mem_props.memoryHeapCount && h < VK_MAX_MEMORY_HEAPS; h++)
if (c->mem_props.memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT)
usage += c->has_memory_budget ? bp.heapUsage[h] : vb[h].statistics.blockBytes,
blocks += vb[h].statistics.blockBytes, allocs += vb[h].statistics.allocationBytes;
plat_mutex_lock_shared(&pool_lock);
for (int i = 0; i < nphys; i++)
if (phys_list[i]->alloc) phys += phys_list[i]->size;
plat_mutex_unlock_shared(&pool_lock);
uint64_t mir = (uint64_t)__atomic_load_n(&m_live_bytes, __ATOMIC_RELAXED), held = vk_headroom_held(),
scratch = vk_astc_scratch_bytes();
uint64_t known = t->bytes_live + t->bytes_released + mir + phys + held + scratch;
int64_t dlss = vk_dlaa_vram();
char ds[48] = "n/a";
if (dlss >= 0) snprintf(ds, sizeof ds, "%.0f MiB", dlss / 1048576.0);
const double M = 1048576.0;
rt_log("mem: vram by use frame %llu: in use %.0f MiB = VMA blocks %.0f MiB (textures %.0f, released textures %.0f, "
"VRAM mirrors %.0f, physical pools %.0f, ASTC scratch %.0f, headroom spare %.0f, other %.0f, free inside the "
"blocks %.0f) + outside VMA %.0f MiB (swapchain, DLSS and frame generation, pipelines, driver; DLSS reports %s)",
(unsigned long long)frame, usage / M, blocks / M, t->bytes_live / M, t->bytes_released / M, mir / M, phys / M,
scratch / M, held / M, allocs > known ? (allocs - known) / M : 0.0, blocks > allocs ? (blocks - allocs) / M : 0.0,
usage > blocks ? (usage - blocks) / M : 0.0, ds);
}
// With NX2NAVI_TRACE=perf every 60 frames: a line when the device-local memory VMA did not allocate (vram_use_log's
// "outside VMA") moved by OUTSIDE_STEP or more since the last line, so the log places it beside what did it. In one
// session it grew from 0.6 to 2.6 GiB at the title screen, in two steps of about 1 GiB while the boot's compute
// pipelines were being built (vk_pipeline.c, stage library prebuild), and DLSS and frame generation add theirs when
// they start.
#define OUTSIDE_STEP (256ull << 20)
static void outside_watch(uint64_t frame) {
VkCtx *c = vk_ctx();
if (!c || !c->has_memory_budget) return;
VkPhysicalDeviceMemoryBudgetPropertiesEXT bp = {.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_BUDGET_PROPERTIES_EXT};
VkPhysicalDeviceMemoryProperties2 p2 = {.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_PROPERTIES_2, .pNext = &bp};
vkGetPhysicalDeviceMemoryProperties2(c->phys, &p2);
VmaBudget vb[VK_MAX_MEMORY_HEAPS];
vmaGetHeapBudgets(c->vma, vb);
uint64_t usage = 0, blocks = 0;
for (uint32_t h = 0; h < c->mem_props.memoryHeapCount && h < VK_MAX_MEMORY_HEAPS; h++)
if (c->mem_props.memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT)
usage += bp.heapUsage[h], blocks += vb[h].statistics.blockBytes;
uint64_t out = usage > blocks ? usage - blocks : 0;
static uint64_t last;
static int seen;
if (seen && (out > last ? out - last : last - out) < OUTSIDE_STEP) return;
const double M = 1048576.0;
rt_log("mem: frame %llu: device memory outside VMA (swapchain, DLSS and frame generation, pipelines, driver) %s%.0f "
"MiB (%s%.0f MiB)", (unsigned long long)frame, seen ? "now " : "", out / M, out >= last ? "+" : "-",
(out >= last ? out - last : last - out) / M);
last = out, seen = 1;
}
// With NX2NAVI_TRACE=perf every 600 frames: what holds device-local memory, to see what grows over a long session.
// VMA's blocks and allocations on the device-local heaps, texture entries with their own memory, the VRAM mirrors and
// the objects waiting for their last submission to retire.
static void vram_log(uint64_t frame) {
VkCtx *c = vk_ctx();
if (!c) return;
VmaBudget vb[VK_MAX_MEMORY_HEAPS];
vmaGetHeapBudgets(c->vma, vb);
uint64_t blocks = 0, block_bytes = 0, allocs = 0, alloc_bytes = 0;
for (uint32_t h = 0; h < c->mem_props.memoryHeapCount && h < VK_MAX_MEMORY_HEAPS; h++)
if (c->mem_props.memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT)
blocks += vb[h].statistics.blockCount, block_bytes += vb[h].statistics.blockBytes,
allocs += vb[h].statistics.allocationCount, alloc_bytes += vb[h].statistics.allocationBytes;
VkTexCensus t;
vk_texture_census(&t);
const double M = 1048576.0;
rt_log("mem: vram frame %llu: vma %llu blocks %.0f MiB, %llu allocations %.0f MiB; textures live %llu (own memory %llu, "
"%.0f MiB), released %llu (own memory %llu, %.0f MiB); mirrors %lld %.0f MiB; releases pending %d",
(unsigned long long)frame, (unsigned long long)blocks, block_bytes / M, (unsigned long long)allocs,
alloc_bytes / M, (unsigned long long)t.live, (unsigned long long)t.own_live, t.bytes_live / M,
(unsigned long long)t.released, (unsigned long long)t.own_released, t.bytes_released / M, (long long)m_live_n,
m_live_bytes / M, dead_pending());
vk_texture_census_log(frame);
pool_census_log(frame);
heap_log(c, frame, vb);
vram_use_log(c, frame, vb, &t);
}
// Device class: INTEGRATED -> uma; a discrete GPU with a DEVICE_LOCAL heap of 1 GiB or more beside a
// heap that is not device-local (host memory) -> discrete; otherwise uma (import everywhere). Sets the mirror mode.
void vk_mem_class_log(const VkPhysicalDeviceProperties *p, const VkPhysicalDeviceMemoryProperties *m) {
int big_local = -1, host_heap = -1;
for (uint32_t h = 0; h < m->memoryHeapCount; h++) {
if (m->memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT) {
if (m->memoryHeaps[h].size >= (1ull << 30) && big_local < 0) big_local = (int)h;
} else if (host_heap < 0) host_heap = (int)h;
}
static const char *const types[] = {"other", "integrated", "discrete", "virtual", "cpu"};
int t = (int)p->deviceType;
const char *tn = t >= 0 && t <= 4 ? types[t] : "unknown";
const char *cls, *why;
if (p->deviceType == VK_PHYSICAL_DEVICE_TYPE_INTEGRATED_GPU) cls = "uma", why = "integrated gpu";
else if (p->deviceType != VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU) cls = "uma", why = "not a discrete gpu";
else if (big_local < 0) cls = "uma", why = "no device-local heap of 1 GiB or more";
else if (host_heap < 0) cls = "uma", why = "every heap is device-local";
else cls = "discrete", why = "device-local heap separate from host memory";
int discrete = cls[0] == 'd';
char line[2048];
size_t k = (size_t)snprintf(line, sizeof line, "mem: class %s (%s; device type %d %s, %s)", cls, why, t, tn, p->deviceName);
for (uint32_t h = 0; h < m->memoryHeapCount; h++)
k += (size_t)snprintf(line + k, k < sizeof line ? sizeof line - k : 0, "; heap %u %.0f MiB flags %#x%s%s", h,
m->memoryHeaps[h].size / 1048576.0, m->memoryHeaps[h].flags,
m->memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT ? " DEVICE_LOCAL" : "",
m->memoryHeaps[h].flags & VK_MEMORY_HEAP_MULTI_INSTANCE_BIT ? " MULTI_INSTANCE" : "");
rt_log("%s", line);
k = (size_t)snprintf(line, sizeof line, "mem: types");
for (uint32_t i = 0; i < m->memoryTypeCount; i++) {
VkMemoryPropertyFlags f = m->memoryTypes[i].propertyFlags;
k += (size_t)snprintf(line + k, k < sizeof line ? sizeof line - k : 0, "; %u heap %u flags %#x%s%s%s%s%s", i,
m->memoryTypes[i].heapIndex, f, f & VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT ? " DEVICE_LOCAL" : "",
f & VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT ? " HOST_VISIBLE" : "",
f & VK_MEMORY_PROPERTY_HOST_COHERENT_BIT ? " HOST_COHERENT" : "",
f & VK_MEMORY_PROPERTY_HOST_CACHED_BIT ? " HOST_CACHED" : "",
f & VK_MEMORY_PROPERTY_LAZILY_ALLOCATED_BIT ? " LAZY" : "");
}
rt_log("%s", line);
mirror_mode(discrete);
}
// ---------------------------------------------------------------- upload batches and deferred destruction
typedef struct { VkBuffer buf; VmaAllocation alloc; VkDeviceSize size; } Staging; // size: the VkBuffer's own
typedef struct { VkBuffer buf; VmaAllocation alloc; uint8_t *map; } StgChunk; // a staging arena chunk
typedef struct {
VkCommandBuffer cb;
VkFence fence;
uint64_t seq;
Staging *stg;
int nstg, cap;
StgChunk *chunk; // arena chunks this batch's uploads took ranges of (vk_staging_at), the last one open
int nchunk, cap_chunk;
VkDeviceSize chunk_used; // bytes of the last chunk handed out
} Batch;
enum { DEAD_IMAGE = 1, DEAD_IMAGE_VMA, DEAD_VIEW, DEAD_BUFFER_VMA, DEAD_HOST };
typedef struct { int kind; uint64_t handle; VmaAllocation alloc; uint64_t seq, tl; } Dead;
static PlatRMutex cmd_lock;
static PlatOnce cmd_once = PLAT_ONCE_INIT;
static VkCommandPool cmd_pool;
static Batch open_b;
static int open_active;
static Batch *inflight;
static int ninflight, cap_inflight;
static Batch *spare; // completed batches whose command buffer and fence are reused
static int nspare, cap_spare;
static uint64_t next_seq, done_seq; // batches handed out / completed (all batches up to done_seq)
static Dead *dead;
static int ndead, cap_dead;
static uint64_t tl_noted; // newest timeline value vk_sync.c reported submitted
static void cmd_init(void *p) {
(void)p;
return;
}
static void grow(void **arr, int *cap, size_t elem) {
*cap = *cap ? *cap * 2 : 16;
void *n = realloc(*arr, (size_t)*cap * elem);
if (!n) rt_fatal("vk: out of memory (upload bookkeeping)");
*arr = n;
}
// Staging buffers of a finished batch go back to a cache instead of the allocator: an upload-heavy scene creates and
// retires the same few sizes thousands of times a frame (a texture re-upload, a texel copy, an ASTC decode), and each
// vmaCreateBuffer/vmaDestroyBuffer pair walks the pool's block list under its lock. Bounded in count and bytes; a buffer
// that does not fit is destroyed as before. The batch is complete when this runs, so nothing reads the bytes again.
// Sizes are the VkBuffers' own: an allocation can be larger than its buffer, and a copy reads only the buffer's range.
// cmd_lock held (retire, and vk_staging between vk_upload_begin and vk_upload_end).
#define STG_CACHE 24
#define STG_CACHE_BYTES (192ull << 20)
static Staging stg_cache[STG_CACHE];
static int stg_ncache;
static VkDeviceSize stg_cache_bytes;
// The staging arena: an upload of at most ARENA_MAX bytes takes an ARENA_ALIGN-aligned range of a persistent mapped
// chunk (vk_staging_at) instead of a buffer of its own. The chunks belong to the open batch while it records and go
// back to the free list when its fence signals, so a frame's few hundred small texture uploads cost a pointer bump
// each instead of a vmaCreateBuffer and vmaDestroyBuffer under the allocator's lock. ARENA_KEEP free chunks stay; the
// rest are destroyed. 256-byte alignment covers every bufferOffset rule of a buffer-to-image copy. NX2NAVI_STG_ARENA=0
// gives every upload its own buffer again.
#define ARENA_CHUNK (8ull << 20)
#define ARENA_MAX (ARENA_CHUNK / 4)
#define ARENA_ALIGN 256ull
#define ARENA_KEEP 8
static StgChunk arena_free[ARENA_KEEP];
static int narena_free;
static uint64_t arena_made, arena_ranges;
static void free_staging(VkCtx *c, Batch *b) {
for (int i = 0; i < b->nstg; i++) {
if (stg_ncache < STG_CACHE && stg_cache_bytes + b->stg[i].size <= STG_CACHE_BYTES) {
stg_cache[stg_ncache++] = b->stg[i];
stg_cache_bytes += b->stg[i].size;
continue;
}
vmaDestroyBuffer(c->vma, b->stg[i].buf, b->stg[i].alloc);
}
b->nstg = 0;
for (int i = 0; i < b->nchunk; i++) {
if (narena_free < ARENA_KEEP) arena_free[narena_free++] = b->chunk[i];
else vmaDestroyBuffer(c->vma, b->chunk[i].buf, b->chunk[i].alloc);
}
b->nchunk = 0, b->chunk_used = 0;
}
// The cached staging buffer that fits n with the least waste: at most four times the request, so a small upload does not
// take a block-sized buffer another upload wanted.
static int staging_cache_take(VkCtx *c, VkDeviceSize n, Staging *out, void **map) {
int best = -1;
for (int i = 0; i < stg_ncache; i++) {
if (stg_cache[i].size < n || stg_cache[i].size > n * 4 + 65536) continue;
if (best < 0 || stg_cache[i].size < stg_cache[best].size) best = i;
}
if (best < 0) return 0;
VmaAllocationInfo ai;
vmaGetAllocationInfo(c->vma, stg_cache[best].alloc, &ai);
*out = stg_cache[best], *map = ai.pMappedData;
stg_cache_bytes -= stg_cache[best].size;
stg_cache[best] = stg_cache[--stg_ncache];
return 1;
}
// Batches complete in submission order (one queue): retire from the front while their fences are signalled.
static void retire(VkCtx *c, int wait_all) {
int k = 0;
for (; k < ninflight; k++) {
Batch *b = &inflight[k];
if (wait_all) vkWaitForFences(c->device, 1, &b->fence, VK_TRUE, UINT64_MAX);
else if (vkGetFenceStatus(c->device, b->fence) != VK_SUCCESS) break;
free_staging(c, b);
done_seq = b->seq;
if (nspare == cap_spare) grow((void **)&spare, &cap_spare, sizeof *spare);
spare[nspare++] = *b;
}
memmove(inflight, inflight + k, (size_t)(ninflight - k) * sizeof *inflight);
ninflight -= k;
}
static void collect_locked(VkCtx *c) {
retire(c, 0);
uint64_t tl = 0;
if (c->timeline) vkGetSemaphoreCounterValue(c->device, c->timeline, &tl);
int k = 0;
for (int i = 0; i < ndead; i++) {
Dead *d = &dead[i];
if (d->seq > done_seq || d->tl > tl) { dead[k++] = *d; continue; }
if (d->kind == DEAD_HOST) { // not a Vulkan object: nothing for the recorder's handle check
free((void *)(uintptr_t)d->handle);
continue;
}
rec_check_handle(d->kind == DEAD_VIEW ? REC_VIEW : d->kind == DEAD_BUFFER_VMA ? REC_BUFFER : REC_IMAGE, d->handle, 1,
d->tl, tl);
switch (d->kind) {
case DEAD_IMAGE: vkDestroyImage(c->device, (VkImage)(uintptr_t)d->handle, NULL); break;
case DEAD_IMAGE_VMA: vmaDestroyImage(c->vma, (VkImage)(uintptr_t)d->handle, d->alloc); break;
case DEAD_VIEW: vkDestroyImageView(c->device, (VkImageView)(uintptr_t)d->handle, NULL); break;
case DEAD_BUFFER_VMA: vmaDestroyBuffer(c->vma, (VkBuffer)(uintptr_t)d->handle, d->alloc); break;
}
}
ndead = k;
}
VkCommandBuffer vk_upload_begin(void) {
plat_once(&cmd_once, cmd_init, NULL);
plat_rmutex_lock(&cmd_lock);
VkCtx *c = vk_ctx();
if (!c) rt_fatal("vk: upload before vk_init");
if (open_active) return open_b.cb;
if (!cmd_pool) {
VkCommandPoolCreateInfo pi = {.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO,
.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT,
.queueFamilyIndex = c->gfx_family};
if (vkCreateCommandPool(c->device, &pi, NULL, &cmd_pool) != VK_SUCCESS) rt_fatal("vk: upload command pool");
}
retire(c, 0);
memset(&open_b, 0, sizeof open_b);
if (nspare) {
Batch s = spare[--nspare];
open_b.cb = s.cb, open_b.fence = s.fence, open_b.stg = s.stg, open_b.cap = s.cap;
open_b.chunk = s.chunk, open_b.cap_chunk = s.cap_chunk;
vkResetCommandBuffer(open_b.cb, 0);
vkResetFences(c->device, 1, &open_b.fence);
} else {
VkCommandBufferAllocateInfo ai = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO, .commandPool = cmd_pool,
.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY, .commandBufferCount = 1};
VkFenceCreateInfo fi = {.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO};
if (vkAllocateCommandBuffers(c->device, &ai, &open_b.cb) != VK_SUCCESS ||
vkCreateFence(c->device, &fi, NULL, &open_b.fence) != VK_SUCCESS)
rt_fatal("vk: upload command buffer");
}
VkCommandBufferBeginInfo bi = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO,
.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT};
vkBeginCommandBuffer(open_b.cb, &bi);
open_b.seq = ++next_seq;
vk_crumb(open_b.cb, VK_CRUMB_UPLOAD, open_b.seq);
open_active = 1;
return open_b.cb;
}
void vk_upload_end(void) { plat_rmutex_unlock(&cmd_lock); }
// Staging memory comes from a pool of its own that keeps STG_MIN_BLOCKS blocks from its first use (at boot). In the
// default pools VMA creates a block under the lock of its memory type (VmaBlockVector::Allocate), so a block made by
// the headroom thread (vk_headroom.c) or a burst of uploads held every staging allocation of the render thread: one
// took 47-50 ms, for a few bytes, in the first frame after a save load. Larger buffers than a block go to the default
// pool.
#define STG_BLOCK (64ull << 20)
#define STG_MIN_BLOCKS 2
static VmaPool stg_pool;
static int stg_pool_tried;
void *vk_staging(VkDeviceSize size, VkBuffer *buf, VmaAllocation *alloc) {
VkCtx *c = vk_ctx();
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .size = size ? size : 16,
.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT};
VmaAllocationCreateInfo aci = {.usage = VMA_MEMORY_USAGE_AUTO_PREFER_HOST,
.flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_RANDOM_BIT | VMA_ALLOCATION_CREATE_MAPPED_BIT};
if (!__atomic_load_n(&stg_pool_tried, __ATOMIC_ACQUIRE)) {
plat_once(&cmd_once, cmd_init, NULL);
plat_rmutex_lock(&cmd_lock);
if (!stg_pool_tried) {
uint32_t type;
VmaPoolCreateInfo pi = {.blockSize = STG_BLOCK, .minBlockCount = STG_MIN_BLOCKS};
if (vmaFindMemoryTypeIndexForBufferInfo(c->vma, &bi, &aci, &type) == VK_SUCCESS) {
pi.memoryTypeIndex = type;
if (vmaCreatePool(c->vma, &pi, &stg_pool) != VK_SUCCESS) stg_pool = NULL;
}
rt_log("vk: staging pool %s (%d blocks of %llu MiB kept)", stg_pool ? "created" : "unavailable: default pool",
STG_MIN_BLOCKS, (unsigned long long)(STG_BLOCK >> 20));
__atomic_store_n(&stg_pool_tried, 1, __ATOMIC_RELEASE);
}
plat_rmutex_unlock(&cmd_lock);
}
if (stg_pool && bi.size <= STG_BLOCK / 2) aci.pool = stg_pool;
if (!alloc) { // the batch owns it, so a cached one of a fitting size serves
Staging cs;
void *cmap;
if (staging_cache_take(c, bi.size, &cs, &cmap)) {
*buf = cs.buf;
if (open_b.nstg == open_b.cap) grow((void **)&open_b.stg, &open_b.cap, sizeof *open_b.stg);
open_b.stg[open_b.nstg++] = cs;
return cmap;
}
}
VmaAllocationInfo info;
VkBuffer b;
VmaAllocation a;
VkResult r = vmaCreateBuffer(c->vma, &bi, &aci, &b, &a, &info);
if (r != VK_SUCCESS) rt_fatal("vk: staging buffer of %llu bytes failed (%d)", (unsigned long long)size, (int)r);
vk_texstat_add(VK_TS_STG_NEW, 1, bi.size);
*buf = b;
if (alloc) *alloc = a; // the caller destroys it (vmaDestroyBuffer) after vk_upload_flush(1)
else {
if (open_b.nstg == open_b.cap) grow((void **)&open_b.stg, &open_b.cap, sizeof *open_b.stg);
open_b.stg[open_b.nstg++] = (Staging){b, a, bi.size};
}
return info.pMappedData;
}
// A range of the staging arena for an upload of the open batch (between vk_upload_begin and vk_upload_end): *buf and
// *off name it, and the batch owns it until its fence signals. Larger requests, and every request with
// NX2NAVI_STG_ARENA=0, get a buffer of their own at offset 0 (vk_staging).
void *vk_staging_at(VkDeviceSize size, VkBuffer *buf, VkDeviceSize *off) {
static int on = -1;
if (on < 0) {
on = cfg_bool("NX2NAVI_STG_ARENA");
if (!on) rt_log("vk: staging arena off (NX2NAVI_STG_ARENA=0): every upload stages in a buffer of its own");
}
*off = 0;
if (!on || !open_active || size > ARENA_MAX) return vk_staging(size, buf, NULL);
VkCtx *c = vk_ctx();
VkDeviceSize at = (open_b.chunk_used + ARENA_ALIGN - 1) & ~(ARENA_ALIGN - 1);
if (!open_b.nchunk || at + size > ARENA_CHUNK) {
StgChunk k;
if (narena_free) {
k = arena_free[--narena_free];
} else {
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .size = ARENA_CHUNK,
.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT};
VmaAllocationCreateInfo aci = {.usage = VMA_MEMORY_USAGE_AUTO_PREFER_HOST,
.flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_RANDOM_BIT |
VMA_ALLOCATION_CREATE_MAPPED_BIT};
if (stg_pool) aci.pool = stg_pool; // half of a staging block (STG_BLOCK)
VmaAllocationInfo info;
if (vmaCreateBuffer(c->vma, &bi, &aci, &k.buf, &k.alloc, &info) != VK_SUCCESS)
return vk_staging(size, buf, NULL);
k.map = info.pMappedData;
vk_texstat_add(VK_TS_STG_NEW, 1, ARENA_CHUNK);
if (arena_made++ < 16) rt_log("vk: staging arena chunk %llu made (%llu MiB)", (unsigned long long)arena_made,
(unsigned long long)(ARENA_CHUNK >> 20));
}
if (open_b.nchunk == open_b.cap_chunk) grow((void **)&open_b.chunk, &open_b.cap_chunk, sizeof *open_b.chunk);
open_b.chunk[open_b.nchunk++] = k;
at = 0;
}
StgChunk *k = &open_b.chunk[open_b.nchunk - 1];
open_b.chunk_used = at + (size ? size : 16);
arena_ranges++;
*buf = k->buf, *off = at;
return k->map + at;
}
uint64_t vk_upload_flush(int wait) {
uint64_t f0 = kern_now_ns();
plat_once(&cmd_once, cmd_init, NULL);
plat_rmutex_lock(&cmd_lock);
VkCtx *c = vk_ctx();
uint64_t seq = next_seq;
if (c && open_active) {
for (int i = 0; i < open_b.nstg; i++) vmaFlushAllocation(c->vma, open_b.stg[i].alloc, 0, VK_WHOLE_SIZE);
for (int i = 0; i < open_b.nchunk; i++)
vmaFlushAllocation(c->vma, open_b.chunk[i].alloc, 0, i + 1 < open_b.nchunk ? VK_WHOLE_SIZE : open_b.chunk_used);
vk_crumb(open_b.cb, VK_CRUMB_UPLOAD_END, open_b.seq);
vkEndCommandBuffer(open_b.cb);
VkSubmitInfo si = {.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO, .commandBufferCount = 1, .pCommandBuffers = &open_b.cb};
vk_queue_lock();
VkResult r = vkQueueSubmit(c->gfx_queue, 1, &si, open_b.fence);
vk_queue_unlock();
if (r != VK_SUCCESS) {
vk_device_fault_report(r);
rt_fatal("vk: upload submit failed (%d)", (int)r);
}
if (ninflight == cap_inflight) grow((void **)&inflight, &cap_inflight, sizeof *inflight);
inflight[ninflight++] = open_b;
open_active = 0;
}
if (c && wait) {
while (ninflight && done_seq < seq) {
vkWaitForFences(c->device, 1, &inflight[0].fence, VK_TRUE, UINT64_MAX);
retire(c, 0);
}
}
if (c) collect_locked(c);
plat_rmutex_unlock(&cmd_lock);
vk_texstat_time(VK_TS_FLUSH, f0);
return seq;
}
static void defer(int kind, uint64_t handle, VmaAllocation alloc) {
plat_once(&cmd_once, cmd_init, NULL);
plat_rmutex_lock(&cmd_lock);
if (ndead == cap_dead) grow((void **)&dead, &cap_dead, sizeof *dead);
// The open batch may record into the object: it counts as the newest user. The GPU must also be past every
// submission that names it. tl_noted covers the submitted ones; an object the replay lets go of while it records
// a submission (vk_exec_recording_fence) may be named by that one too, not yet submitted: vk_texture.c
// slot_resolve hands a draw a released entry, and entry_of destroys released entries at that address later in
// the same submission. Stamped with the last submitted value, the object went as soon as the GPU reached it,
// before the submission that names it ran, and with the recorder thread (vk_rec.c) behind, before the recorder
// made the call that names it. So every deferral is stamped with the submission being
// recorded. NX2NAVI_DEFER_STAMP=0 brings back the old stamp.
static int stamp = -1;
if (stamp < 0) stamp = rt_knob("DEFER_STAMP", "1") != 0;
uint64_t tl = tl_noted, rf = stamp ? vk_exec_recording_fence() : 0;
// vk_exec_recording_fence knows only the recording thread's own submission. A texture a guest thread releases (the
// synchronous replay: vk_texture_release, entry_of, sweep_released on the game's loader threads) can be named by
// the submission another thread records meanwhile: that one is issued + 1 (vk_sync_submit sets issued after it).
if (stamp && !rf) rf = vk_sync_issued() + 1;
if (rf > tl) tl = rf;
dead[ndead++] = (Dead){kind, handle, alloc, open_active ? open_b.seq : next_seq, tl};
if (kind != DEAD_HOST)
rec_check_handle(kind == DEAD_VIEW ? REC_VIEW : kind == DEAD_BUFFER_VMA ? REC_BUFFER : REC_IMAGE, handle, 0, tl,
vk_sync_fence_done());
plat_rmutex_unlock(&cmd_lock);
}
static int dead_pending(void) {
plat_once(&cmd_once, cmd_init, NULL);
plat_rmutex_lock(&cmd_lock);
int n = ndead;
plat_rmutex_unlock(&cmd_lock);
return n;
}
void vk_defer_image(VkImage img, VmaAllocation alloc) {
if (img) defer(alloc ? DEAD_IMAGE_VMA : DEAD_IMAGE, (uint64_t)(uintptr_t)img, alloc);
}
void vk_defer_view(VkImageView view) { if (view) defer(DEAD_VIEW, (uint64_t)(uintptr_t)view, NULL); }
void vk_defer_free(void *p) { if (p) defer(DEAD_HOST, (uint64_t)(uintptr_t)p, NULL); }
void vk_defer_buffer(VkBuffer buf, VmaAllocation alloc) {
if (buf) defer(DEAD_BUFFER_VMA, (uint64_t)(uintptr_t)buf, alloc);
}
void vk_mem_note_submit(uint64_t timeline_value) {
plat_once(&cmd_once, cmd_init, NULL);
plat_rmutex_lock(&cmd_lock);
if (timeline_value > tl_noted) tl_noted = timeline_value;
plat_rmutex_unlock(&cmd_lock);
}
void vk_mem_collect(void) {
VkCtx *c = vk_ctx();
if (!c) return;
plat_once(&cmd_once, cmd_init, NULL);
plat_rmutex_lock(&cmd_lock);
collect_locked(c);
plat_rmutex_unlock(&cmd_lock);
}
void vk_mem_stats(VkMemStats *s) {
memset(s, 0, sizeof *s);
plat_mutex_lock_shared(&pool_lock);
for (int i = 0; i < pool_index.n; i++) {
const PoolRec *r = pool_index.e[i].buf;
if (r->kind == POOL_IMPORT) s->imported_pools++;
else if (r->kind == POOL_PHYS) s->physical_pools++;
else s->virtual_pools++;
}
plat_mutex_unlock_shared(&pool_lock);
s->import_failures = (uint64_t)n_import_failures;
s->upload_batches = next_seq;
}
void vk_mem_shutdown(void) {
VkCtx *c = vk_ctx();
if (!c) return;
vk_upload_flush(1);
vk_queue_lock();
vkDeviceWaitIdle(c->device);
vk_queue_unlock();
vk_gpa_report();
vk_gpa_shutdown();
vk_texture_shutdown();
plat_rmutex_lock(&cmd_lock);
retire(c, 1);
uint64_t keep = tl_noted;
tl_noted = 0;
for (int i = 0; i < ndead; i++) dead[i].seq = 0, dead[i].tl = 0; // the device is idle
collect_locked(c);
tl_noted = keep;
for (int i = 0; i < nspare; i++) {
vkDestroyFence(c->device, spare[i].fence, NULL);
free(spare[i].stg), free(spare[i].chunk);
}
nspare = 0;
// The staging cache and the arena's free chunks live in stg_pool: before it goes
for (int i = 0; i < stg_ncache; i++) vmaDestroyBuffer(c->vma, stg_cache[i].buf, stg_cache[i].alloc);
stg_ncache = 0, stg_cache_bytes = 0;
for (int i = 0; i < narena_free; i++) vmaDestroyBuffer(c->vma, arena_free[i].buf, arena_free[i].alloc);
narena_free = 0;
if (cmd_pool) vkDestroyCommandPool(c->device, cmd_pool, NULL);
cmd_pool = VK_NULL_HANDLE;
if (stg_pool) vmaDestroyPool(c->vma, stg_pool); // after retire: no staging buffer is left in it
stg_pool = NULL;
stg_pool_tried = 0;
plat_rmutex_unlock(&cmd_lock);
plat_mutex_lock(&pool_lock);
for (int i = 0; i < pool_index.n; i++) destroy_rec(c, pool_index.e[i].buf);
pool_index.n = 0;
nm = npend = nwbl = nsnl = snl_copied = 0; // mirrors went with their pools
plat_mutex_lock(&snap_pool_lock);
for (int i = 0; i < nsnap_pool; i++) vmaDestroyBuffer(c->vma, snap_pool[i].buf, snap_pool[i].alloc);
nsnap_pool = 0, snap_pool_bytes = 0;
plat_mutex_unlock(&snap_pool_lock);
plat_mutex_lock(&shadow_lock);
nstouch = nspend = 0; // shadows too
s_live_n = s_live_bytes = 0;
plat_mutex_unlock(&shadow_lock);
for (int i = 0; i < nretiring; i++) destroy_rec(c, retiring[i]); // finalized ones whose callback never ran
nretiring = 0;
__atomic_add_fetch(&pool_gen, 1, __ATOMIC_RELEASE);
nphys = 0;
plat_mutex_unlock(&pool_lock);
plat_mutex_lock(&gen_lock);
for (uint64_t i = 0; i < GEN_CHUNKS; i++) free(gen_chunk[i]), gen_chunk[i] = NULL, free(cgen_chunk[i]), cgen_chunk[i] = NULL;
plat_mutex_unlock(&gen_lock);
}
Reference in New Issue View Git Blame Copy Permalink