Advance M8-M11 parity workflows

web/app/src/render/nanovdb-volume-renderer.ts

@@ -1,4 +1,18 @@
 import type { NanoVDBGridIR, NanoVDBMaterialIR } from "../../../protocol/volume-vdb";
+import {
+  consumeNanoVDBProgressiveRedrawBudget,
+  NANOVDB_PROGRESSIVE_REDRAW_LIMIT_CODE,
+  NANOVDB_PROGRESSIVE_REDRAW_MAX_FRAMES,
+  validateNanoVDBProgressiveRedrawLimit,
+} from "../../../protocol/nanovdb-progressive-redraw";
+import {
+  createNanoVDBPageFeedbackBuffer,
+  nanoVDBPageFeedbackByteLength,
+  NANOVDB_PAGE_FEEDBACK_RECORD_WGSL,
+  NANOVDB_PAGE_FEEDBACK_WGSL,
+  parseNanoVDBPageFeedbackBatch,
+  type NanoVDBPageFeedbackBatch,
+} from "../../../protocol/nanovdb-page-feedback";
 
 export interface NanoVDBWebGPUCapabilityIR {
   available: boolean;
@@ -23,6 +37,9 @@ export interface NanoVDBWebGPUGrid {
   residentVirtualPages: readonly number[];
   uploadPage(pageIndex: number, data?: ArrayBuffer): void;
   touchPage(pageIndex: number): boolean;
+  beginFrame(): void;
+  pinPage(pageIndex: number): boolean;
+  endFrame(): void;
   evictPage(pageIndex: number): void;
   hasResidentPage(pageIndex: number): boolean;
   dispose(): void;
@@ -44,6 +61,104 @@ export interface NanoVDBGpuPageAllocatorStatsIR {
   keys: string[];
 }
 
+export interface NanoVDBProgressiveRedrawStatsIR {
+  pending: boolean;
+  scheduledCount: number;
+  redrawCount: number;
+  maxRedraws: number;
+  capped: boolean;
+  errorCode: "NANOVDB_PROGRESSIVE_REDRAW_LIMIT" | null;
+}
+
+export interface NanoVDBProgressiveRedrawOptionsIR {
+  maxRedraws?: number;
+}
+
+export class NanoVDBProgressiveRedrawScheduler {
+  private pending = false;
+  private scheduledCount = 0;
+  private redrawCount = 0;
+  private capped = false;
+  private disposed = false;
+  private generation = 0;
+  private readonly maxRedraws: number;
+
+  constructor(
+    private readonly requestFrame: (callback: () => void) => void,
+    private readonly redraw: () => void,
+    options: NanoVDBProgressiveRedrawOptionsIR = {},
+  ) {
+    const maxRedraws = options.maxRedraws ?? NANOVDB_PROGRESSIVE_REDRAW_MAX_FRAMES;
+    this.maxRedraws = validateNanoVDBProgressiveRedrawLimit(maxRedraws);
+  }
+
+  schedule(): boolean {
+    if (this.disposed || this.pending || this.capped) return false;
+    const budget = consumeNanoVDBProgressiveRedrawBudget(this.redrawCount, this.maxRedraws);
+    if (!budget.allowed) {
+      this.capped = budget.capped;
+      return false;
+    }
+    this.pending = true;
+    const generation = this.generation;
+    try {
+      this.requestFrame(() => {
+        if (this.disposed || generation !== this.generation || !this.pending) return;
+        this.pending = false;
+        const consumed = consumeNanoVDBProgressiveRedrawBudget(this.redrawCount, this.maxRedraws);
+        this.redrawCount = consumed.redrawCount;
+        this.capped = consumed.capped;
+        this.redraw();
+      });
+      this.scheduledCount++;
+      return true;
+    }
+    catch (error) {
+      this.pending = false;
+      throw error;
+    }
+  }
+
+  stats(): NanoVDBProgressiveRedrawStatsIR {
+    return {
+      pending: this.pending,
+      scheduledCount: this.scheduledCount,
+      redrawCount: this.redrawCount,
+      maxRedraws: this.maxRedraws,
+      capped: this.capped,
+      errorCode: this.capped ? NANOVDB_PROGRESSIVE_REDRAW_LIMIT_CODE : null,
+    };
+  }
+
+  /** Starts a new render epoch and invalidates callbacks queued by the old epoch. */
+  beginRender(): void {
+    if (this.disposed) return;
+    this.generation++;
+    this.pending = false;
+    this.scheduledCount = 0;
+    this.redrawCount = 0;
+    this.capped = false;
+  }
+
+  dispose(): void {
+    this.disposed = true;
+    this.generation++;
+    this.pending = false;
+  }
+}
+
+export class NanoVDBProgressivePageUploader {
+  constructor(
+    private readonly grid: NanoVDBWebGPUGrid,
+    private readonly redraw: NanoVDBProgressiveRedrawScheduler,
+  ) {}
+
+  upload(pageId: number, data: ArrayBuffer): { pageId: number; redrawScheduled: boolean } {
+    this.grid.uploadPage(pageId, data);
+    return { pageId, redrawScheduled: this.redraw.schedule() };
+  }
+}
+
 export interface NanoVDBDeviceLossIR { reason?: string; message: string }
 
 const traversalWGSL = /* wgsl */`
@@ -57,7 +172,7 @@ fn word(byte_offset: u32) -> u32 {
   let page = byte_offset / params.page_bytes;
   if (page >= params.page_count) { return 0u; }
   let slot = page_table[page];
-  if (slot == 0xffffffffu || slot >= params.resident_pages) { return 0u; }
+  if (slot == 0xffffffffu || slot >= params.resident_pages) { /* NANOVDB_PAGE_FAULT */ return 0u; }
   let physical = slot * params.page_bytes + (byte_offset % params.page_bytes);
   if (physical > params.resident_pages * params.page_bytes - 4u) { return 0u; }
   return grid[physical >> 2u];
@@ -148,6 +263,11 @@ fn sample_density_linear(position: vec3<f32>) -> vec2<f32> {
 }
 `;
 
+const pageFeedbackTraversalWGSL = traversalWGSL.replaceAll(
+  "/* NANOVDB_PAGE_FAULT */",
+  "nanovdb_record_page_fault(page);",
+);
+
 function specializeFloatTraversal(prefix: string, gridName: string, pageTableName: string, parameterPrefix: string): string {
   let source = traversalWGSL
     .replaceAll("grid[", `${gridName}[`)
@@ -255,9 +375,10 @@ export function uploadNanoVDBFloat32Grid(device: GPUDevice, payload: ArrayBuffer
   const buffer = device.createBuffer({ label: "NanoVDB Float32 grid", size: payload.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, mappedAtCreation: true });
   new Uint8Array(buffer.getMappedRange()).set(new Uint8Array(payload));
   buffer.unmap();
-  const pageTable = device.createBuffer({ label: "NanoVDB direct page table", size: 4, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, mappedAtCreation: true });
+  const pageTable = device.createBuffer({ label: "NanoVDB direct page table", size: 4, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC, mappedAtCreation: true });
   new Uint32Array(pageTable.getMappedRange())[0] = 0;
   pageTable.unmap();
+  let disposed = false;
   return {
     buffer,
     pageTable,
@@ -275,9 +396,17 @@ export function uploadNanoVDBFloat32Grid(device: GPUDevice, payload: ArrayBuffer
       if (pageIndex !== 0 || (data && data.byteLength !== payload.byteLength)) throw new Error("NANOVDB_STREAM_INCOMPLETE: direct NanoVDB grid has one immutable page");
     },
     touchPage: (pageIndex) => pageIndex === 0,
+    beginFrame: () => undefined,
+    pinPage: (pageIndex) => pageIndex === 0,
+    endFrame: () => undefined,
     evictPage: () => { throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: direct NanoVDB grid cannot evict its only page"); },
     hasResidentPage: (pageIndex) => pageIndex === 0,
-    dispose: () => { buffer.destroy(); pageTable.destroy(); },
+    dispose: () => {
+      if (disposed) return;
+      disposed = true;
+      buffer.destroy();
+      pageTable.destroy();
+    },
   };
 }
 
@@ -322,7 +451,7 @@ export function createNanoVDBFloat32GridPaged(
   const pageTableBytes = Math.max(4, pageCount * 4);
   try {
     buffer = device.createBuffer({ label: "NanoVDB paged Float32 grid", size: physicalBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
-    pageTable = device.createBuffer({ label: "NanoVDB page table", size: pageTableBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, mappedAtCreation: true });
+    pageTable = device.createBuffer({ label: "NanoVDB page table", size: pageTableBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC, mappedAtCreation: true });
     new Uint32Array(pageTable.getMappedRange()).fill(0xffffffff);
     pageTable.unmap();
   }
@@ -335,6 +464,9 @@ export function createNanoVDBFloat32GridPaged(
   const lastUsed = new Map<number, number>();
   let clock = 0;
   let evictions = 0;
+  let frameActive = false;
+  let disposed = false;
+  const framePins = new Set<number>();
   const writePageTable = (pageIndex: number, slot: number): void => {
     device.queue.writeBuffer(pageTable, pageIndex * 4, new Uint32Array([slot]));
   };
@@ -356,8 +488,12 @@ export function createNanoVDBFloat32GridPaged(
     const existingSlot = resident.get(pageIndex);
     let slot = existingSlot ?? [...Array(residentPageCount).keys()].find((candidate) => !residentHasSlot(candidate));
     if (slot === undefined) {
-      const oldest = [...lastUsed].sort((left, right) => left[1] - right[1] || left[0] - right[0])[0];
-      if (!oldest) throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: no resident NanoVDB page slot is available");
+      const oldest = [...lastUsed]
+        .filter(([candidate]) => !framePins.has(candidate))
+        .sort((left, right) => left[1] - right[1] || left[0] - right[0])[0];
+      if (!oldest) throw new Error(frameActive
+        ? "NANOVDB_GPU_BUDGET_EXCEEDED: all resident NanoVDB pages are pinned"
+        : "NANOVDB_GPU_BUDGET_EXCEEDED: no resident NanoVDB page slot is available");
       slot = resident.get(oldest[0]);
       evict(oldest[0], true);
     }
@@ -386,9 +522,25 @@ export function createNanoVDBFloat32GridPaged(
       lastUsed.set(pageIndex, ++clock);
       return true;
     },
+    beginFrame: () => { frameActive = true; framePins.clear(); },
+    pinPage: (pageIndex) => {
+      if (!frameActive || !resident.has(pageIndex)) return false;
+      framePins.add(pageIndex);
+      return true;
+    },
+    endFrame: () => { frameActive = false; framePins.clear(); },
     evictPage: evict,
     hasResidentPage: (pageIndex) => resident.has(pageIndex),
-    dispose: () => { resident.clear(); lastUsed.clear(); buffer.destroy(); pageTable.destroy(); },
+    dispose: () => {
+      if (disposed) return;
+      disposed = true;
+      frameActive = false;
+      framePins.clear();
+      resident.clear();
+      lastUsed.clear();
+      buffer.destroy();
+      pageTable.destroy();
+    },
   };
 }
 
@@ -499,7 +651,97 @@ function paramsBuffer(device: GPUDevice, values: Uint32Array): GPUBuffer {
   return buffer;
 }
 
-export async function readNanoVDBWordsWebGPU(device: GPUDevice, uploaded: NanoVDBWebGPUGrid, byteOffsets: readonly number[]): Promise<number[]> {
+export function createNanoVDBPageFeedbackGPUBuffer(device: GPUDevice, capacity = 1024): GPUBuffer {
+  const initial = createNanoVDBPageFeedbackBuffer(capacity);
+  let buffer: GPUBuffer | undefined;
+  try {
+    buffer = device.createBuffer({
+      label: "NanoVDB page feedback",
+      size: initial.byteLength,
+      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
+    });
+    device.queue.writeBuffer(buffer, 0, initial);
+    return buffer;
+  }
+  catch (error) {
+    buffer?.destroy();
+    throw error;
+  }
+}
+
+export interface NanoVDBPagedRenderResources {
+  grid: NanoVDBWebGPUGrid;
+  feedbackBuffer: GPUBuffer;
+  feedbackCapacity: number;
+  dispose(): void;
+}
+
+export function createNanoVDBPagedRenderResources(
+  device: GPUDevice,
+  byteLength: number,
+  pageByteLength: number,
+  maxResidentBytes: number,
+  feedbackCapacity = 1024,
+): NanoVDBPagedRenderResources {
+  let grid: NanoVDBWebGPUGrid | undefined;
+  let feedbackBuffer: GPUBuffer | undefined;
+  try {
+    grid = createNanoVDBFloat32GridPaged(device, byteLength, pageByteLength, maxResidentBytes);
+    feedbackBuffer = createNanoVDBPageFeedbackGPUBuffer(device, feedbackCapacity);
+  }
+  catch (error) {
+    feedbackBuffer?.destroy();
+    grid?.dispose();
+    throw error;
+  }
+  let disposed = false;
+  return {
+    grid,
+    feedbackBuffer,
+    feedbackCapacity,
+    dispose: () => {
+      if (disposed) return;
+      disposed = true;
+      feedbackBuffer.destroy();
+      grid.dispose();
+    },
+  };
+}
+
+export async function readNanoVDBPageFeedbackGPUBuffer(
+  device: GPUDevice,
+  feedback: GPUBuffer,
+  capacity: number,
+  pageCount: number,
+  renderRevision: number,
+): Promise<NanoVDBPageFeedbackBatch> {
+  const byteLength = nanoVDBPageFeedbackByteLength(capacity);
+  const readback = device.createBuffer({
+    label: `NanoVDB page feedback readback revision ${renderRevision}`,
+    size: byteLength,
+    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
+  });
+  const encoder = device.createCommandEncoder();
+  encoder.copyBufferToBuffer(feedback, 0, readback, 0, byteLength);
+  device.queue.submit([encoder.finish()]);
+  await readback.mapAsync(GPUMapMode.READ);
+  const buffer = readback.getMappedRange().slice(0);
+  readback.unmap();
+  readback.destroy();
+  return parseNanoVDBPageFeedbackBatch(buffer, pageCount, renderRevision);
+}
+
+function resolveNanoVDBPageFeedbackGPUBuffer(device: GPUDevice, value: GPUBuffer | undefined): { buffer: GPUBuffer; owned: boolean } {
+  if (value) return { buffer: value, owned: false };
+  return { buffer: createNanoVDBPageFeedbackGPUBuffer(device, 1), owned: true };
+}
+
+export async function readNanoVDBWordsWebGPU(
+  device: GPUDevice,
+  uploaded: NanoVDBWebGPUGrid,
+  byteOffsets: readonly number[],
+  pageFeedback?: GPUBuffer,
+): Promise<number[]> {
   if (byteOffsets.length < 1 || byteOffsets.length > 4096 || byteOffsets.some((offset) => !Number.isSafeInteger(offset) || offset < 0 || offset > uploaded.byteLength - 4 || offset % 4 !== 0)) {
     throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: paged word read offsets");
   }
@@ -510,6 +752,7 @@ export async function readNanoVDBWordsWebGPU(device: GPUDevice, uploaded: NanoVD
   const resultBuffer = device.createBuffer({ size: resultBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC });
   const readback = device.createBuffer({ size: resultBytes, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
   const params = paramsBuffer(device, new Uint32Array([uploaded.byteLength, byteOffsets.length, 0, 0, uploaded.pageByteLength, uploaded.pageCount, uploaded.residentPageCapacity, uploaded.paged ? 1 : 0]));
+  const feedback = resolveNanoVDBPageFeedbackGPUBuffer(device, pageFeedback);
   const module = device.createShaderModule({ label: "NanoVDB paged word reader", code: /* wgsl */`
 struct Params { data_bytes: u32, count: u32, width: u32, height: u32, page_bytes: u32, page_count: u32, resident_pages: u32, paged: u32 }
 @group(0) @binding(0) var<storage, read> grid: array<u32>;
@@ -517,7 +760,10 @@ struct Params { data_bytes: u32, count: u32, width: u32, height: u32, page_bytes
 @group(0) @binding(2) var<storage, read_write> results: array<u32>;
 @group(0) @binding(3) var<uniform> params: Params;
 @group(0) @binding(4) var<storage, read> page_table: array<u32>;
-${traversalWGSL}
+${NANOVDB_PAGE_FEEDBACK_WGSL}
+@group(0) @binding(5) var<storage, read_write> nanovdb_page_feedback: NanoVDBPageFeedback;
+${NANOVDB_PAGE_FEEDBACK_RECORD_WGSL}
+${pageFeedbackTraversalWGSL}
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) id: vec3<u32>) {
   if (id.x < params.count) { results[id.x] = word(offsets[id.x]); }
@@ -527,6 +773,7 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>) {
     { binding: 0, resource: { buffer: uploaded.buffer } }, { binding: 1, resource: { buffer: offsetBuffer } },
     { binding: 2, resource: { buffer: resultBuffer } }, { binding: 3, resource: { buffer: params } },
     { binding: 4, resource: { buffer: uploaded.pageTable } },
+    { binding: 5, resource: { buffer: feedback.buffer } },
   ] });
   const encoder = device.createCommandEncoder();
   const pass = encoder.beginComputePass();
@@ -537,10 +784,16 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>) {
   const result = [...new Uint32Array(readback.getMappedRange().slice(0))];
   readback.unmap();
   offsetBuffer.destroy(); resultBuffer.destroy(); readback.destroy(); params.destroy();
+  if (feedback.owned) feedback.buffer.destroy();
   return result;
 }
 
-export async function sampleNanoVDBFloat32WebGPU(device: GPUDevice, uploaded: NanoVDBWebGPUGrid, coordinates: Array<readonly [number, number, number]>): Promise<Array<{ value: number; active: boolean; valid: boolean }>> {
+export async function sampleNanoVDBFloat32WebGPU(
+  device: GPUDevice,
+  uploaded: NanoVDBWebGPUGrid,
+  coordinates: Array<readonly [number, number, number]>,
+  pageFeedback?: GPUBuffer,
+): Promise<Array<{ value: number; active: boolean; valid: boolean }>> {
   if (coordinates.length < 1 || coordinates.length > 4096) throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: sample count");
   const coordinateData = new Int32Array(coordinates.length * 4);
   coordinates.forEach((coord, index) => coordinateData.set(coord, index * 4));
@@ -550,6 +803,7 @@ export async function sampleNanoVDBFloat32WebGPU(device: GPUDevice, uploaded: Na
   const resultBuffer = device.createBuffer({ size: resultBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC });
   const readback = device.createBuffer({ size: resultBytes, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
   const params = paramsBuffer(device, new Uint32Array([uploaded.byteLength, coordinates.length, 0, 0, uploaded.pageByteLength, uploaded.pageCount, uploaded.residentPageCapacity, uploaded.paged ? 1 : 0]));
+  const feedback = resolveNanoVDBPageFeedbackGPUBuffer(device, pageFeedback);
   const module = device.createShaderModule({ label: "NanoVDB Float32 sampler", code: /* wgsl */`
 struct Params { data_bytes: u32, count: u32, width: u32, height: u32, page_bytes: u32, page_count: u32, resident_pages: u32, paged: u32 }
 @group(0) @binding(0) var<storage, read> grid: array<u32>;
@@ -557,7 +811,10 @@ struct Params { data_bytes: u32, count: u32, width: u32, height: u32, page_bytes
 @group(0) @binding(2) var<storage, read_write> results: array<vec4<f32>>;
 @group(0) @binding(3) var<uniform> params: Params;
 @group(0) @binding(4) var<storage, read> page_table: array<u32>;
-${traversalWGSL}
+${NANOVDB_PAGE_FEEDBACK_WGSL}
+@group(0) @binding(5) var<storage, read_write> nanovdb_page_feedback: NanoVDBPageFeedback;
+${NANOVDB_PAGE_FEEDBACK_RECORD_WGSL}
+${pageFeedbackTraversalWGSL}
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) id: vec3<u32>) {
   if (id.x >= params.count) { return; }
@@ -569,6 +826,7 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>) {
     { binding: 0, resource: { buffer: uploaded.buffer } }, { binding: 1, resource: { buffer: coordinateBuffer } },
     { binding: 2, resource: { buffer: resultBuffer } }, { binding: 3, resource: { buffer: params } },
     { binding: 4, resource: { buffer: uploaded.pageTable } },
+    { binding: 5, resource: { buffer: feedback.buffer } },
   ] });
   const encoder = device.createCommandEncoder();
   const pass = encoder.beginComputePass();
@@ -579,6 +837,7 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>) {
   const values = new Float32Array(readback.getMappedRange().slice(0));
   readback.unmap();
   coordinateBuffer.destroy(); resultBuffer.destroy(); readback.destroy(); params.destroy();
+  if (feedback.owned) feedback.buffer.destroy();
   return coordinates.map((_coord, index) => ({ value: values[index * 4], active: values[index * 4 + 1] > 0.5, valid: values[index * 4 + 1] >= 0 }));
 }