Mirrors/totk-pc-recomp
1
0 Fork 0
mirror of https://github.com/valentderah/totk-pc-recomp.git synced 2026-11-12 10:21:02 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
4c1762104513f7b94de2f236083547bfd9ecfd12
totk-pc-recomp/runtime/vk_memory.c
T
Valiantsin DzerakhandClaude Opus 5.5 43704c0d9f Turn off the recorder thread on AMD GPUs and GPL on AMD's RDNA 1/2 driver, and log GPU breadcrumbs on a lost device 
AMD drivers lost the device at boot or at menu opens with the recorder thread; vk_recorder = 0 ended it for three
RX 9070 XT users (issue #5), and RADV saw an illegal opcode at the end of a command buffer (issue #40). It now
defaults to off on vendor 0x1002; TOTK_VK_RECORDER=1 turns it back on.

On AMD's separate driver branch for RDNA 1/2 ("AMD proprietary shader compiler") every device loss came with a cold
pipeline cache and draws waiting on fast-linked pipelines, and the recorder setting did not help there. TOTK_GPL now
has no fixed default: off on that branch, on elsewhere.

VK_AMD_buffer_marker breadcrumbs (vk_crumbs.c) mark the start and end of each submission part, upload batch, ASTC
decode chunk and present; on device lost the log names the last point the GPU reached and the last one it finished.

vk_fsr.c ensure_images held no queue lock around vkDeviceWaitIdle, which needs every queue externally synchronized.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-30 19:48:38 +03:00

1840 lines
91 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.
//
// Behaviour ported from MacNativeGames games/totk/third_party/d2r_gfx/gfx/nvn/totk_vpool.c (request checks, unmap
// semantics, 64 KiB pages) and nvn_dirty.c (generation stamps: a texture is stale when a covered page carries a
// generation newer than its upload), @ b75be8cf (MIT, Copyright (c) 2026 jhaizhou; origin d2r-mobile); and from LAS
// Recomp runtime/nvn/vulkan/exec.c:2094-2131 (vk_pages_changed: every write-watch query resets the OS bits, so each
// query stamps the pages it saw with a new generation shared by all users) and textures.c:240-270 (deferred
// destruction) @ da35b83a (MIT, Copyright (c) 2026 Peiikai). MNG aliases host VM (mach_vm_remap) and protects pages
// (fault handler); here virtual pools are a GPU-side page table over device-local memory and dirty pages come from
// GetWriteWatch (mem_written_pages), so neither is needed.
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <windows.h>
#include "cfg.h"
#include "gfx.h"
#include "host.h"
#include "nvn_api.h"
#include "nvn_pool_index.h"
#include "vk.h"
// Raw performance counter ticks for the tracking cost (no division on the replay's path); qpc_ms converts.
static uint64_t qpc(void) {
LARGE_INTEGER t;
QueryPerformanceCounter(&t);
return (uint64_t)t.QuadPart;
}
static double qpc_ms(uint64_t ticks) {
static double per_ms;
if (!per_ms) {
LARGE_INTEGER f;
QueryPerformanceFrequency(&f);
per_ms = (double)f.QuadPart / 1000.0;
}
return (double)ticks / per_ms;
}
// ---------------------------------------------------------------- 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
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 long used; // resolved at least once: the GPU may still read it at finalize
// TOTK_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
VkBuffer mbuf; // MS_ON: device-local copy of [import_lo, import_hi), same offsets as buf
VmaAllocation malloc_;
uint32_t acc_sub; // replay: the last submission that synced the mirror (mirror_sync)
uint32_t wb_sub; // replay: the last submission that wrote mirror bytes back to guest memory
uint32_t mgen; // gen_lock: every page write stamped up to this generation is in the mirror
// (or in a copy queued before its next use)
uint64_t fl_lo, fl_hi; // flush_lock: nvnMemoryPoolFlushMappedRange bytes not copied yet (fl_hi 0: none)
} PoolRec;
static void mirror_forget(PoolRec *r); // pool_lock exclusive: off the mirrored list
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 SRWLOCK pool_lock = SRWLOCK_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 long 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;
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 (!p_host_props)
p_host_props = (PFN_vkGetMemoryHostPointerPropertiesEXT)vkGetDeviceProcAddr(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};
VkMemoryAllocateInfo ai = {.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO, .pNext = &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;
}
InterlockedIncrement(&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("TOTK_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->buf) vkDestroyBuffer(c->device, r->buf, NULL);
if (r->mem) vkFreeMemory(c->device, r->mem, NULL);
if (r->alloc) vmaFreeMemory(c->vma, r->alloc);
free(r->pt);
free(r);
}
int vk_pool_init(pool_t *p) {
VkCtx *c = vk_ctx();
if (!c) { rt_log("vk: MemoryPoolInitialize before vk_init"); return 0; }
// The GPU never sees a GPU_NO_ACCESS pool, except a physical one (0x189; 1.4.3 also 0x18c, CPU-visible) through the
// virtual pools that map it. Either gets device-local memory: 1.4.3's CPU never writes its 0x18c pool's storage
// (unchanged over 3600 frames of the intro), so the guest storage and the GPU copy need no sync.
if ((p->flags & NVN_MEMORY_POOL_FLAGS_GPU_NO_ACCESS) && !(p->flags & NVN_MEMORY_POOL_FLAGS_PHYSICAL)) return 1;
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 = InterlockedIncrement(&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; }
AcquireSRWLockExclusive(&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;
}
ReleaseSRWLockExclusive(&pool_lock);
return 1;
}
void vk_pool_fini(pool_t *p) {
VkCtx *c = vk_ctx();
if (!c) return;
pidx_ent e;
AcquireSRWLockExclusive(&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;
}
}
ReleaseSRWLockExclusive(&pool_lock);
if (!r) return;
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 destroy_rec(c, r);
}
// Requests were checked by the front end (nvn_api_core.c, MNG totk_vpool.c); apply them in order.
int vk_pool_map_virtual(pool_t *v, int n, const void *requests) {
const nvn_mapping_t *rq = requests;
AcquireSRWLockExclusive(&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};
}
ReleaseSRWLockExclusive(&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. TOTK_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("TOTK_RESOLVE_CHECK");
on = e && *e && strcmp(e, "0");
}
return on;
}
static int resolve_locked(uint64_t addr, uint64_t size, VkMemSpan *out);
int vk_mem_resolve(uint64_t addr, uint64_t size, VkMemSpan *out) {
memset(out, 0, sizeof *out);
uint64_t g = __atomic_load_n(&pool_gen, __ATOMIC_ACQUIRE), end = addr + size;
const ResolveHit *h = rhit_slot(addr);
if (end >= addr && h->gen == g && addr >= h->lo && end <= h->hi) {
out->buffer = h->buf;
out->offset = addr - h->base;
out->avail = h->end - addr;
out->alloc = h->alloc;
out->cls = h->cls;
out->pool = h->pool;
if (h->alloc) out->alloc_offset = out->offset;
if (resolve_check_on()) {
VkMemSpan s;
if (!resolve_locked(addr, size, &s) || memcmp(&s, out, sizeof s))
rt_fatal("vk: resolve cache of 0x%llx differs from the pool index", (unsigned long long)addr);
}
return 1;
}
return resolve_locked(addr, size, out);
}
// The pool index lookup under the lock; a non-overlapped import or physical pool goes into the address's cache slot.
static int resolve_locked(uint64_t addr, uint64_t size, VkMemSpan *out) {
memset(out, 0, sizeof *out);
AcquireSRWLockShared(&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;
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) InterlockedExchange(&ph->used, 1); // read first: the flag is only ever set
}
}
if (ok && !r->used) InterlockedExchange(&r->used, 1);
ReleaseSRWLockShared(&pool_lock);
if (!ok) memset(out, 0, sizeof *out);
return ok;
}
uint64_t vk_mem_gen(void) {
return __atomic_load_n(&pool_gen, __ATOMIC_ACQUIRE) + __atomic_load_n(&mirror_gen, __ATOMIC_ACQUIRE);
}
int vk_mem_class_of(uint64_t addr) {
VkMemSpan m;
return vk_mem_resolve(addr, 1, &m) ? m.cls : VK_MCLS_NONE;
}
// ---------------------------------------------------------------- census
static volatile LONG64 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) + (LONG64)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 TOTK_TRACE=perf every 1800 frames.
static void dirty_frame(uint64_t frame);
static void mirror_log(uint64_t frame);
static int dead_pending(void);
static void vram_log(uint64_t frame);
void vk_mem_census(uint64_t frame) {
static uint64_t window_start;
if ((rt_trace & RT_TRACE_PERF) && frame % 600 == 0) vram_log(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); // and 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++) InterlockedExchange64(&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};
AcquireSRWLockShared(&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]++;
}
ReleaseSRWLockShared(&pool_lock);
rt_log("mem: census at frame %llu: import %llu pools %.1f MiB (%ld 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] = InterlockedExchange64(&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;
}
int vk_mem_vram(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;
}
// ---------------------------------------------------------------- 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 SRWLOCK gen_lock = SRWLOCK_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];
// 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 *gen_slot(uint64_t page_addr, int create) {
uint64_t ci = page_addr >> GEN_CHUNK_SHIFT;
if (ci >= GEN_CHUNKS) return NULL;
if (!gen_chunk[ci]) {
if (!create) return NULL;
gen_chunk[ci] = calloc(GEN_PER_CHUNK, sizeof(uint32_t));
if (!gen_chunk[ci]) rt_fatal("vk: out of memory for page generations");
}
return &gen_chunk[ci][(page_addr >> GEN_PAGE_SHIFT) & (GEN_PER_CHUNK - 1)];
}
// gen_lock held exclusively; g is the newest generation handed out.
static void stamp_page(uint64_t a, uint32_t g) {
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;
}
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;
AcquireSRWLockExclusive(&gen_lock);
stamp_range(lo, hi, ++gen_now);
ReleaseSRWLockExclusive(&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;
AcquireSRWLockExclusive(&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;
}
ReleaseSRWLockExclusive(&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;
AcquireSRWLockShared(&gen_lock);
for (uint64_t r = r0; r <= r1; r++)
if (snap_at[r] != epoch) { excl = 1; break; }
if (excl) {
ReleaseSRWLockShared(&gen_lock);
AcquireSRWLockExclusive(&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) ReleaseSRWLockExclusive(&gen_lock);
else ReleaseSRWLockShared(&gen_lock);
return untracked ? -1 : changed;
}
void vk_mem_snap_stats(uint64_t *checks, uint64_t *calls, uint64_t *untracked) {
*checks = __atomic_exchange_n(&n_snap_checks, 0, __ATOMIC_RELAXED);
*calls = __atomic_exchange_n(&n_snap_calls, 0, __ATOMIC_RELAXED);
*untracked = __atomic_exchange_n(&n_snap_untracked, 0, __ATOMIC_RELAXED);
}
// ---------------------------------------------------------------- dirty census (TOTK_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 SRWLOCK dirty_lock = SRWLOCK_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;
}
}
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;
AcquireSRWLockShared(&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;
}
AcquireSRWLockExclusive(&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;
}
ReleaseSRWLockExclusive(&gen_lock);
ReleaseSRWLockShared(&pool_lock);
uint64_t tot = 0;
for (int k = 0; k < DC_N; k++) tot += eager[k];
AcquireSRWLockExclusive(&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;
ReleaseSRWLockExclusive(&dirty_lock);
}
static void dirty_frame(uint64_t frame) {
AcquireSRWLockExclusive(&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) { ReleaseSRWLockExclusive(&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);
ReleaseSRWLockExclusive(&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, TOTK_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 stamped after the mirror's mgen goes into a copy; so do the
// nvnMemoryPoolFlushMappedRange ranges. 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).
// - TOTK_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 16
enum { MR_BUDGET, MR_UNTRACKED, MR_OVERLAP, MR_ALLOC, MR_N };
int vk_mirror_on;
static int mirror_verify;
static int mirror_gap = 2; // pre-snap: hot pools at most this many 2 MiB regions apart share one query
static PoolRec **mlist; // pool_lock: mirrored pools by import_lo
static int nm, mcap;
static uint32_t msub; // replay: number of the submission being recorded
static uint32_t mepoch; // replay: its write-watch snapshot epoch (vk_texture_epoch)
typedef struct { PoolRec *r; uint64_t off, size; } MRange;
static MRange *pend, *wbl; // replay: copies for the next prologue; GPU writes into mirrors this submission
static int npend, cpend, nwbl, cwbl;
static VkBufferCopy *mregions; // replay: vkCmdCopyBuffer regions
static int cregions;
static SRWLOCK flush_lock = SRWLOCK_INIT;
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;
} MirrorAcc;
static MirrorAcc mrep; // replay: this submission
static MirrorAcc macc; // mirror_stat_lock: since the last summary
static SRWLOCK mirror_stat_lock = SRWLOCK_INIT;
static volatile LONG64 m_live_n, m_live_bytes, m_refused_n[MR_N], m_refused_bytes[MR_N];
static volatile LONG64 m_flush_calls, m_flush_bytes, v_pages, v_skipped, v_mismatch;
static void mirror_mode(int discrete) {
const char *e = cfg_value("TOTK_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 = "TOTK_MEMORY=import";
else if (!strcmp(e, "mirror")) mode = 1, why = "TOTK_MEMORY=mirror";
else if (!strcmp(e, "verify")) mode = 2, why = "TOTK_MEMORY=verify";
else rt_log("mem: TOTK_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("TOTK_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.
const char *fi = cfg_get("TOTK_FREEZE_IDLE");
freeze_idle = fi && *fi ? atoi(fi) : 60;
if (freeze_idle < 0) freeze_idle = 0;
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");
}
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;
uint64_t n = r->import_hi - r->import_lo;
if (InterlockedIncrement64(&m_refused_n[why]) <= 4) {
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);
}
InterlockedAdd64(&m_refused_bytes[why], (LONG64)n);
}
// 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;
AcquireSRWLockShared(&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);
ReleaseSRWLockShared(&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;
if (vk_mem_vram(&used, &budget, &exact) && used + n + budget / 8 > budget) {
char why[128];
snprintf(why, sizeof why, "its %.1f MiB mirror does not fit the VRAM budget (%.0f of %.0f MiB in use)",
n / 1048576.0, used / 1048576.0, budget / 1048576.0);
return refuse(r, MR_BUDGET, why);
}
AcquireSRWLockExclusive(&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
ReleaseSRWLockExclusive(&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);
}
AcquireSRWLockExclusive(&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++;
ReleaseSRWLockExclusive(&pool_lock);
r->acc_sub = msub;
r->mstate = MS_ON;
range_add(&pend, &npend, &cpend, r, 0, n);
mrep.first_n++, mrep.first_bytes += n;
InterlockedIncrement64(&m_live_n);
InterlockedAdd64(&m_live_bytes, (LONG64)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--;
InterlockedDecrement64(&m_live_n);
InterlockedAdd64(&m_live_bytes, -(LONG64)(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 SRWLOCK retire_lock = SRWLOCK_INIT;
static void mirror_retiring(PoolRec *r) {
AcquireSRWLockExclusive(&retire_lock);
if (nretiring == cretiring) grow((void **)&retiring, &cretiring, sizeof *retiring);
retiring[nretiring++] = r;
ReleaseSRWLockExclusive(&retire_lock);
}
static void mirror_retire(void *arg) {
AcquireSRWLockExclusive(&retire_lock);
for (int i = 0; i < nretiring; i++)
if (retiring[i] == arg) { retiring[i] = retiring[--nretiring]; break; }
ReleaseSRWLockExclusive(&retire_lock);
destroy_rec(vk_ctx(), arg);
}
// Replay thread, the first access of this submission: queues the copies of the pages written since the last sync.
static void mirror_sync(PoolRec *r) {
uint64_t t0 = qpc();
r->acc_sub = msub;
uint64_t lo = r->import_lo, hi = r->import_hi;
uint64_t r0 = lo >> SNAP_SHIFT, r1 = (hi - 1) >> SNAP_SHIFT;
int excl = 0;
AcquireSRWLockShared(&gen_lock);
for (uint64_t k = r0; k <= r1; k++)
if (snap_at[k] != mepoch) { excl = 1; break; }
if (excl) {
ReleaseSRWLockShared(&gen_lock);
AcquireSRWLockExclusive(&gen_lock);
mrep.sync_queries += snap_regions(r0, r1, mepoch);
}
uint32_t since = r->mgen;
for (uint64_t k = r0; k <= r1; k++) {
if (reg_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 = gen_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);
}
r->mgen = gen_now;
if (excl) ReleaseSRWLockExclusive(&gen_lock);
else ReleaseSRWLockShared(&gen_lock);
if (__atomic_load_n(&r->fl_hi, __ATOMIC_ACQUIRE)) {
AcquireSRWLockExclusive(&flush_lock);
uint64_t a = r->fl_lo, b = r->fl_hi;
r->fl_hi = 0;
ReleaseSRWLockExclusive(&flush_lock);
a = a > lo ? a & ~0xfffull : lo;
b = b < hi ? (b + 0xfff) & ~0xfffull : hi;
if (a < b) range_add(&pend, &npend, &cpend, r, a - lo, b - a);
}
mrep.syncs++;
mrep.sync_t += qpc() - t0;
}
void vk_mem_use(VkMemSpan *sp) {
PoolRec *r = sp->pool;
if (r->mstate == MS_NONE) mirror_create(r);
if (r->mstate != MS_ON) return;
if (r->acc_sub != msub) mirror_sync(r);
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;
}
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;
}
void vk_mem_gpu_wrote(const VkMemSpan *sp, uint64_t size) {
PoolRec *r = sp->pool;
if (!r || r->mstate != MS_ON || sp->buffer != r->mbuf || !size) return;
range_add(&wbl, &nwbl, &cwbl, r, sp->offset, size);
}
void nvn_pool_flushed(const pool_t *p, uint64_t off, uint64_t size) {
if (!vk_mirror_on || !size || !p->mem) return;
InterlockedIncrement64(&m_flush_calls);
InterlockedAdd64(&m_flush_bytes, (LONG64)size);
AcquireSRWLockShared(&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, b = a + size;
AcquireSRWLockExclusive(&flush_lock);
if (!r->fl_hi) r->fl_lo = a, r->fl_hi = b;
else {
if (a < r->fl_lo) r->fl_lo = a;
if (b > r->fl_hi) r->fl_hi = b;
}
ReleaseSRWLockExclusive(&flush_lock);
}
ReleaseSRWLockShared(&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;
AcquireSRWLockExclusive(&gen_lock);
ww_ep++;
freeze_pass();
ReleaseSRWLockExclusive(&gen_lock);
return;
}
uint64_t t0 = qpc();
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;
AcquireSRWLockShared(&pool_lock);
AcquireSRWLockExclusive(&gen_lock);
ww_ep++;
uint64_t b0 = ww_bytes;
for (int i = 0; i < nm; i++) {
const PoolRec *r = mlist[i];
if (r->acc_sub != prev) continue;
uint64_t a = r->import_lo >> SNAP_SHIFT, b = (r->import_hi - 1) >> SNAP_SHIFT;
if (have && a <= run1 + 1 + (uint64_t)mirror_gap) {
if (b > run1) run1 = b;
continue;
}
if (have) q += snap_regions(run0, run1, mepoch);
run0 = a, run1 = b, have = 1;
}
if (have) q += snap_regions(run0, run1, mepoch);
mrep.presnap_bytes += ww_bytes - b0;
if (freeze_idle) freeze_pass();
ReleaseSRWLockExclusive(&gen_lock);
ReleaseSRWLockShared(&pool_lock);
mrep.subs++;
mrep.presnap_queries += q;
mrep.presnap_t += qpc() - t0;
}
int vk_mem_mirror_pending(void) { return npend; }
void vk_mem_mirror_prologue(VkCommandBuffer cb) {
if (!npend) 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);
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 LONG 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;
if (page_dirty(a, s->mgen)) { skipped++; continue; }
bad++;
if (InterlockedIncrement(&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);
}
}
}
InterlockedAdd64(&v_pages, (LONG64)pages);
InterlockedAdd64(&v_skipped, (LONG64)skipped);
InterlockedAdd64(&v_mismatch, (LONG64)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;
AcquireSRWLockShared(&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;
}
ReleaseSRWLockShared(&pool_lock);
if (!j->n) { free(j); return 0; }
VkBufferCreateInfo bi = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, .size = dst,
.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT};
VmaAllocationCreateInfo ai = {.usage = VMA_MEMORY_USAGE_AUTO_PREFER_HOST,
.flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_RANDOM_BIT | VMA_ALLOCATION_CREATE_MAPPED_BIT};
VmaAllocationInfo info;
if (vmaCreateBuffer(vk_ctx()->vma, &bi, &ai, &j->buf, &j->alloc, &info) != VK_SUCCESS) { free(j); return 0; }
j->map = info.pMappedData;
j->sub = msub;
mbarrier(cb, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT, VK_ACCESS_2_MEMORY_WRITE_BIT, VK_PIPELINE_STAGE_2_COPY_BIT,
VK_ACCESS_2_TRANSFER_READ_BIT);
for (int i = 0; i < j->n; i++) {
VkBufferCopy c = {j->s[i].off, j->s[i].dst, j->s[i].len};
vkCmdCopyBuffer(cb, j->s[i].r->mbuf, j->buf, 1, &c);
}
mbarrier(cb, VK_PIPELINE_STAGE_2_COPY_BIT, VK_ACCESS_2_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_2_HOST_BIT,
VK_ACCESS_2_HOST_READ_BIT);
vk_completion_add(fence, verify_cb, j);
return j->n + 2;
}
int vk_mem_mirror_end(VkCommandBuffer cb, uint64_t fence) {
if (!vk_mirror_on) return 0;
int rec = 0;
if (nwbl) {
mbarrier(cb, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT, VK_ACCESS_2_MEMORY_WRITE_BIT, VK_PIPELINE_STAGE_2_COPY_BIT,
VK_ACCESS_2_TRANSFER_READ_BIT);
for (int i = 0; i < nwbl; i++) wbl[i].r->wb_sub = msub;
mrep.wb_ranges += (uint64_t)nwbl;
rec += copy_ranges(cb, wbl, nwbl, 0, &mrep.wb_bytes, &mrep.wb_copies) + 2;
mbarrier(cb, VK_PIPELINE_STAGE_2_COPY_BIT, VK_ACCESS_2_TRANSFER_WRITE_BIT,
VK_PIPELINE_STAGE_2_HOST_BIT | VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
VK_ACCESS_2_HOST_READ_BIT | VK_ACCESS_2_MEMORY_READ_BIT | VK_ACCESS_2_MEMORY_WRITE_BIT);
nwbl = 0;
}
if (mirror_verify) rec += verify_record(cb, fence);
AcquireSRWLockExclusive(&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];
ReleaseSRWLockExclusive(&mirror_stat_lock);
memset(&mrep, 0, sizeof mrep);
return rec;
}
// One summary line: with TOTK_TRACE=perf every 60 frames, else at frames 600 and 1800.
static void mirror_log(uint64_t frame) {
static uint64_t last, ww_last;
AcquireSRWLockExclusive(&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;
ReleaseSRWLockExclusive(&gen_lock);
double wf = frame > ww_last ? (double)(frame - ww_last) : 1.0;
ww_last = frame;
uint64_t fz, ff, ft, fo;
mem_freeze_stats(&fz, &ff, &ft, &fo);
if (rt_trace & RT_TRACE_PERF)
rt_log("mem: write-watch frame %llu: per frame, every consumer: %.1f queries over %.1f MiB, %.3f ms (%.1f found "
"%.0f written pages, %.3f ms); frozen %.0f MiB; since the last line: freezes %llu, thaws by a write %llu, by a file read "
"or mapping change %llu, refused %llu",
(unsigned long long)frame, wc / wf, wb / wf / 1048576.0, qpc_ms(wt) / wf, dc / wf, wp / wf, qpc_ms(dt) / wf,
fz * 2.0, (unsigned long long)ff, (unsigned long long)ft, (unsigned long long)fo, (unsigned long long)nr);
if (!vk_mirror_on) return;
AcquireSRWLockExclusive(&mirror_stat_lock);
MirrorAcc a = macc;
memset(&macc, 0, sizeof macc);
ReleaseSRWLockExclusive(&mirror_stat_lock);
double f = frame > last ? (double)(frame - last) : 1.0, M = 1048576.0;
last = frame;
uint64_t fc = (uint64_t)InterlockedExchange64(&m_flush_calls, 0), fb = (uint64_t)InterlockedExchange64(&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)InterlockedExchange64(&v_pages, 0), (long long)InterlockedExchange64(&v_skipped, 0),
(long long)InterlockedExchange64(&v_mismatch, 0));
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, flush calls %.1f (%.3f MiB), queries "
"%.1f pre-snap (%.1f MiB) + %.1f sync, prologues %.2f, tracking cpu %.3f ms (pre-snap %.3f, syncs %.3f)%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, 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, ver);
}
// With TOTK_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());
}
// 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; } Staging;
typedef struct {
VkCommandBuffer cb;
VkFence fence;
uint64_t seq;
Staging *stg;
int nstg, cap;
} Batch;
enum { DEAD_IMAGE = 1, DEAD_IMAGE_VMA, DEAD_VIEW, DEAD_BUFFER_VMA };
typedef struct { int kind; uint64_t handle; VmaAllocation alloc; uint64_t seq, tl; } Dead;
static CRITICAL_SECTION cmd_lock;
static INIT_ONCE cmd_once = INIT_ONCE_STATIC_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 BOOL CALLBACK cmd_init(PINIT_ONCE o, PVOID p, PVOID *ctx) {
(void)o, (void)p, (void)ctx;
InitializeCriticalSection(&cmd_lock);
return TRUE;
}
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;
}
static void free_staging(VkCtx *c, Batch *b) {
for (int i = 0; i < b->nstg; i++) vmaDestroyBuffer(c->vma, b->stg[i].buf, b->stg[i].alloc);
b->nstg = 0;
}
// Batches complete in submission order (one queue): retire from the front while their fences are signalled.
static void retire(VkCtx *c, int wait_all) {
int k = 0;
for (; k < ninflight; k++) {
Batch *b = &inflight[k];
if (wait_all) vkWaitForFences(c->device, 1, &b->fence, VK_TRUE, UINT64_MAX);
else if (vkGetFenceStatus(c->device, b->fence) != VK_SUCCESS) break;
free_staging(c, b);
done_seq = b->seq;
if (nspare == cap_spare) grow((void **)&spare, &cap_spare, sizeof *spare);
spare[nspare++] = *b;
}
memmove(inflight, inflight + k, (size_t)(ninflight - k) * sizeof *inflight);
ninflight -= k;
}
static void collect_locked(VkCtx *c) {
retire(c, 0);
uint64_t tl = 0;
if (c->timeline) vkGetSemaphoreCounterValue(c->device, c->timeline, &tl);
int k = 0;
for (int i = 0; i < ndead; i++) {
Dead *d = &dead[i];
if (d->seq > done_seq || d->tl > tl) { dead[k++] = *d; continue; }
switch (d->kind) {
case DEAD_IMAGE: vkDestroyImage(c->device, (VkImage)(uintptr_t)d->handle, NULL); break;
case DEAD_IMAGE_VMA: vmaDestroyImage(c->vma, (VkImage)(uintptr_t)d->handle, d->alloc); break;
case DEAD_VIEW: vkDestroyImageView(c->device, (VkImageView)(uintptr_t)d->handle, NULL); break;
case DEAD_BUFFER_VMA: vmaDestroyBuffer(c->vma, (VkBuffer)(uintptr_t)d->handle, d->alloc); break;
}
}
ndead = k;
}
VkCommandBuffer vk_upload_begin(void) {
InitOnceExecuteOnce(&cmd_once, cmd_init, NULL, NULL);
EnterCriticalSection(&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;
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) { LeaveCriticalSection(&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)) {
InitOnceExecuteOnce(&cmd_once, cmd_init, NULL, NULL);
EnterCriticalSection(&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);
}
LeaveCriticalSection(&cmd_lock);
}
if (stg_pool && bi.size <= STG_BLOCK / 2) aci.pool = stg_pool;
VmaAllocationInfo info;
VkBuffer b;
VmaAllocation a;
VkResult r = vmaCreateBuffer(c->vma, &bi, &aci, &b, &a, &info);
if (r != VK_SUCCESS) rt_fatal("vk: staging buffer of %llu bytes failed (%d)", (unsigned long long)size, (int)r);
*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};
}
return info.pMappedData;
}
uint64_t vk_upload_flush(int wait) {
InitOnceExecuteOnce(&cmd_once, cmd_init, NULL, NULL);
EnterCriticalSection(&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);
vk_crumb(open_b.cb, VK_CRUMB_UPLOAD_END, open_b.seq);
vkEndCommandBuffer(open_b.cb);
VkSubmitInfo si = {.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO, .commandBufferCount = 1, .pCommandBuffers = &open_b.cb};
vk_queue_lock();
VkResult r = vkQueueSubmit(c->gfx_queue, 1, &si, open_b.fence);
vk_queue_unlock();
if (r != VK_SUCCESS) {
vk_device_fault_report(r);
rt_fatal("vk: upload submit failed (%d)", (int)r);
}
if (ninflight == cap_inflight) grow((void **)&inflight, &cap_inflight, sizeof *inflight);
inflight[ninflight++] = open_b;
open_active = 0;
}
if (c && wait) {
while (ninflight && done_seq < seq) {
vkWaitForFences(c->device, 1, &inflight[0].fence, VK_TRUE, UINT64_MAX);
retire(c, 0);
}
}
if (c) collect_locked(c);
LeaveCriticalSection(&cmd_lock);
return seq;
}
static void defer(int kind, uint64_t handle, VmaAllocation alloc) {
InitOnceExecuteOnce(&cmd_once, cmd_init, NULL, NULL);
EnterCriticalSection(&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.
dead[ndead++] = (Dead){kind, handle, alloc, open_active ? open_b.seq : next_seq, tl_noted};
LeaveCriticalSection(&cmd_lock);
}
static int dead_pending(void) {
InitOnceExecuteOnce(&cmd_once, cmd_init, NULL, NULL);
EnterCriticalSection(&cmd_lock);
int n = ndead;
LeaveCriticalSection(&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_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) {
InitOnceExecuteOnce(&cmd_once, cmd_init, NULL, NULL);
EnterCriticalSection(&cmd_lock);
if (timeline_value > tl_noted) tl_noted = timeline_value;
LeaveCriticalSection(&cmd_lock);
}
void vk_mem_collect(void) {
VkCtx *c = vk_ctx();
if (!c) return;
InitOnceExecuteOnce(&cmd_once, cmd_init, NULL, NULL);
EnterCriticalSection(&cmd_lock);
collect_locked(c);
LeaveCriticalSection(&cmd_lock);
}
void vk_mem_stats(VkMemStats *s) {
memset(s, 0, sizeof *s);
AcquireSRWLockShared(&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++;
}
ReleaseSRWLockShared(&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_texture_shutdown();
EnterCriticalSection(&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);
}
nspare = 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;
LeaveCriticalSection(&cmd_lock);
AcquireSRWLockExclusive(&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 = 0; // mirrors went with their pools
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;
ReleaseSRWLockExclusive(&pool_lock);
AcquireSRWLockExclusive(&gen_lock);
for (uint64_t i = 0; i < GEN_CHUNKS; i++) free(gen_chunk[i]), gen_chunk[i] = NULL;
ReleaseSRWLockExclusive(&gen_lock);
}
Reference in new issue View Git Blame Copy Permalink