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
73abd6577f76cbbc4ac869aed2e83ee6a462f085
NX2navi/runtime/vk_memory.c
chrg6 73abd6577f Try a smaller allocation before giving up on a texture upload when the video memory is used up. 
On an integrated Radeon 890M the ASTC scratch buffer failed to allocate and then a staging buffer ended the run. The scratch asks for 32 MiB whatever the upload needs, and the staging pool grows only in 64 MiB blocks of one memory type. When the device reports it is out of memory, the scratch is now tried again at the size the upload needs, and a staging buffer that does not fit the pool is tried again outside it with the same mapping, where a smaller block or another host-visible type may still have room. The menu's own upload flushes what it wrote, since such a buffer may lie in memory that is not coherent.

Built and run through the memory probe (96 GPU checks pass); the out-of-memory paths themselves were not reached here.

Co-Authored-By: Claude Fable 5.2 <noreply@anthropic.com>
2026-11-10 23:14:00 -05:00

3920 lines
209 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_lazy_choice.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 pm_frame; // replay: the frame (pframe) of pm_cur
uint8_t pm_cur, pm_prev; // replay: the places in their frame (vk_sub_in_frame, 7 the rest) of the submissions
// that synced it, in frame pm_frame and in the one before (NX2NAVI_PRESNAP_PLACE)
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
// shadow_lock: base copies of the pages a noted GPU write covers (merge_page), open addressing by page address;
// csb a power of two, page 0 an empty slot
struct SBase { uint64_t page; uint8_t *base; } *sb;
uint32_t nsb, csb;
// 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);
for (uint32_t i = 0; i < r->csb; i++) free(r->sb[i].base); // shadow base copies shadow_forget left (shutdown)
free(r->sb);
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 pool's pages go into the device address table: shaders that read a pointer from memory reach them there. An
// import and a shadow copy both have buf; only the import owns mem.
if (r->kind == POOL_IMPORT && 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 touch);
// touch: the GPU may read the pool's buffer in the submission being recorded, so its shadow goes on the touched list
// (shadow storage). A lookup (vk_mem_lookup) leaves it off: the caller reads guest memory itself or only asks what
// backs the address. Shaders that reach a pool through the device address table resolve nothing: the part they run in
// brings every pool of the table up to date (vk_mem_gpa_sync).
static int mem_resolve(uint64_t addr, uint64_t size, VkMemSpan *out, int touch) {
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 && touch && h->pool && !((PoolRec *)h->pool)->stouched) shadow_touch(h->pool);
if (resolve_check_on()) {
VkMemSpan s;
if (!resolve_locked(addr, size, &s, touch) || 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, touch);
}
int vk_mem_resolve(uint64_t addr, uint64_t size, VkMemSpan *out) { return mem_resolve(addr, size, out, 1); }
int vk_mem_lookup(uint64_t addr, uint64_t size, VkMemSpan *out) { return mem_resolve(addr, size, out, 0); }
// 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, int touch) {
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 && touch && !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_lookup(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_gen_lock_shared(void) { plat_mutex_lock_shared(&gen_lock); }
void vk_mem_gen_unlock_shared(void) { plat_mutex_unlock_shared(&gen_lock); }
void vk_mem_gen_lock(void) { plat_mutex_lock(&gen_lock); }
void vk_mem_gen_unlock(void) { plat_mutex_unlock(&gen_lock); }
int vk_mem_unchanged_why_locked(uint64_t addr, uint64_t size, uint32_t since, uint32_t epoch) {
if (!size || !epoch) return 4;
uint64_t lo = addr & ~0xfffull, hi = (addr + size + 0xfff) & ~0xfffull;
uint64_t r0 = lo >> SNAP_SHIFT, r1 = (hi - 1) >> SNAP_SHIFT;
if (r1 >= SNAP_REGIONS) return 4;
for (uint64_t r = r0; r <= r1; r++) { // REG_SHIFT is SNAP_SHIFT: reg_max is the newest stamp in the region
if (snap_at[r] != epoch) return 1;
if (snap_untracked[r]) return 2;
if (reg_max[r] > since) return 3;
}
return 0;
}
int vk_mem_unchanged_locked(uint64_t addr, uint64_t size, uint32_t since, uint32_t epoch) {
return !vk_mem_unchanged_why_locked(addr, size, since, epoch);
}
uint64_t vk_mem_snap_range_locked(uint64_t addr, uint64_t size, uint32_t epoch) {
if (!size || !epoch) return 0;
uint64_t lo = addr & ~0xfffull, hi = (addr + size + 0xfff) & ~0xfffull;
uint64_t r0 = lo >> SNAP_SHIFT, r1 = (hi - 1) >> SNAP_SHIFT;
if (r1 >= SNAP_REGIONS) return 0;
for (uint64_t r = r0; r <= r1; r++)
if (snap_at[r] != epoch) {
uint64_t q = snap_regions(r0, r1, epoch);
n_snap_calls += q;
return q;
}
return 0;
}
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. A part whose shaders read
// through the device address table gets every pool the table names brought up to date the same way (vk_mem_gpa_sync).
// - 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 the GPU's bytes back into guest memory.
// - Pages both write: a GPU write is noted with the whole range its binding names, and the CPU writes other fields of the
// same records meanwhile. Each page such a range covers keeps a base copy from its first note on, the bytes guest
// memory and the shadow last agreed on. A copy-in takes only the bytes where guest memory differs from the base (the
// CPU's writes) and the write-back at the fence only those where the shadow differs from it (the GPU's), and either
// moves the base along (merge_page). A byte both changed takes the CPU's value when a copy-in comes first and the
// GPU's when the fence does. Whole-range copies both ways lost the CPU's bytes: Ocarina writes a lighting parameter
// into a buffer whose binding its compute passes write every frame, the GPU only ever read zero there, the lighting
// coefficients and then the volumetric fog came out NaN, and the fog pass drew the 3D scene black.
// NX2NAVI_SHADOW=1/0 forces the mode on or off; the default is on for Linux and Wine.
#define SB_PAGE 4096ull
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 uint64_t s_sb_pages; // shadow_lock: base copies kept (sbase_make)
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 uint64_t s_wb_seen, s_wb_ticks; // shadow_lock: bytes the fence write-backs compared, their time
static uint64_t s_q_ticks; // shadow_lock: the copy-ins' write-watch queries
static int wb_no_reup(void);
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);
}
// ---- base copies of the pages both write ("Pages both write" above). shadow_lock held throughout.
static uint32_t sbase_hash(uint64_t pg) { return (uint32_t)((pg >> 12) * 0x9E3779B97F4A7C15ull >> 32); }
// The base copy of guest page pg of r, NULL when it has none.
static uint8_t *sbase_find(const PoolRec *r, uint64_t pg) {
if (!r->nsb) return NULL;
for (uint32_t i = sbase_hash(pg) & (r->csb - 1);; i = (i + 1) & (r->csb - 1)) {
if (r->sb[i].page == pg) return r->sb[i].base;
if (!r->sb[i].page) return NULL;
}
}
static void sbase_insert(struct SBase *t, uint32_t cap, uint64_t pg, uint8_t *base) {
uint32_t i = sbase_hash(pg) & (cap - 1);
while (t[i].page) i = (i + 1) & (cap - 1);
t[i].page = pg, t[i].base = base;
}
// Page pg of r gets a base copy unless it has one. Taken from the shadow: no GPU write was noted on the page before, so
// the shadow holds what guest memory held at its last copy-in, and guest memory differs from it by the CPU's writes since.
static void sbase_make(PoolRec *r, uint64_t pg) {
if (sbase_find(r, pg)) return;
if ((r->nsb + 1) * 4 > r->csb * 3) { // at most three quarters full
uint32_t cap = r->csb ? r->csb * 2 : 16;
struct SBase *t = calloc(cap, sizeof *t);
if (!t) rt_fatal("vk: out of memory for the shadow's base copies");
for (uint32_t i = 0; i < r->csb; i++)
if (r->sb[i].page) sbase_insert(t, cap, r->sb[i].page, r->sb[i].base);
free(r->sb);
r->sb = t, r->csb = cap;
}
uint8_t *b = malloc(SB_PAGE);
if (!b) rt_fatal("vk: out of memory for the shadow's base copies");
memcpy(b, r->smap + (pg - r->import_lo), SB_PAGE);
sbase_insert(r->sb, r->csb, pg, b);
r->nsb++;
s_sb_pages++;
}
static void sbase_free(PoolRec *r) {
for (uint32_t i = 0; i < r->csb; i++) free(r->sb[i].base);
free(r->sb);
s_sb_pages -= r->nsb;
r->sb = NULL, r->nsb = r->csb = 0;
}
// 0xff in each byte of x that is not zero, 0 in the others.
static inline uint64_t byte_mask(uint64_t x) {
x |= x >> 4, x |= x >> 2, x |= x >> 1;
return (x & 0x0101010101010101ull) * 0xff;
}
// The bytes of v that m selects into dst (little-endian), the other bytes of dst left as they are: the other side may
// be writing them at the same time.
static inline void store_bytes(uint8_t *dst, uint64_t v, uint64_t m) {
if (m == ~0ull) {
memcpy(dst, &v, 8);
return;
}
for (int k = 0; k < 8; k++)
if (m >> (8 * k) & 0xff) dst[k] = (uint8_t)(v >> (8 * k));
}
// n bytes of one page with base copy b: the bytes where src differs from the base go into dst and into the base. src
// and dst are guest memory and the shadow, one way or the other; each src byte that moves is read once, as its side
// may be writing it. Returns the bytes moved.
static uint64_t merge_page(uint8_t *dst, const uint8_t *src, uint8_t *b, uint64_t n) {
uint64_t moved = 0, i = 0;
for (; i < n && ((uintptr_t)(src + i) & 7); i++) {
uint8_t v = ((const volatile uint8_t *)src)[i];
if (v != b[i]) dst[i] = b[i] = v, moved++;
}
for (; i + 8 <= n; i += 8) {
// Most of a page is unchanged: a 64-byte block equal to its base is passed over with one vector compare. A
// byte its side writes after the compare is a later write, which the next copy-in or write-back takes.
if (!((uintptr_t)(src + i) & 63) && i + 64 <= n) {
uint64_t x = 0;
for (int k = 0; k < 64; k += 8) {
uint64_t sv, ov;
memcpy(&sv, src + i + k, 8), memcpy(&ov, b + i + k, 8);
x |= sv ^ ov;
}
if (!x) {
i += 56;
continue;
}
}
uint64_t s = *(const volatile uint64_t *)(src + i), o;
memcpy(&o, b + i, 8);
uint64_t m = byte_mask(s ^ o);
if (!m) continue;
store_bytes(dst + i, s, m);
memcpy(b + i, &s, 8); // the bytes m leaves out are equal in both
moved += (uint64_t)__builtin_popcountll(m) / 8;
}
for (; i < n; i++) {
uint8_t v = ((const volatile uint8_t *)src)[i];
if (v != b[i]) dst[i] = b[i] = v, moved++;
}
return moved;
}
// 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);
sbase_free(r);
plat_mutex_unlock(&shadow_lock);
}
// shadow_lock held: copies guest [a, b) into the shadow; of a page with a base copy only the CPU's bytes (merge_page).
static void shadow_copy(PoolRec *r, uint64_t a, uint64_t b) {
while (a < b) {
uint64_t e = b;
uint8_t *base = NULL;
if (r->nsb) { // the plain run up to the next page with a base copy
for (e = a; e < b; e = (e & ~(SB_PAGE - 1)) + SB_PAGE)
if ((base = sbase_find(r, e & ~(SB_PAGE - 1)))) break;
if (e > b) e = b;
if (e == a) { // a starts on such a page: its part of [a, b)
uint64_t pg = a & ~(SB_PAGE - 1), pe = pg + SB_PAGE < b ? pg + SB_PAGE : b;
s_copied += merge_page(r->smap + (a - r->import_lo), (const uint8_t *)(uintptr_t)a, base + (a - pg), pe - a);
a = pe;
continue;
}
}
memcpy(r->smap + (a - r->import_lo), (const void *)(uintptr_t)a, e - a);
s_copied += e - a;
a = e;
}
}
// shadow_lock and gen_lock held: stamps the pool's written pages and copies those newer than its shadow. The pool's
// regions are queried through the write-watch snapshots of the guest submission (snap_regions, epoch), which the
// texture checks and the mirror syncs share: each 2 MiB region at most once per guest submission. A query of every
// touched pool at every part's submit was about 380 write-watch calls and 10 ms a frame of the submitting thread in
// Kokiri Forest, while the GPU waited for the next part (RTX 4090, 2026-11-09). A page the CPU writes after its
// region's snapshot goes in at the next guest submission's, as it does into a VRAM mirror.
// cpu (wb_no_reup): the pages are read from the tables without the stamps of a mirror's write-back (stamp_page_as),
// since that write-back went through the shadow, which holds those bytes already; with the other tables every storage
// buffer the game's compute shaders wrote was compared again at the next copy-in.
static void shadow_sync_pool(PoolRec *r, uint32_t epoch, int cpu) {
uint64_t lo = r->import_lo, hi = r->import_hi;
uint64_t r0 = lo >> SNAP_SHIFT, r1 = (hi - 1) >> SNAP_SHIFT;
int tracked = 1;
uint64_t t0 = qpc();
if (epoch && r1 < SNAP_REGIONS) {
s_queries += snap_regions(r0, r1, epoch);
for (uint64_t k = r0; k <= r1 && tracked; k++) tracked = !snap_untracked[k];
} else { // outside the snapshot regions, or before the first guest submission: one query per unfrozen run
uint32_t g = ++gen_now;
for (uint64_t a = lo; a < hi;) {
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;
}
}
s_q_ticks += qpc() - t0;
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);
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 = (cpu ? cgen_chunk : gen_chunk)[a >> GEN_CHUNK_SHIFT]; // a 2 MiB region lies in one chunk
if ((cpu ? creg_max : reg_max)[k] <= since || !ch) {
if (run) shadow_copy(r, run, a), 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), run = 0;
}
}
if (run) shadow_copy(r, run, hi);
}
void vk_mem_shadow_sync(void) {
if (!shadow_on) return;
uint64_t t0 = qpc();
uint32_t epoch = vk_texture_epoch();
int cpu = vk_mirror_on && wb_no_reup(); // before the locks: its first call logs
plat_mutex_lock(&shadow_lock);
if (nstouch) {
plat_mutex_lock(&gen_lock);
for (int i = 0; i < nstouch; i++) shadow_sync_pool(stouch[i], epoch, cpu), 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);
}
// Record time: the GPU may write [addr, addr+size) of CPU-visible memory. With shadow storage the pages it covers get
// their base copies before the submission's copy-in (vk_mem_shadow_sync) can reach them.
void vk_mem_gpu_pending(uint64_t addr, uint64_t size) {
if (!shadow_on || !size) return;
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)
for (uint64_t pg = addr & ~(SB_PAGE - 1); pg < addr + size; pg += SB_PAGE) sbase_make(r, pg);
plat_mutex_unlock(&shadow_lock);
plat_mutex_unlock_shared(&pool_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;
}
// NX2NAVI_SHADOW_BULK_WRITE=1 (default 0): before a write-back stores into guest memory, the write watch's protection
// of the whole range is cleared in one call (plat_mem_prepare_write). Where the watch protects pages (Linux), the
// first store into each page otherwise takes a fault of its own; the pages count as written either way.
static int shadow_bulk_write(void) {
static int on = -1;
if (on < 0) {
on = cfg_bool("NX2NAVI_SHADOW_BULK_WRITE");
if (on) rt_log("mem: shadow write-backs clear the write watch's protection of their range in one call "
"(NX2NAVI_SHADOW_BULK_WRITE)");
}
return on;
}
// Completion thread: the GPU's bytes of [addr, addr+size) from the shadow into guest memory (merge_page against the
// base copies; a page without one whole). Returns the bytes moved.
static uint64_t shadow_write_back(uint64_t addr, uint64_t size) {
uint64_t moved = 0;
int bulk = shadow_bulk_write(); // before the locks: its first call logs
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) {
uint64_t t0 = qpc();
mem_thaw(addr, size);
if (bulk) plat_mem_prepare_write(addr, size);
for (uint64_t a = addr, end = addr + size; a < end;) {
uint64_t pg = a & ~(SB_PAGE - 1), pe = pg + SB_PAGE < end ? pg + SB_PAGE : end;
uint8_t *base = sbase_find(r, pg), *s = r->smap + (a - r->import_lo);
if (base) moved += merge_page((uint8_t *)(uintptr_t)a, s, base + (a - pg), pe - a);
else memcpy((void *)(uintptr_t)a, s, pe - a), moved += pe - a; // a write never noted as pending
a = pe;
}
s_wb_bytes += moved, s_wb_seen += size, s_wb_ticks += qpc() - t0;
}
plat_mutex_unlock(&shadow_lock);
plat_mutex_unlock_shared(&pool_lock);
return moved;
}
static void gpu_written(uint64_t addr, uint64_t size, int back) {
// back: the bytes are in a VRAM mirror, and its write-back brings them into the shadow; that submission's
// shadow_wb_cb, after the replay's own callbacks, then copies them into guest memory.
if (shadow_on && size && !back) shadow_write_back(addr, size);
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); }
static int wb_lazy(void);
static int lazy_clean(uint64_t lo, uint64_t hi, void *ctx);
void vk_gpu_written_back(uint64_t addr, uint64_t size) {
gpu_written(addr, size, 1);
// NX2NAVI_WB_PROBE: the bytes are in guest memory now; note who reads them next. With shadow storage they get there
// later, in the submission's shadow_wb_cb, whose thaw would end the probe at once: that job arms it instead.
if (!shadow_on) mem_probe(addr, size);
VkMemSpan sp;
if (size && wb_lazy() && vk_mem_lookup(addr, size, &sp) && sp.pool) { // or newer ones stay in the mirror from here on
PoolRec *r = sp.pool;
if (r->mstate == MS_ON && r->mbuf && r->buf && addr >= r->import_lo && addr + size <= r->import_hi)
mem_lazy_arm(addr, size, (uint64_t)(uintptr_t)r->mbuf, (uint64_t)(uintptr_t)r->buf, addr - r->import_lo, r);
}
}
int vk_mem_mirrored(uint64_t addr, uint64_t size) {
VkMemSpan sp;
if (!vk_mirror_on || !size || !vk_mem_lookup(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), %.2f ms (queries %.2f); %.1f MiB written back of %.1f MiB the GPU may have written, %.2f ms; %llu "
"pages the GPU writes keep a base copy",
(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, qpc_ms(s_ticks),
qpc_ms(s_q_ticks), s_wb_bytes / 1048576.0, s_wb_seen / 1048576.0, qpc_ms(s_wb_ticks), (unsigned long long)s_sb_pages);
s_syncs = s_queries = s_copied = s_untracked = s_wb_bytes = s_ticks = s_wb_seen = s_wb_ticks = s_q_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 uint32_t pframe; // replay: frames begun (vk_mem_mirror_begin: a place not after the last one's)
static int pplace = 99; // replay: the place of the last submission vk_mem_mirror_begin saw
// NX2NAVI_PRESNAP_PLACE=1 (default), 0, ab: a submission's pre-snap (vk_mem_mirror_begin) takes the write-watch
// snapshots of the pools the submission at its place in the last frame synced, not those the submission before it
// synced. In Kokiri Forest the frame's first submission (#0) and its fourth (#3) use no pool, and their pre-snaps
// covered the pools of the submission before them, #5's and #2's: 36.8 queries over 282 MiB a frame, 0.23-0.31 ms; then
// #2 and #5, whose submissions before them use no pool either, queried their pools one at a time at their first use,
// 63.5 queries over 254 MiB.
static int presnap_place(void) {
static int mode = -1;
if (mode < 0) mode = rt_knob("PRESNAP_PLACE", "1");
return RT_KNOB_ON(mode);
}
static int sub_place(void) { return vk_sub_in_frame >= 0 && vk_sub_in_frame < 7 ? vk_sub_in_frame : 7; }
static void pm_note(PoolRec *r) { // replay: r synced in the submission being recorded
if (r->pm_frame != pframe) {
r->pm_prev = r->pm_frame + 1 == pframe ? r->pm_cur : 0;
r->pm_cur = 0, r->pm_frame = pframe;
}
r->pm_cur |= (uint8_t)(1u << sub_place());
}
static uint8_t pm_last(const PoolRec *r) { // the places that synced r in the last frame
return r->pm_frame == pframe ? r->pm_prev : r->pm_frame + 1 == pframe ? r->pm_cur : 0;
}
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 MIRROR_SHARE 2 // the mirrors' share of the budget that the budget check never refuses: one in this many
#define RETRY_SOON 16
static int mirror_evict_on = -1;
static int64_t m_evicted_n, m_evicted_bytes, m_retries;
static volatile int64_t m_over_n, m_over_bytes; // mirrors made although the budget check failed (mirror_create)
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;
mem_lazy_release(r->import_lo, r->import_hi - r->import_lo); // what only the mirror holds comes down first
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;
}
// Over the budget still. The mirrors keep their place up to MIRROR_SHARE of the budget whatever else fills the
// VRAM: on a 6 GiB card a loaded area's textures alone pass the budget before the game's buffers are first
// used, every pool was refused, and the GPU read vertices, indices and storage over PCIe at each draw: 70 to
// 250 ms of GPU time a frame, 5 to 13 fps, in two reports from RTX 3060 Laptop cards. A mirror the driver cannot
// keep in VRAM is paged to system memory, where the imported pool already is, so a mirror never reads slower
// than its pool, and what is drawn from every frame is not what the system pages out first. The textures are
// still asked for the room. Past that share a pool is refused as before.
if (used + n + budget / 8 > budget) {
uint64_t over = used + n + budget / 8 - budget;
if (evict && (uint64_t)m_live_bytes + n <= budget / MIRROR_SHARE) {
vk_texture_want_room(over);
plat_inc64(&m_over_n);
plat_add64(&m_over_bytes, (int64_t)n);
if (m_over_n <= 4 || !(m_over_n & 255))
rt_log("mem: mirror of %.1f MiB made over the VRAM budget (%.0f of %.0f MiB in use, mirrors %.0f MiB): "
"%lld such so far, %.0f MiB; idle textures are asked for the room",
n / 1048576.0, used / 1048576.0, budget / 1048576.0, m_live_bytes / 1048576.0,
(long long)m_over_n, m_over_bytes / 1048576.0);
} else {
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(over);
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;
pm_note(r);
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;
}
// ---- shadow storage: the device address table's readers
// A shader that follows a pointer through the table (vk_gpa.c) can reach any pool the table names, and the replay
// never resolves that pool, so nothing put it on the touched list and its shadow kept the bytes of its last use as a
// buffer. vk_exec_draw.c marks a part whose pipelines declare the table (vk_mem_gpa_use: at the draw or dispatch, not
// where descriptors are written, which are reused), and the part's prologue (vk_sync.c) calls vk_mem_gpa_sync: the
// shadows of every pool in the table get the pages the CPU wrote, and those of them that have a VRAM mirror, which
// the table then names, queue the same pages for upload unless this submission synced the mirror already. No mirror
// is made for a reader of the table. Without shadow storage nothing is done: an imported pool is guest memory, and
// mirrors stay unsynced for the table's readers (vk_exec_draw.c ssbo_info).
static int gpa_part_used; // replay thread: the part being recorded uses the table
void vk_mem_gpa_use(void) { gpa_part_used = 1; }
int vk_mem_shadow_storage(void) { return shadow_on; }
static int gpa_shadow_pool(const PoolRec *r) { return r->kind == POOL_IMPORT && r->gpa_rec && r->smap; }
void vk_mem_gpa_sync(void) {
if (!gpa_part_used) return;
gpa_part_used = 0;
if (!shadow_on) return;
uint64_t t0 = qpc(), run0 = 0, run1 = 0;
uint32_t epoch = vk_texture_epoch();
int cpu = vk_mirror_on && wb_no_reup(), have = 0; // before the locks: its first call logs
plat_mutex_lock_shared(&pool_lock);
plat_mutex_lock(&shadow_lock);
plat_mutex_lock(&gen_lock);
// one write-watch query per run of regions the pools cover (the index is sorted by address) instead of one a pool;
// shadow_sync_pool then finds its regions snapped
for (int i = 0; epoch && i < pool_index.n; i++) {
const PoolRec *r = pool_index.e[i].buf;
if (!gpa_shadow_pool(r)) continue;
uint64_t a = r->import_lo >> SNAP_SHIFT, b = (r->import_hi - 1) >> SNAP_SHIFT;
if (b >= SNAP_REGIONS) continue;
if (have && a <= run1 + 1) {
if (b > run1) run1 = b;
} else {
if (have) s_queries += snap_regions(run0, run1, epoch);
run0 = a, run1 = b, have = 1;
}
}
if (have) s_queries += snap_regions(run0, run1, epoch);
for (int i = 0; i < pool_index.n; i++) {
PoolRec *r = pool_index.e[i].buf;
if (gpa_shadow_pool(r)) shadow_sync_pool(r, epoch, cpu), s_syncs++;
}
plat_mutex_unlock(&gen_lock);
// The uploads read the shadows just brought up to date, and the prologue records them ahead of the part, after
// the write-backs of the submissions before it.
for (int i = 0; i < pool_index.n; i++) {
PoolRec *r = pool_index.e[i].buf;
if (gpa_shadow_pool(r) && r->mstate == MS_ON && r->acc_sub != msub) mirror_sync(r);
}
s_ticks += qpc() - t0;
plat_mutex_unlock(&shadow_lock);
plat_mutex_unlock_shared(&pool_lock);
}
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[5];
for (int i = 0; i < 5; 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 = 5, .pBindings = b};
VkPushConstantRange pr = {VK_SHADER_STAGE_COMPUTE_BIT, 0, 16};
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. host: 0 device memory (snapshots), CHUNK_LIST host memory the CPU writes (range lists),
// CHUNK_MARKS host memory the CPU reads back (the merge's page marks, shadow_wb_cb).
enum { CHUNK_LIST = 1, CHUNK_MARKS = 2 };
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 == CHUNK_MARKS ? VMA_ALLOCATION_CREATE_HOST_ACCESS_RANDOM_BIT |
VMA_ALLOCATION_CREATE_MAPPED_BIT
: 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. marks: the chunk
// of the page marks, its first at marks_off, or -1 when the dispatch marks none (the range list stands in for the
// binding, which the shader then does not write).
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, int marks, uint64_t marks_off) {
uint64_t npg = ((base & (SB_PAGE - 1)) + span + SB_PAGE - 1) / SB_PAGE;
VkDescriptorBufferInfo bi[5] = {{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},
marks >= 0 ? (VkDescriptorBufferInfo){chunks[marks].buf, marks_off, npg * 4}
: (VkDescriptorBufferInfo){chunks[list].buf, list_off, (uint64_t)n * 16}};
VkWriteDescriptorSet w[5];
for (int i = 0; i < 5; 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, 5, w); // in call order (vk_rec.h)
uint32_t pc[4] = {(uint32_t)n, total, (uint32_t)(base & (SB_PAGE - 1)), marks >= 0};
vkCmdPushConstants(cb, WM.layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof pc, pc);
vkCmdDispatch(cb, (total + 63) / 64, 1, 1);
}
// ---- shadow storage: the GPU's writes into VRAM mirrors back into guest memory
// The mirror's write-back (vk_mem_mirror_end) merges or copies them into the shadow, and the submission's job, added to
// its fence after the replay's own callbacks, copies them on into guest memory: of a merge dispatch only the pages its
// shader marked (wb_merge.comp), of a copy every page, compared with its base copy (merge_page). The fence callback of
// each noted range (gpu_written) compared all of it: 65 MiB of storage buffers a frame in Kokiri Forest, 9 ms a frame
// of the completion thread, behind which the game's fence waits sat (RTX 4090, 2026-11-09).
typedef struct { uint64_t addr, size; uint32_t mark0; } SwbRange; // guest range; its first page's mark, or SWB_ALL
#define SWB_ALL UINT32_MAX
typedef struct { VmaAllocation marks_alloc; const uint32_t *marks; int n; SwbRange g[]; } SwbJob;
static SwbRange *swl; // replay: this submission's merge dispatches
static int nswl, cswl;
static int swl_marks = -1; // replay: their page marks' chunk, -1 none
static void swl_add(uint64_t addr, uint64_t size, uint32_t mark0) {
if (nswl == cswl) grow((void **)&swl, &cswl, sizeof *swl);
swl[nswl++] = (SwbRange){addr, size, mark0};
}
// Completion thread: [a, b) into guest memory, and the pages that changed stamped as the mirror's write-back.
static void swb_run(uint64_t a, uint64_t b) {
if (b <= a || !shadow_write_back(a, b - a) || nvn_gpu_only(a)) return;
int back = wb_no_reup();
plat_mutex_lock(&gen_lock);
uint32_t g = ++gen_now;
for (uint64_t x = a & ~0xfffull; x < b; x += 1ull << GEN_PAGE_SHIFT) stamp_page_as(x, g, back);
plat_mutex_unlock(&gen_lock);
}
static void shadow_wb_cb(void *arg) { // completion thread
SwbJob *j = arg;
if (j->marks_alloc) vmaInvalidateAllocation(vk_ctx()->vma, j->marks_alloc, 0, VK_WHOLE_SIZE);
for (int i = 0; i < j->n; i++) {
const SwbRange *g = &j->g[i];
uint64_t end = g->addr + g->size;
if (g->mark0 == SWB_ALL) {
swb_run(g->addr, end);
continue;
}
uint64_t run = 0; // start of the run of marked pages; guest address 0 is never a pool's
uint32_t k = g->mark0;
for (uint64_t pg = g->addr & ~(SB_PAGE - 1); pg < end; pg += SB_PAGE, k++) {
uint64_t a = pg > g->addr ? pg : g->addr;
if (j->marks[k]) {
if (!run) run = a;
} else if (run) {
swb_run(run, a);
run = 0;
}
}
if (run) swb_run(run, end);
}
// NX2NAVI_WB_PROBE (vk_gpu_written_back): every range is in guest memory now; a later range's thaw would end the
// probe of a region an earlier one armed, so the probes come last
for (int i = 0; i < j->n; i++) mem_probe(j->g[i].addr, j->g[i].size);
free(j);
}
// vk_mem_mirror_end, before the snapshot chunks go back at the fence (snap_free reuses the marks' chunk): the job of
// this submission's merge dispatches and copies.
static void swb_job(uint64_t fence) {
SwbJob *j = malloc(sizeof *j + sizeof(SwbRange) * (size_t)(nswl + nwbl));
if (!j) rt_fatal("vk: out of memory");
j->marks_alloc = NULL, j->marks = NULL;
if (swl_marks >= 0) {
VmaAllocationInfo info;
vmaGetAllocationInfo(vk_ctx()->vma, chunks[swl_marks].alloc, &info);
j->marks_alloc = chunks[swl_marks].alloc, j->marks = info.pMappedData;
}
memcpy(j->g, swl, sizeof *swl * (size_t)nswl);
j->n = nswl;
for (int i = 0; i < nwbl; i++) j->g[j->n++] = (SwbRange){wbl[i].r->import_lo + wbl[i].off, wbl[i].size, SWB_ALL};
nswl = 0, swl_marks = -1;
vk_completion_add(fence, shadow_wb_cb, j);
}
// 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). With shadow storage each
// dispatch also marks the pages it changed, for swb_job.
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);
struct MGroup { int i, j; uint64_t base, end; uint32_t total; } *gr = malloc(sizeof *gr * (size_t)nsnl);
if (!s || !gr) rt_fatal("vk: out of memory");
memcpy(s, snl, sizeof *s * (size_t)nsnl);
qsort(s, (size_t)nsnl, sizeof *s, snap_cmp);
// the dispatches: ranges of one pool and snapshot chunk, within MERGE_SPAN bytes and MERGE_WORDS words
int ng = 0;
uint64_t nmarks = 0; // bytes of page marks, each dispatch's at an aligned offset
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;
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;
total += (uint32_t)(s[j].size / 4);
}
gr[ng++] = (struct MGroup){i, j, base, end, total};
nmarks += (((base & (SB_PAGE - 1)) + end - base + SB_PAGE - 1) / SB_PAGE * 4 + align - 1) & ~(align - 1);
i = j;
}
// 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, CHUNK_LIST);
if (list < 0) {
free(s), free(gr);
return 0;
}
chunks[list].used = UINT64_MAX; // not a snapshot chunk
int marks = shadow_on ? chunk_new(nmarks, CHUNK_MARKS) : -1;
uint32_t *mm = NULL;
if (marks >= 0) {
chunks[marks].used = UINT64_MAX;
VmaAllocationInfo mi;
vmaGetAllocationInfo(c->vma, chunks[marks].alloc, &mi);
mm = mi.pMappedData;
memset(mm, 0, nmarks);
vmaFlushAllocation(c->vma, chunks[marks].alloc, 0, VK_WHOLE_SIZE);
swl_marks = marks;
}
VmaAllocationInfo li;
vmaGetAllocationInfo(c->vma, chunks[list].alloc, &li);
uint32_t *lm = li.pMappedData;
uint64_t lo = 0, mo = 0;
int rec = 0;
vkCmdBindPipeline(cb, VK_PIPELINE_BIND_POINT_COMPUTE, WM.pipe);
for (int k = 0; k < ng; k++) {
const struct MGroup *d = &gr[k];
lo = (lo + align - 1) & ~(align - 1);
uint64_t l0 = lo;
uint32_t first = 0; // the range's first invocation
for (int j = d->i; j < d->j; j++) {
uint32_t *q = lm + lo / 4;
q[0] = (uint32_t)((s[j].off - d->base) / 4), q[1] = (uint32_t)(s[j].soff / 4), q[2] = (uint32_t)(s[j].size / 4);
q[3] = first;
first += (uint32_t)(s[j].size / 4);
lo += 16;
mrep.merge_bytes += s[j].size;
}
merge_dispatch(cb, s + d->i, d->j - d->i, d->base, d->end - d->base, list, l0, d->total, marks, mo);
if (shadow_on) {
uint64_t ga = s[d->i].r->import_lo + d->base;
if (marks >= 0) swl_add(ga, d->end - d->base, (uint32_t)(mo / 4));
else
for (int j = d->i; j < d->j; j++) swl_add(s[j].r->import_lo + s[j].off, s[j].size, SWB_ALL);
}
mo += (((d->base & (SB_PAGE - 1)) + d->end - d->base + SB_PAGE - 1) / SB_PAGE * 4 + align - 1) & ~(align - 1);
mrep.merge_dispatches++;
rec += 3;
}
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), free(gr);
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);
}
// ---- lazy write-back (NX2NAVI_WB_LAZY; the default by platform and driver, wb_lazy_choice.h; memory.c "Lazy write-back")
// A range a shader wrote goes back to guest memory once and is then armed (vk_gpu_written_back, mem_lazy_arm). While it
// stays armed, gpu_wrote leaves the later writes of its pages in the mirror: no snapshot, no merge, no copy. When the
// CPU touches a page, memory.c calls lazy_fetch with the mirror and the imported buffer of the range: a copy of its
// own, sent after every part handed over so far (the held ones first), and waited for. The part that is still being
// recorded when that happens has already left its writes out: at its end (lazy_part_check, from vk_mem_mirror_end or
// from the prologue of a part split off earlier) it writes back, as ever, what it skipped and finds no longer armed;
// the snapshot for that merge is taken in the part's own prologue, before its GPU work. From that check to the part's
// hand-over (vk_mem_lazy_handed) memory.c's window keeps a fault from releasing a range unseen.
static VkCommandPool lz_pool;
static VkCommandBuffer lz_cb;
static VkFence lz_fence;
static volatile int32_t lz_state; // 0 not asked, 1 on, 2 off
static PlatMutex lz_init_lock = PLAT_MUTEX_INIT;
typedef struct { VkMemSpan sp; uint64_t size; int shader; } LzSkip;
static LzSkip *lz_skip; // replay: what the part being recorded left in the mirror
static int lz_nskip, lz_cskip;
static int lz_win; // replay: the part's window is open (lazy_part_check)
static uint64_t lz_skipped; // replay: bytes left in the mirror since the last log line
static void mbarrier(VkCommandBuffer cb, VkPipelineStageFlags2 ss, VkAccessFlags2 sa, VkPipelineStageFlags2 ds,
VkAccessFlags2 da);
// memory.c's fetch (frz_lock held there, so one at a time): [off, off + size) of the mirror h0 into the imported
// buffer h1, after all the GPU work handed over so far.
static void lazy_fetch(uint64_t h0, uint64_t h1, uint64_t off, uint64_t size) {
VkCtx *c = vk_ctx();
VkBufferCopy reg = {off, off, size};
VkCommandBufferBeginInfo bi = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO,
.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT};
vkResetCommandPool(c->device, lz_pool, 0);
vkBeginCommandBuffer(lz_cb, &bi);
mbarrier(lz_cb, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT, VK_ACCESS_2_MEMORY_WRITE_BIT, VK_PIPELINE_STAGE_2_COPY_BIT,
VK_ACCESS_2_TRANSFER_READ_BIT | VK_ACCESS_2_TRANSFER_WRITE_BIT);
vkCmdCopyBuffer(lz_cb, (VkBuffer)(uintptr_t)h0, (VkBuffer)(uintptr_t)h1, 1, &reg);
mbarrier(lz_cb, VK_PIPELINE_STAGE_2_COPY_BIT, VK_ACCESS_2_TRANSFER_WRITE_BIT,
VK_PIPELINE_STAGE_2_HOST_BIT | VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
VK_ACCESS_2_HOST_READ_BIT | VK_ACCESS_2_MEMORY_READ_BIT | VK_ACCESS_2_MEMORY_WRITE_BIT);
vkEndCommandBuffer(lz_cb);
VkSubmitInfo si = {.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO, .commandBufferCount = 1, .pCommandBuffers = &lz_cb};
vk_queue_lock();
vk_sync_flush_held(); // the held submissions wrote first
VkResult r = vkQueueSubmit(c->gfx_queue, 1, &si, lz_fence);
vk_queue_unlock();
if (r == VK_SUCCESS) r = vkWaitForFences(c->device, 1, &lz_fence, VK_TRUE, 10000000000ull);
if (r != VK_SUCCESS) {
vk_device_fault_report(r);
rt_fatal("vk: bringing %llu bytes the GPU wrote down to guest memory failed (%d)", (unsigned long long)size, (int)r);
}
vkResetFences(c->device, 1, &lz_fence);
}
// memory.c's test before it arms [lo, hi) of the pool ctx (frz_lock held there, the pages already inaccessible): the
// mirror holds every CPU write to them. A fetch copies the mirror over the whole range, not the merge of the bytes the
// GPU changed that a write-back does, so a CPU write that came after the range's write-back reached guest memory and
// before this arm, and that no sync of the mirror has carried up yet, would be overwritten by the mirror's older bytes:
// a page stamped after the mirror's last sync (the CPU-write table, mgen), or a write-watch bit nobody has read yet.
// No lock: the pages are inaccessible, so nothing stamps them meanwhile; the mirror's mgen only grows.
static int lazy_clean(uint64_t lo, uint64_t hi, void *ctx) {
const PoolRec *r = ctx;
uint32_t mgen = __atomic_load_n(&r->mgen, __ATOMIC_ACQUIRE);
for (uint64_t a = lo; a < hi; a += 1ull << GEN_PAGE_SHIFT) {
const uint32_t *ch = __atomic_load_n(&cgen_chunk[a >> GEN_CHUNK_SHIFT], __ATOMIC_ACQUIRE);
if (ch && __atomic_load_n(&ch[(a >> GEN_PAGE_SHIFT) & (GEN_PER_CHUNK - 1)], __ATOMIC_RELAXED) > mgen) return 0;
}
uint64_t p[1];
return mem_written_peek(lo, hi - lo, p, 1) == 0; // UINT64_MAX (not watched) counts as written
}
// Whether the lazy write-back is wanted on this device, decided and logged once, when the driver is known
// (vk_device.c create_device, after vk_caps_init): the player's setting, else the default of the platform and driver
// (wb_lazy_choice.h). Never in import mode, where no pool has a mirror; with shadow storage, where a fetch would copy
// the mirror into the pool's buffer, which is guest memory only when the pool is imported (a shadow's bytes reach
// guest memory in the submission's shadow_wb_cb); or with the write-back probe, which watches the same pages.
_Static_assert(WBL_DRIVER_AMD == VK_DRIVER_ID_AMD_PROPRIETARY, "wb_lazy_choice.h names AMD's own driver");
static int lz_wanted = -1; // -1 not decided, 0 off, 1 on
static const char *lz_why = ""; // the reason the log gave
void vk_mem_lazy_decide(void) {
const char *set = cfg_is_set("NX2NAVI_WB_LAZY") ? cfg_get("NX2NAVI_WB_LAZY") : NULL;
WblChoice ch = wb_lazy_choice(set, WBL_ON_WINDOWS, vk_caps.driver_id);
if (!vk_mirror_on) ch = (WblChoice){0, "no pool has a mirror in import mode"};
else if (shadow_on) ch = (WblChoice){0, "shadow storage keeps it off"};
else if (ch.on && cfg_bool("NX2NAVI_WB_PROBE")) ch = (WblChoice){0, "the write-back probe watches the same pages"};
lz_why = ch.why;
__atomic_store_n(&lz_wanted, ch.on, __ATOMIC_RELEASE);
rt_log("mem: lazy write-back %s: %s", ch.on ? "on" : "off", ch.why);
}
static int wb_lazy(void) {
int s = __atomic_load_n(&lz_state, __ATOMIC_ACQUIRE);
if (s) return s == 1;
plat_mutex_lock(&lz_init_lock);
if (!lz_state) {
VkCtx *c = vk_ctx();
int on = c && __atomic_load_n(&lz_wanted, __ATOMIC_ACQUIRE) == 1;
if (on) {
VkCommandPoolCreateInfo pi = {.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO,
.queueFamilyIndex = c->gfx_family};
VkFenceCreateInfo fi = {.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO};
on = vkCreateCommandPool(c->device, &pi, NULL, &lz_pool) == VK_SUCCESS;
VkCommandBufferAllocateInfo ai = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO, .commandPool = lz_pool,
.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY, .commandBufferCount = 1};
on = on && vkAllocateCommandBuffers(c->device, &ai, &lz_cb) == VK_SUCCESS &&
vkCreateFence(c->device, &fi, NULL, &lz_fence) == VK_SUCCESS;
on = on && mem_lazy_enable(lazy_fetch, lazy_clean);
}
rt_log("mem: what shaders write into mirrored pools %s (lazy write-back: %s)",
on ? "goes back to guest memory once and then only when the CPU touches it"
: "goes back to guest memory at the end of every submission",
on || lz_wanted != 1 ? lz_why : "its command buffer and fence could not be made");
__atomic_store_n(&lz_state, on ? 1 : 2, __ATOMIC_RELEASE);
}
plat_mutex_unlock(&lz_init_lock);
return lz_state == 1;
}
static void gpu_wrote_range(const VkMemSpan *sp, uint64_t size, int shader);
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 || !wb_lazy()) {
gpu_wrote_range(sp, size, shader);
return;
}
uint64_t g0 = r->import_lo + sp->offset, g1 = g0 + size, run = g0, out[16][2];
int k = mem_lazy_skip(g0, size, out, 16);
for (int i = 0; i <= k; i++) {
uint64_t lo = i < k ? out[i][0] : g1, hi = i < k ? out[i][1] : g1;
if (lo > run) { // not armed: written back as ever
VkMemSpan part = *sp;
part.offset = sp->offset + (run - g0);
gpu_wrote_range(&part, lo - run, shader);
}
if (i < k) {
if (lz_nskip == lz_cskip) grow((void **)&lz_skip, &lz_cskip, sizeof *lz_skip);
LzSkip k = {*sp, hi - lo, shader};
k.sp.offset = sp->offset + (lo - g0);
lz_skip[lz_nskip++] = k;
lz_skipped += hi - lo;
}
run = hi;
}
}
// The end of the part being recorded: what it left in the mirror for a range that is no longer armed (the CPU touched
// it meanwhile) is written back after all. Opens memory.c's window, which vk_mem_lazy_handed closes.
static void lazy_part_check(void) {
if (lz_win || !lz_nskip) return;
mem_lazy_window(1);
lz_win = 1;
for (int i = 0; i < lz_nskip; i++) {
const PoolRec *r = lz_skip[i].sp.pool;
if (!mem_lazy_armed(r->import_lo + lz_skip[i].sp.offset, lz_skip[i].size))
gpu_wrote_range(&lz_skip[i].sp, lz_skip[i].size, lz_skip[i].shader);
}
}
// vk_sync.c, before a part's prologue is recorded: the check for a part that was split off (the last part's ran in
// vk_mem_mirror_end already).
void vk_mem_lazy_part_end(void) { lazy_part_check(); }
// vk_sync.c, once the part is held or handed to the queue: a fetch from here on has its writes.
void vk_mem_lazy_handed(void) {
lz_nskip = 0;
if (lz_win) mem_lazy_window(0), lz_win = 0;
}
static void gpu_wrote_range(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;
int pl = sub_place(), pp = presnap_place();
if (pl <= pplace) pframe++; // a new frame, though its first submissions may have had nothing to record
pplace = pl;
for (int i = 0; i < nm; i++) {
const PoolRec *r = mlist[i];
if (pp ? !(pm_last(r) >> pl & 1) : 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_wb_pending(void) { return vk_mirror_on && (nsnl || nwbl); }
int vk_mem_mirror_end(VkCommandBuffer cb, uint64_t fence) {
if (!vk_mirror_on) return 0;
int rec = 0;
lazy_part_check(); // before the merges are recorded: what the last part skipped and has to write back after all
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);
vk_gpu_gap_note(VK_GAP_WB_MERGE, 0);
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 (shadow_on && (nswl || nwbl)) swb_job(fence); // before snap_free takes back the marks' chunk
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;
vk_gpu_gap_note(VK_GAP_WB_COPY, 0);
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);
mem_probe_report(frame);
mem_lazy_report(frame, lz_skipped);
lz_skipped = 0;
// NX2NAVI_WB_LAZY_TOUCH=<n> (test, 0 off): at each of these lines this thread reads a byte of n armed ranges, as
// the game's CPU might, so the fault, the fetch and the hold run without waiting for the game to do it.
static int touch = -1;
if (touch < 0) touch = cfg_int("NX2NAVI_WB_LAZY_TOUCH");
if (touch > 0) {
uint64_t at[16];
int k = mem_lazy_pick(at, touch > 16 ? 16 : touch);
for (int i = 0; i < k; i++) {
volatile uint8_t v = *(volatile const uint8_t *)(uintptr_t)at[i];
(void)v;
}
}
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;
}
// The overlay's "game" figure beside the card's memory in use (perfmon.c, its sampling thread): what the driver
// counts for this process on the device-local heaps, the same figure vram_use_log splits by use.
// With the device lost report (vk_device.c): the video memory in use and the budget the driver gives the game. A lost
// device with the memory at or over its budget points at the memory; one far below it does not, and the report of a
// card with 8 GB that ended this way held neither number.
void vk_mem_lost_report(void) {
uint64_t used, budget;
int exact;
if (!vram_query(&used, &budget, &exact)) return;
rt_log("vk: video memory when the device was lost: %llu MiB in use of a budget of %llu MiB (%s)",
(unsigned long long)(used >> 20), (unsigned long long)(budget >> 20),
exact ? "the driver's figures, every program counted in the budget" : "this game's own allocations only");
}
// The settings that change how the game's memory reaches the GPU, as they are in effect, for a fatal error's summary
// line (vk_device.c fatal_note): the memory mode and the lazy write-back with its reason and state. No lock, any thread.
void vk_mem_fatal_note(char *out, size_t n) {
char lazy[160];
mem_lazy_note(lazy, sizeof lazy);
int st = __atomic_load_n(&lz_state, __ATOMIC_ACQUIRE), want = __atomic_load_n(&lz_wanted, __ATOMIC_ACQUIRE);
snprintf(out, n, "memory mode %s%s, lazy write-back %s (%s%s%s)", mirror_verify ? "verify" : vk_mirror_on ? "mirror" : "import",
shadow_on ? " with shadow storage" : "", want < 0 ? "not decided" : st == 2 && want ? "failed to start" : want ? "on" : "off",
lz_why, *lazy ? "; " : "", lazy);
}
float vk_mem_device_used_mb(void) {
VkCtx *c = vk_ctx();
if (!c || !c->has_memory_budget) return -1;
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);
uint64_t usage = 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];
return (float)(usage / 1048576.0);
}
// One line per boot stage (vk_pipeline.c calls it after the pipeline cache load, the shader prewarm, the recorded
// pipelines and the compute pipeline prebuild): the device-local memory in use, the part of it VMA did not allocate
// (outside_watch's "outside VMA": pipelines and the driver's own, before the swapchain and the upscaler exist), and the
// process's commit and working set. Subtracting two lines gives what a stage left behind: on an installed copy the
// compute pipeline prebuild alone left 2.1 GiB of device memory on an RTX 5090, held by twenty very large compute
// shaders no run dispatches (vk_pipeline.c, "compute pipelines the boot keeps"), which this line was added to place.
void vk_mem_stage_log(const char *stage) {
VkCtx *c = vk_ctx();
double proc, used, limit;
float ram_used, ram_total, ws;
plat_commit_info(&proc, &used, &limit);
plat_mem_info(&ram_used, &ram_total, &ws);
const double M = 1048576.0;
uint64_t usage = 0, blocks = 0;
int budget = c && c->has_memory_budget;
if (budget) {
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] = {0};
if (c->vma) vmaGetHeapBudgets(c->vma, vb);
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;
}
char dev[96] = "device memory unknown (no VK_EXT_memory_budget)";
if (budget)
snprintf(dev, sizeof dev, "device memory in use %.0f MiB, outside VMA %.0f MiB", usage / M,
usage > blocks ? (usage - blocks) / M : 0.0);
rt_log("mem: boot stage, after %s: %s; process commit %.0f MB, working set %.0f MB; system commit %.0f of %.0f MB",
stage, dev, proc, (double)ws, used, limit);
}
// 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 destroy_dead(VkCtx *c, const Dead *d) {
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;
}
}
// The objects the GPU is done with: destroyed here (out != NULL: moved to *out, for the caller to destroy without
// cmd_lock; *nout of them).
static void collect_locked_out(VkCtx *c, Dead **out, int *nout, int *cout) {
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);
if (!out) { destroy_dead(c, d); continue; }
if (*nout == *cout) grow((void **)out, cout, sizeof **out);
(*out)[(*nout)++] = *d;
}
ndead = k;
}
static void collect_locked(VkCtx *c) { collect_locked_out(c, NULL, NULL, NULL); }
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_ERROR_OUT_OF_DEVICE_MEMORY && aci.pool) {
// The staging pool can grow only in 64 MiB blocks of its one memory type. A smaller allocation may
// still fit in a default block or another host-visible type. Keep the same mapping requirements.
aci.pool = NULL;
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) {
// a lost device ends the wait at once and leaves the fence unsignalled: retire would take nothing and this
// loop would spin for good, holding the upload lock, instead of the run ending with the device's report
VkResult w = vkWaitForFences(c->device, 1, &inflight[0].fence, VK_TRUE, UINT64_MAX);
if (w != VK_SUCCESS) {
vk_device_fault_report(w);
rt_fatal("vk: the wait for an upload failed (%d)", (int)w);
}
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);
}
// The objects go once cmd_lock is released: destroying an image and its memory takes up to milliseconds (the
// collect's 0.3 slow ones a frame took a median 0.26-0.43 ms, up to 1.9 ms), and the replay's uploads
// (vk_upload_begin) wait for cmd_lock.
void vk_mem_collect(void) {
VkCtx *c = vk_ctx();
if (!c) return;
plat_once(&cmd_once, cmd_init, NULL);
Dead *out = NULL;
int nout = 0, cout = 0;
plat_rmutex_lock(&cmd_lock);
collect_locked_out(c, &out, &nout, &cout);
plat_rmutex_unlock(&cmd_lock);
for (int i = 0; i < nout; i++) destroy_dead(c, &out[i]);
free(out);
}
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 = nswl = 0; // mirrors went with their pools
swl_marks = -1;
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 = 0, s_sb_pages = 0; // shadows too (their base copies went with their records)
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