Advance WebGPU volume and bounded workflows
This commit is contained in:
mes123456
1
web/app/src/render/RenderAssets.ts
Normal file
1
web/app/src/render/RenderAssets.ts
Normal file
@@ -0,0 +1 @@
|
||||
export * from "../../../protocol/render-assets";
|
||||
615
web/app/src/render/nanovdb-volume-renderer.ts
Normal file
615
web/app/src/render/nanovdb-volume-renderer.ts
Normal file
@@ -0,0 +1,615 @@
|
||||
import type { NanoVDBGridIR, NanoVDBMaterialIR } from "../../../protocol/volume-vdb";
|
||||
|
||||
export interface NanoVDBWebGPUCapabilityIR {
|
||||
available: boolean;
|
||||
reason?: string;
|
||||
maxStorageBufferBindingSize?: number;
|
||||
maxBufferSize?: number;
|
||||
}
|
||||
|
||||
export interface NanoVDBWebGPUGrid {
|
||||
buffer: GPUBuffer;
|
||||
pageTable: GPUBuffer;
|
||||
byteLength: number;
|
||||
pageByteLength: number;
|
||||
pageCount: number;
|
||||
residentPageCount: number;
|
||||
residentPageCapacity: number;
|
||||
paged: boolean;
|
||||
residentVirtualPages: readonly number[];
|
||||
uploadPage(pageIndex: number, data?: ArrayBuffer): void;
|
||||
evictPage(pageIndex: number): void;
|
||||
hasResidentPage(pageIndex: number): boolean;
|
||||
dispose(): void;
|
||||
}
|
||||
|
||||
export interface NanoVDBMaterialGridUploadsIR {
|
||||
temperature?: NanoVDBWebGPUGrid;
|
||||
color?: NanoVDBWebGPUGrid;
|
||||
emission?: NanoVDBWebGPUGrid;
|
||||
}
|
||||
|
||||
export interface NanoVDBGpuPageAllocatorStatsIR {
|
||||
residentBytes: number;
|
||||
maxResidentBytes: number;
|
||||
residentPages: number;
|
||||
evictions: number;
|
||||
keys: string[];
|
||||
}
|
||||
|
||||
export interface NanoVDBDeviceLossIR { reason?: string; message: string }
|
||||
|
||||
const traversalWGSL = /* wgsl */`
|
||||
fn in_range(byte_offset: u32, byte_length: u32) -> bool {
|
||||
return byte_offset <= params.data_bytes && byte_length <= params.data_bytes - byte_offset;
|
||||
}
|
||||
fn word(byte_offset: u32) -> u32 {
|
||||
if ((byte_offset & 3u) != 0u || !in_range(byte_offset, 4u)) { return 0u; }
|
||||
if (params.paged == 0u) { return grid[byte_offset >> 2u]; }
|
||||
if (params.page_bytes == 0u) { return 0u; }
|
||||
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; }
|
||||
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];
|
||||
}
|
||||
fn scalar(byte_offset: u32) -> f32 { return bitcast<f32>(word(byte_offset)); }
|
||||
fn mask_on(byte_offset: u32, index: u32) -> bool {
|
||||
let address = byte_offset + (index >> 5u) * 4u;
|
||||
return in_range(address, 4u) && (word(address) & (1u << (index & 31u))) != 0u;
|
||||
}
|
||||
fn valid_grid() -> bool {
|
||||
return params.data_bytes >= 736u && word(0u) == 0x6f6e614eu && word(4u) == 0x31424456u &&
|
||||
(word(16u) >> 21u) == 32u && word(32u) == params.data_bytes && word(36u) == 0u;
|
||||
}
|
||||
fn root_key(coord: vec3<i32>) -> vec2<u32> {
|
||||
let x = bitcast<u32>(coord.x) >> 12u;
|
||||
let y = bitcast<u32>(coord.y) >> 12u;
|
||||
let z = bitcast<u32>(coord.z) >> 12u;
|
||||
return vec2<u32>(z | ((y & 0x7ffu) << 21u), (y >> 11u) | (x << 10u));
|
||||
}
|
||||
fn key_less(a: vec2<u32>, b: vec2<u32>) -> bool { return a.y < b.y || (a.y == b.y && a.x < b.x); }
|
||||
fn child_address(parent: u32, offset_address: u32, child_bytes: u32) -> u32 {
|
||||
let low = word(offset_address);
|
||||
let high = word(offset_address + 4u);
|
||||
if (low == 0u || high != 0u || low > params.data_bytes || parent > params.data_bytes - low) { return 0xffffffffu; }
|
||||
let child = parent + low;
|
||||
if (!in_range(child, child_bytes)) { return 0xffffffffu; }
|
||||
return child;
|
||||
}
|
||||
fn sample_density(coord: vec3<i32>) -> vec2<f32> {
|
||||
if (!valid_grid()) { return vec2<f32>(0.0, -1.0); }
|
||||
let tree = 672u;
|
||||
let root = child_address(tree, tree + 24u, 64u);
|
||||
if (root == 0xffffffffu) { return vec2<f32>(0.0, -1.0); }
|
||||
let count = word(root + 24u);
|
||||
if (count > (params.data_bytes - root - 64u) / 32u) { return vec2<f32>(0.0, -1.0); }
|
||||
let wanted = root_key(coord);
|
||||
var low = 0u;
|
||||
var high = count;
|
||||
var tile = 0xffffffffu;
|
||||
for (var iteration = 0u; iteration < 32u && low < high; iteration++) {
|
||||
let middle = low + (high - low) / 2u;
|
||||
let address = root + 64u + middle * 32u;
|
||||
let candidate = vec2<u32>(word(address), word(address + 4u));
|
||||
if (all(candidate == wanted)) { tile = address; break; }
|
||||
if (key_less(wanted, candidate)) { low = middle + 1u; } else { high = middle; }
|
||||
}
|
||||
if (tile == 0xffffffffu) { return vec2<f32>(scalar(root + 28u), 0.0); }
|
||||
let root_child_low = word(tile + 8u);
|
||||
let root_child_high = word(tile + 12u);
|
||||
if (root_child_low == 0u && root_child_high == 0u) { return vec2<f32>(scalar(tile + 20u), select(0.0, 1.0, word(tile + 16u) != 0u)); }
|
||||
let upper = child_address(root, tile + 8u, 270400u);
|
||||
if (upper == 0xffffffffu) { return vec2<f32>(0.0, -1.0); }
|
||||
let ux = (bitcast<u32>(coord.x) & 4095u) >> 7u;
|
||||
let uy = (bitcast<u32>(coord.y) & 4095u) >> 7u;
|
||||
let uz = (bitcast<u32>(coord.z) & 4095u) >> 7u;
|
||||
let upper_index = (ux << 10u) | (uy << 5u) | uz;
|
||||
if (!mask_on(upper + 4128u, upper_index)) { return vec2<f32>(scalar(upper + 8256u + upper_index * 8u), select(0.0, 1.0, mask_on(upper + 32u, upper_index))); }
|
||||
let lower = child_address(upper, upper + 8256u + upper_index * 8u, 33856u);
|
||||
if (lower == 0xffffffffu) { return vec2<f32>(0.0, -1.0); }
|
||||
let lx = (bitcast<u32>(coord.x) & 127u) >> 3u;
|
||||
let ly = (bitcast<u32>(coord.y) & 127u) >> 3u;
|
||||
let lz = (bitcast<u32>(coord.z) & 127u) >> 3u;
|
||||
let lower_index = (lx << 8u) | (ly << 4u) | lz;
|
||||
if (!mask_on(lower + 544u, lower_index)) { return vec2<f32>(scalar(lower + 1088u + lower_index * 8u), select(0.0, 1.0, mask_on(lower + 32u, lower_index))); }
|
||||
let leaf = child_address(lower, lower + 1088u + lower_index * 8u, 2144u);
|
||||
if (leaf == 0xffffffffu) { return vec2<f32>(0.0, -1.0); }
|
||||
let voxel = ((bitcast<u32>(coord.x) & 7u) << 6u) | ((bitcast<u32>(coord.y) & 7u) << 3u) | (bitcast<u32>(coord.z) & 7u);
|
||||
return vec2<f32>(scalar(leaf + 96u + voxel * 4u), select(0.0, 1.0, mask_on(leaf + 16u, voxel)));
|
||||
}
|
||||
fn sample_density_linear(position: vec3<f32>) -> vec2<f32> {
|
||||
let base = vec3<i32>(floor(position));
|
||||
let fraction = position - vec3<f32>(base);
|
||||
var value = 0.0;
|
||||
var activity = 0.0;
|
||||
for (var x = 0i; x < 2i; x += 1i) {
|
||||
for (var y = 0i; y < 2i; y += 1i) {
|
||||
for (var z = 0i; z < 2i; z += 1i) {
|
||||
let sample = sample_density(base + vec3<i32>(x, y, z));
|
||||
if (sample.y < 0.0) { return vec2<f32>(0.0, -1.0); }
|
||||
let offset = vec3<f32>(f32(x), f32(y), f32(z));
|
||||
let weight3 = select(vec3<f32>(1.0) - fraction, fraction, offset == vec3<f32>(1.0));
|
||||
value += sample.x * weight3.x * weight3.y * weight3.z;
|
||||
activity = max(activity, sample.y);
|
||||
}
|
||||
}
|
||||
}
|
||||
return vec2<f32>(value, activity);
|
||||
}
|
||||
`;
|
||||
|
||||
function specializeFloatTraversal(prefix: string, gridName: string, pageTableName: string, parameterPrefix: string): string {
|
||||
let source = traversalWGSL
|
||||
.replaceAll("grid[", `${gridName}[`)
|
||||
.replaceAll("page_table[page]", pageTableName ? `${pageTableName}[page]` : "page")
|
||||
.replaceAll("params.data_bytes", `params.${parameterPrefix}_data_bytes`)
|
||||
.replaceAll("params.page_bytes", `params.${parameterPrefix}_page_bytes`)
|
||||
.replaceAll("params.page_count", `params.${parameterPrefix}_page_count`)
|
||||
.replaceAll("params.resident_pages", `params.${parameterPrefix}_resident_pages`)
|
||||
.replaceAll("params.paged", `params.${parameterPrefix}_paged`);
|
||||
for (const name of ["sample_density_linear", "sample_density", "child_address", "valid_grid", "root_key", "key_less", "in_range", "mask_on", "scalar", "word"]) {
|
||||
source = source.replace(new RegExp(`\\b${name}\\b`, "g"), `${prefix}_${name}`);
|
||||
}
|
||||
return source;
|
||||
}
|
||||
|
||||
const temperatureTraversalWGSL = specializeFloatTraversal("temperature", "temperature_grid", "", "temperature");
|
||||
const emissionTraversalWGSL = specializeFloatTraversal("emission", "emission_grid", "", "emission");
|
||||
|
||||
const vec3TraversalWGSL = /* wgsl */`
|
||||
fn color_in_range(byte_offset: u32, byte_length: u32) -> bool {
|
||||
return byte_offset <= params.color_data_bytes && byte_length <= params.color_data_bytes - byte_offset;
|
||||
}
|
||||
fn color_word(byte_offset: u32) -> u32 {
|
||||
if ((byte_offset & 3u) != 0u || !color_in_range(byte_offset, 4u) || params.color_page_bytes == 0u) { return 0u; }
|
||||
let page = byte_offset / params.color_page_bytes;
|
||||
if (page >= params.color_page_count) { return 0u; }
|
||||
let slot = page;
|
||||
if (slot == 0xffffffffu || slot >= params.color_resident_pages) { return 0u; }
|
||||
let physical = slot * params.color_page_bytes + (byte_offset % params.color_page_bytes);
|
||||
if (physical > params.color_resident_pages * params.color_page_bytes - 4u) { return 0u; }
|
||||
return color_grid[physical >> 2u];
|
||||
}
|
||||
fn color_scalar(byte_offset: u32) -> f32 { return bitcast<f32>(color_word(byte_offset)); }
|
||||
fn color_vec3(byte_offset: u32) -> vec3<f32> { return vec3<f32>(color_scalar(byte_offset), color_scalar(byte_offset + 4u), color_scalar(byte_offset + 8u)); }
|
||||
fn color_mask_on(byte_offset: u32, index: u32) -> bool { return (color_word(byte_offset + (index >> 5u) * 4u) & (1u << (index & 31u))) != 0u; }
|
||||
fn color_valid_grid() -> bool {
|
||||
return params.color_data_bytes >= 768u && color_word(0u) == 0x6f6e614eu && color_word(4u) == 0x31424456u &&
|
||||
(color_word(16u) >> 21u) == 32u && color_word(32u) == params.color_data_bytes && color_word(36u) == 0u;
|
||||
}
|
||||
fn color_root_key(coord: vec3<i32>) -> vec2<u32> {
|
||||
let x = bitcast<u32>(coord.x) >> 12u; let y = bitcast<u32>(coord.y) >> 12u; let z = bitcast<u32>(coord.z) >> 12u;
|
||||
return vec2<u32>(z | ((y & 0x7ffu) << 21u), (y >> 11u) | (x << 10u));
|
||||
}
|
||||
fn color_key_less(a: vec2<u32>, b: vec2<u32>) -> bool { return a.y < b.y || (a.y == b.y && a.x < b.x); }
|
||||
fn color_child_address(parent: u32, offset_address: u32, child_bytes: u32) -> u32 {
|
||||
let low = color_word(offset_address); let high = color_word(offset_address + 4u);
|
||||
if (low == 0u || high != 0u || low > params.color_data_bytes || parent > params.color_data_bytes - low) { return 0xffffffffu; }
|
||||
let child = parent + low; if (!color_in_range(child, child_bytes)) { return 0xffffffffu; } return child;
|
||||
}
|
||||
fn sample_color(coord: vec3<i32>) -> vec4<f32> {
|
||||
if (!color_valid_grid()) { return vec4<f32>(0.0, 0.0, 0.0, -1.0); }
|
||||
let tree = 672u; let root = color_child_address(tree, tree + 24u, 96u);
|
||||
if (root == 0xffffffffu) { return vec4<f32>(0.0, 0.0, 0.0, -1.0); }
|
||||
let count = color_word(root + 24u);
|
||||
if (count > (params.color_data_bytes - root - 96u) / 32u) { return vec4<f32>(0.0, 0.0, 0.0, -1.0); }
|
||||
let wanted = color_root_key(coord); var low = 0u; var high = count; var tile = 0xffffffffu;
|
||||
for (var iteration = 0u; iteration < 32u && low < high; iteration++) {
|
||||
let middle = low + (high - low) / 2u; let address = root + 96u + middle * 32u;
|
||||
let candidate = vec2<u32>(color_word(address), color_word(address + 4u));
|
||||
if (all(candidate == wanted)) { tile = address; break; }
|
||||
if (color_key_less(wanted, candidate)) { low = middle + 1u; } else { high = middle; }
|
||||
}
|
||||
if (tile == 0xffffffffu) { return vec4<f32>(0.0); }
|
||||
let root_child_low = color_word(tile + 8u); let root_child_high = color_word(tile + 12u);
|
||||
if (root_child_low == 0u && root_child_high == 0u) { return vec4<f32>(color_vec3(tile + 20u), select(0.0, 1.0, color_word(tile + 16u) != 0u)); }
|
||||
let upper = color_child_address(root, tile + 8u, 532544u); if (upper == 0xffffffffu) { return vec4<f32>(0.0, 0.0, 0.0, -1.0); }
|
||||
let upper_index = (((bitcast<u32>(coord.x) & 4095u) >> 7u) << 10u) | (((bitcast<u32>(coord.y) & 4095u) >> 7u) << 5u) | ((bitcast<u32>(coord.z) & 4095u) >> 7u);
|
||||
if (!color_mask_on(upper + 4128u, upper_index)) { return vec4<f32>(color_vec3(upper + 8256u + upper_index * 16u), select(0.0, 1.0, color_mask_on(upper + 32u, upper_index))); }
|
||||
let lower = color_child_address(upper, upper + 8256u + upper_index * 16u, 66624u); if (lower == 0xffffffffu) { return vec4<f32>(0.0, 0.0, 0.0, -1.0); }
|
||||
let lower_index = (((bitcast<u32>(coord.x) & 127u) >> 3u) << 8u) | (((bitcast<u32>(coord.y) & 127u) >> 3u) << 4u) | ((bitcast<u32>(coord.z) & 127u) >> 3u);
|
||||
if (!color_mask_on(lower + 544u, lower_index)) { return vec4<f32>(color_vec3(lower + 1088u + lower_index * 16u), select(0.0, 1.0, color_mask_on(lower + 32u, lower_index))); }
|
||||
let leaf = color_child_address(lower, lower + 1088u + lower_index * 16u, 6272u); if (leaf == 0xffffffffu) { return vec4<f32>(0.0, 0.0, 0.0, -1.0); }
|
||||
let voxel = ((bitcast<u32>(coord.x) & 7u) << 6u) | ((bitcast<u32>(coord.y) & 7u) << 3u) | (bitcast<u32>(coord.z) & 7u);
|
||||
return vec4<f32>(color_vec3(leaf + 128u + voxel * 12u), select(0.0, 1.0, color_mask_on(leaf + 16u, voxel)));
|
||||
}
|
||||
fn sample_color_linear(position: vec3<f32>) -> vec4<f32> {
|
||||
let base = vec3<i32>(floor(position)); let fraction = position - vec3<f32>(base); var value = vec3<f32>(0.0); var activity = 0.0;
|
||||
for (var x = 0i; x < 2i; x += 1i) { for (var y = 0i; y < 2i; y += 1i) { for (var z = 0i; z < 2i; z += 1i) {
|
||||
let sample = sample_color(base + vec3<i32>(x, y, z)); if (sample.w < 0.0) { return vec4<f32>(0.0, 0.0, 0.0, -1.0); }
|
||||
let offset = vec3<f32>(f32(x), f32(y), f32(z)); let weight3 = select(vec3<f32>(1.0) - fraction, fraction, offset == vec3<f32>(1.0));
|
||||
value += sample.xyz * weight3.x * weight3.y * weight3.z; activity = max(activity, sample.w);
|
||||
}}}
|
||||
return vec4<f32>(value, activity);
|
||||
}
|
||||
`;
|
||||
|
||||
export async function probeNanoVDBWebGPU(requiredBytes = 1): Promise<{ capability: NanoVDBWebGPUCapabilityIR; adapter?: GPUAdapter; device?: GPUDevice }> {
|
||||
if (!navigator.gpu) return { capability: { available: false, reason: "WebGPU is unavailable" } };
|
||||
const adapter = await navigator.gpu.requestAdapter({ powerPreference: "high-performance" });
|
||||
if (!adapter) return { capability: { available: false, reason: "No WebGPU adapter is available" } };
|
||||
const maxStorageBufferBindingSize = Number(adapter.limits.maxStorageBufferBindingSize);
|
||||
const maxBufferSize = Number(adapter.limits.maxBufferSize);
|
||||
if (requiredBytes > maxStorageBufferBindingSize || requiredBytes > maxBufferSize) return { capability: { available: false, reason: "NanoVDB grid exceeds WebGPU adapter limits", maxStorageBufferBindingSize, maxBufferSize } };
|
||||
const device = await adapter.requestDevice({ requiredLimits: { maxStorageBufferBindingSize: requiredBytes, maxBufferSize: requiredBytes } });
|
||||
return { capability: { available: true, maxStorageBufferBindingSize, maxBufferSize }, adapter, device };
|
||||
}
|
||||
|
||||
export function uploadNanoVDBFloat32Grid(device: GPUDevice, payload: ArrayBuffer): NanoVDBWebGPUGrid {
|
||||
if (payload.byteLength === 0 || payload.byteLength % 32 !== 0 || payload.byteLength > device.limits.maxStorageBufferBindingSize || payload.byteLength > device.limits.maxBufferSize) throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: grid payload cannot be uploaded");
|
||||
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 });
|
||||
new Uint32Array(pageTable.getMappedRange())[0] = 0;
|
||||
pageTable.unmap();
|
||||
return {
|
||||
buffer,
|
||||
pageTable,
|
||||
byteLength: payload.byteLength,
|
||||
pageByteLength: payload.byteLength,
|
||||
pageCount: 1,
|
||||
residentPageCount: 1,
|
||||
residentPageCapacity: 1,
|
||||
paged: false,
|
||||
residentVirtualPages: [0],
|
||||
uploadPage: (pageIndex, data) => {
|
||||
if (pageIndex !== 0 || (data && data.byteLength !== payload.byteLength)) throw new Error("NANOVDB_STREAM_INCOMPLETE: direct NanoVDB grid has one immutable page");
|
||||
},
|
||||
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(); },
|
||||
};
|
||||
}
|
||||
|
||||
export function uploadNanoVDBFloat32GridPaged(
|
||||
device: GPUDevice,
|
||||
payload: ArrayBuffer,
|
||||
pageByteLength: number,
|
||||
maxResidentBytes: number,
|
||||
): NanoVDBWebGPUGrid {
|
||||
if (payload.byteLength === 0 || payload.byteLength % 32 !== 0 || !Number.isSafeInteger(pageByteLength) || pageByteLength < 64 * 1024 || pageByteLength % 32 !== 0) {
|
||||
throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: invalid paged grid layout");
|
||||
}
|
||||
if (!Number.isSafeInteger(maxResidentBytes) || maxResidentBytes < pageByteLength) {
|
||||
throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: invalid paged resident budget");
|
||||
}
|
||||
const pageCount = Math.ceil(payload.byteLength / pageByteLength);
|
||||
const residentPageCount = Math.min(pageCount, Math.max(1, Math.floor(maxResidentBytes / pageByteLength)));
|
||||
const physicalBytes = residentPageCount * pageByteLength;
|
||||
if (pageCount > 8192 || physicalBytes > device.limits.maxStorageBufferBindingSize || physicalBytes > device.limits.maxBufferSize) {
|
||||
throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: paged grid exceeds the resident or adapter budget");
|
||||
}
|
||||
const buffer = device.createBuffer({ label: "NanoVDB paged Float32 grid", size: physicalBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
|
||||
const pageTableBytes = Math.max(4, pageCount * 4);
|
||||
const pageTable = device.createBuffer({ label: "NanoVDB page table", size: pageTableBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, mappedAtCreation: true });
|
||||
const table = new Uint32Array(pageTable.getMappedRange());
|
||||
table.fill(0xffffffff);
|
||||
pageTable.unmap();
|
||||
const resident = new Map<number, number>();
|
||||
const upload = (pageIndex: number, data = payload.slice(pageIndex * pageByteLength, Math.min(payload.byteLength, (pageIndex + 1) * pageByteLength))): void => {
|
||||
if (!Number.isSafeInteger(pageIndex) || pageIndex < 0 || pageIndex >= pageCount || data.byteLength === 0 || data.byteLength > pageByteLength) {
|
||||
throw new Error("NANOVDB_STREAM_INCOMPLETE: NanoVDB page is outside the virtual grid");
|
||||
}
|
||||
const existingSlot = resident.get(pageIndex);
|
||||
const slot = existingSlot ?? [...Array(residentPageCount).keys()].find((candidate) => !residentHasSlot(candidate));
|
||||
if (slot === undefined) throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: no resident NanoVDB page slot is available");
|
||||
device.queue.writeBuffer(buffer, slot * pageByteLength, data);
|
||||
table[pageIndex] = slot;
|
||||
device.queue.writeBuffer(pageTable, pageIndex * 4, new Uint32Array([slot]));
|
||||
resident.set(pageIndex, slot);
|
||||
};
|
||||
const residentHasSlot = (slot: number): boolean => {
|
||||
for (const current of resident.values()) if (current === slot) return true;
|
||||
return false;
|
||||
};
|
||||
const initialPages = Math.min(pageCount, residentPageCount);
|
||||
for (let page = 0; page < initialPages; page++) upload(page);
|
||||
return {
|
||||
buffer,
|
||||
pageTable,
|
||||
byteLength: payload.byteLength,
|
||||
pageByteLength,
|
||||
pageCount,
|
||||
residentPageCount: resident.size,
|
||||
residentPageCapacity: residentPageCount,
|
||||
paged: true,
|
||||
get residentVirtualPages() { return [...resident.keys()].sort((a, b) => a - b); },
|
||||
uploadPage: upload,
|
||||
evictPage: (pageIndex) => {
|
||||
if (!resident.delete(pageIndex)) return;
|
||||
table[pageIndex] = 0xffffffff;
|
||||
device.queue.writeBuffer(pageTable, pageIndex * 4, new Uint32Array([0xffffffff]));
|
||||
},
|
||||
hasResidentPage: (pageIndex) => resident.has(pageIndex),
|
||||
dispose: () => { resident.clear(); buffer.destroy(); pageTable.destroy(); },
|
||||
};
|
||||
}
|
||||
|
||||
export class NanoVDBGpuPageAllocator {
|
||||
private readonly pages = new Map<string, { buffer: GPUBuffer; bytes: number; used: number }>();
|
||||
private clock = 0;
|
||||
private evictions = 0;
|
||||
|
||||
constructor(
|
||||
private readonly device: GPUDevice,
|
||||
readonly pageByteLength: number,
|
||||
readonly maxResidentBytes: number,
|
||||
) {
|
||||
if (!Number.isSafeInteger(pageByteLength) || pageByteLength < 64 * 1024 || pageByteLength % 32 !== 0 ||
|
||||
!Number.isSafeInteger(maxResidentBytes) || maxResidentBytes < pageByteLength) {
|
||||
throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: invalid GPU page allocator budget");
|
||||
}
|
||||
}
|
||||
|
||||
upload(key: string, data: ArrayBuffer): GPUBuffer {
|
||||
if (!key || data.byteLength === 0 || data.byteLength > this.pageByteLength) throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: invalid GPU page");
|
||||
const existing = this.pages.get(key);
|
||||
if (existing) { existing.used = ++this.clock; return existing.buffer; }
|
||||
while (this.residentBytes() + this.pageByteLength > this.maxResidentBytes) this.evictOldest();
|
||||
const buffer = this.device.createBuffer({ label: `NanoVDB page ${key}`, size: this.pageByteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, mappedAtCreation: true });
|
||||
new Uint8Array(buffer.getMappedRange()).set(new Uint8Array(data));
|
||||
buffer.unmap();
|
||||
this.pages.set(key, { buffer, bytes: this.pageByteLength, used: ++this.clock });
|
||||
return buffer;
|
||||
}
|
||||
|
||||
touch(key: string): boolean {
|
||||
const page = this.pages.get(key);
|
||||
if (!page) return false;
|
||||
page.used = ++this.clock;
|
||||
return true;
|
||||
}
|
||||
|
||||
has(key: string): boolean { return this.pages.has(key); }
|
||||
|
||||
stats(): NanoVDBGpuPageAllocatorStatsIR {
|
||||
return { residentBytes: this.residentBytes(), maxResidentBytes: this.maxResidentBytes, residentPages: this.pages.size, evictions: this.evictions, keys: [...this.pages.keys()].sort() };
|
||||
}
|
||||
|
||||
dispose(): void {
|
||||
for (const page of this.pages.values()) page.buffer.destroy();
|
||||
this.pages.clear();
|
||||
}
|
||||
|
||||
private residentBytes(): number { return [...this.pages.values()].reduce((sum, page) => sum + page.bytes, 0); }
|
||||
|
||||
private evictOldest(): void {
|
||||
const oldest = [...this.pages].sort((left, right) => left[1].used - right[1].used || left[0].localeCompare(right[0]))[0];
|
||||
if (!oldest) throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: no GPU page can be evicted");
|
||||
oldest[1].buffer.destroy();
|
||||
this.pages.delete(oldest[0]);
|
||||
this.evictions++;
|
||||
}
|
||||
}
|
||||
|
||||
export class NanoVDBWebGPUDeviceSession {
|
||||
device?: GPUDevice;
|
||||
generation = 0;
|
||||
status: "idle" | "ready" | "lost" | "disposed" = "idle";
|
||||
private loss?: Promise<NanoVDBDeviceLossIR>;
|
||||
private readonly lossListeners = new Set<(loss: NanoVDBDeviceLossIR) => void>();
|
||||
|
||||
onDeviceLost(listener: (loss: NanoVDBDeviceLossIR) => void): () => void {
|
||||
this.lossListeners.add(listener);
|
||||
return () => this.lossListeners.delete(listener);
|
||||
}
|
||||
|
||||
async open(requiredBytes: number): Promise<GPUDevice> {
|
||||
if (this.status === "disposed") throw new Error("VOLUME_SHADER_UNAVAILABLE: WebGPU session is disposed");
|
||||
const probe = await probeNanoVDBWebGPU(requiredBytes);
|
||||
if (!probe.capability.available || !probe.device) throw new Error(`VOLUME_SHADER_UNAVAILABLE: ${probe.capability.reason ?? "WebGPU unavailable"}`);
|
||||
this.device = probe.device;
|
||||
this.generation++;
|
||||
this.status = "ready";
|
||||
this.loss = probe.device.lost.then((info: NanoVDBDeviceLossIR) => {
|
||||
if (this.device === probe.device && this.status !== "disposed") this.status = "lost";
|
||||
for (const listener of this.lossListeners) listener(info);
|
||||
return info;
|
||||
});
|
||||
return probe.device;
|
||||
}
|
||||
|
||||
async waitForLoss(): Promise<NanoVDBDeviceLossIR> {
|
||||
if (!this.loss) throw new Error("VOLUME_SHADER_UNAVAILABLE: WebGPU session has not opened");
|
||||
return this.loss;
|
||||
}
|
||||
|
||||
async recover(requiredBytes: number): Promise<GPUDevice> {
|
||||
this.device?.destroy();
|
||||
return this.open(requiredBytes);
|
||||
}
|
||||
|
||||
dispose(): void {
|
||||
this.status = "disposed";
|
||||
this.device?.destroy();
|
||||
this.device = undefined;
|
||||
}
|
||||
}
|
||||
|
||||
function paramsBuffer(device: GPUDevice, values: Uint32Array): GPUBuffer {
|
||||
const buffer = device.createBuffer({ size: Math.max(16, values.byteLength), usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST });
|
||||
device.queue.writeBuffer(buffer, 0, values);
|
||||
return buffer;
|
||||
}
|
||||
|
||||
export async function sampleNanoVDBFloat32WebGPU(device: GPUDevice, uploaded: NanoVDBWebGPUGrid, coordinates: Array<readonly [number, number, number]>): 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));
|
||||
const coordinateBuffer = device.createBuffer({ size: coordinateData.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
|
||||
device.queue.writeBuffer(coordinateBuffer, 0, coordinateData);
|
||||
const resultBytes = coordinates.length * 16;
|
||||
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 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>;
|
||||
@group(0) @binding(1) var<storage, read> coords: array<vec4<i32>>;
|
||||
@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}
|
||||
@compute @workgroup_size(64)
|
||||
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
|
||||
if (id.x >= params.count) { return; }
|
||||
let sample = sample_density(coords[id.x].xyz);
|
||||
results[id.x] = vec4<f32>(sample.x, sample.y, 0.0, 0.0);
|
||||
}` });
|
||||
const pipeline = device.createComputePipeline({ layout: "auto", compute: { module, entryPoint: "main" } });
|
||||
const bindGroup = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [
|
||||
{ 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 } },
|
||||
] });
|
||||
const encoder = device.createCommandEncoder();
|
||||
const pass = encoder.beginComputePass();
|
||||
pass.setPipeline(pipeline); pass.setBindGroup(0, bindGroup); pass.dispatchWorkgroups(Math.ceil(coordinates.length / 64)); pass.end();
|
||||
encoder.copyBufferToBuffer(resultBuffer, 0, readback, 0, resultBytes);
|
||||
device.queue.submit([encoder.finish()]);
|
||||
await readback.mapAsync(GPUMapMode.READ);
|
||||
const values = new Float32Array(readback.getMappedRange().slice(0));
|
||||
readback.unmap();
|
||||
coordinateBuffer.destroy(); resultBuffer.destroy(); readback.destroy(); params.destroy();
|
||||
return coordinates.map((_coord, index) => ({ value: values[index * 4], active: values[index * 4 + 1] > 0.5, valid: values[index * 4 + 1] >= 0 }));
|
||||
}
|
||||
|
||||
export async function renderNanoVDBFloat32WebGPU(
|
||||
device: GPUDevice,
|
||||
uploaded: NanoVDBWebGPUGrid,
|
||||
gridDefinition: NanoVDBGridIR,
|
||||
material: NanoVDBMaterialIR,
|
||||
width = 96,
|
||||
height = 96,
|
||||
materialGrids: NanoVDBMaterialGridUploadsIR = {},
|
||||
): Promise<Uint8Array> {
|
||||
if (gridDefinition.valueType !== "FLOAT32" || width < 1 || height < 1 || width > 2048 || height > 2048) throw new Error("NANOVDB_GRID_UNSUPPORTED: bounded Float32 render input required");
|
||||
const outputBytes = width * height * 4;
|
||||
const output = device.createBuffer({ size: outputBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC });
|
||||
const readback = device.createBuffer({ size: outputBytes, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
|
||||
const paramsData = new ArrayBuffer(224);
|
||||
const u32 = new Uint32Array(paramsData);
|
||||
const i32 = new Int32Array(paramsData);
|
||||
const f32 = new Float32Array(paramsData);
|
||||
u32.set([uploaded.byteLength, material.interpolation === "LINEAR" ? 1 : 0, width, height], 0);
|
||||
u32.set([uploaded.pageByteLength, uploaded.pageCount, uploaded.residentPageCapacity, uploaded.paged ? 1 : 0], 4);
|
||||
const uploadFields = (grid: NanoVDBWebGPUGrid | undefined): [number, number, number, number, number, number, number, number] => grid
|
||||
? [grid.byteLength, grid.pageByteLength, grid.pageCount, grid.residentPageCapacity, grid.paged ? 1 : 0, 0, 0, 0]
|
||||
: [0, 0, 0, 0, 0, 0, 0, 0];
|
||||
u32.set(uploadFields(materialGrids.temperature), 8);
|
||||
u32.set(uploadFields(materialGrids.color), 16);
|
||||
u32.set(uploadFields(materialGrids.emission), 24);
|
||||
i32.set([...gridDefinition.indexBounds.min, 0], 32);
|
||||
i32.set([...gridDefinition.indexBounds.max, 0], 36);
|
||||
f32.set([material.densityScale, material.emissionScale, material.anisotropy, Math.max(0.01, gridDefinition.voxelSize[2])], 40);
|
||||
f32.set([...(material.color ?? [0.72, 0.78, 0.86]), 1], 44);
|
||||
f32.set([...(material.emissionColor ?? [1, 1, 1]), 1], 48);
|
||||
f32.set([materialGrids.temperature ? 1 : 0, materialGrids.color ? 1 : 0, materialGrids.emission ? 1 : 0, material.temperatureScale], 52);
|
||||
const params = device.createBuffer({ size: paramsData.byteLength, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST });
|
||||
device.queue.writeBuffer(params, 0, paramsData);
|
||||
const temperatureSource = materialGrids.temperature ? temperatureTraversalWGSL : /* wgsl */`
|
||||
fn temperature_sample_density(coord: vec3<i32>) -> vec2<f32> { return vec2<f32>(0.0); }
|
||||
fn temperature_sample_density_linear(position: vec3<f32>) -> vec2<f32> { return vec2<f32>(0.0); }
|
||||
`;
|
||||
const colorSource = materialGrids.color ? vec3TraversalWGSL : /* wgsl */`
|
||||
fn sample_color(coord: vec3<i32>) -> vec4<f32> { return vec4<f32>(0.0); }
|
||||
fn sample_color_linear(position: vec3<f32>) -> vec4<f32> { return vec4<f32>(0.0); }
|
||||
`;
|
||||
const emissionSource = materialGrids.emission ? emissionTraversalWGSL : /* wgsl */`
|
||||
fn emission_sample_density(coord: vec3<i32>) -> vec2<f32> { return vec2<f32>(0.0); }
|
||||
fn emission_sample_density_linear(position: vec3<f32>) -> vec2<f32> { return vec2<f32>(0.0); }
|
||||
`;
|
||||
const module = device.createShaderModule({ label: "NanoVDB bounded volume integrator", code: /* wgsl */`
|
||||
struct Params {
|
||||
data_bytes: u32, interpolation: u32, width: u32, height: u32,
|
||||
page_bytes: u32, page_count: u32, resident_pages: u32, paged: u32,
|
||||
temperature_data_bytes: u32, temperature_page_bytes: u32, temperature_page_count: u32, temperature_resident_pages: u32,
|
||||
temperature_paged: u32, temperature_pad0: u32, temperature_pad1: u32, temperature_pad2: u32,
|
||||
color_data_bytes: u32, color_page_bytes: u32, color_page_count: u32, color_resident_pages: u32,
|
||||
color_paged: u32, color_pad0: u32, color_pad1: u32, color_pad2: u32,
|
||||
emission_data_bytes: u32, emission_page_bytes: u32, emission_page_count: u32, emission_resident_pages: u32,
|
||||
emission_paged: u32, emission_pad0: u32, emission_pad1: u32, emission_pad2: u32,
|
||||
index_min: vec4<i32>, index_max: vec4<i32>, material: vec4<f32>, color: vec4<f32>, emission_color: vec4<f32>, material_grids: vec4<f32>
|
||||
}
|
||||
@group(0) @binding(0) var<storage, read> grid: array<u32>;
|
||||
@group(0) @binding(1) var<storage, read_write> pixels: array<u32>;
|
||||
@group(0) @binding(2) var<uniform> params: Params;
|
||||
@group(0) @binding(3) var<storage, read> page_table: array<u32>;
|
||||
@group(0) @binding(4) var<storage, read> temperature_grid: array<u32>;
|
||||
@group(0) @binding(5) var<storage, read> color_grid: array<u32>;
|
||||
@group(0) @binding(6) var<storage, read> emission_grid: array<u32>;
|
||||
${traversalWGSL}
|
||||
${temperatureSource}
|
||||
${colorSource}
|
||||
${emissionSource}
|
||||
fn blackbody_color(kelvin: f32) -> vec3<f32> {
|
||||
let t = smoothstep(800.0, 12000.0, clamp(kelvin, 800.0, 12000.0));
|
||||
return mix(vec3<f32>(1.0, 0.11, 0.015), vec3<f32>(0.62, 0.8, 1.0), t);
|
||||
}
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
|
||||
if (id.x >= params.width || id.y >= params.height) { return; }
|
||||
let extent = vec2<f32>(params.index_max.xy - params.index_min.xy + vec2<i32>(1));
|
||||
let uv = (vec2<f32>(id.xy) + vec2<f32>(0.5)) / vec2<f32>(f32(params.width), f32(params.height));
|
||||
let xy_position = vec2<f32>(params.index_min.xy) + uv * extent - vec2<f32>(0.5);
|
||||
let xy = vec2<i32>(round(xy_position));
|
||||
let z_count = max(1, params.index_max.z - params.index_min.z + 1);
|
||||
let stride = max(1, (z_count + 255) / 256);
|
||||
let g = clamp(params.material.z, -0.99, 0.99);
|
||||
let phase = (1.0 - g * g) / (12.5663706 * pow(max(0.0001, 1.0 + g * g), 1.5));
|
||||
var transmittance = 1.0;
|
||||
var radiance = vec3<f32>(0.0);
|
||||
for (var z = params.index_min.z; z <= params.index_max.z; z += stride) {
|
||||
var sample = sample_density(vec3<i32>(xy, z));
|
||||
if (params.interpolation == 1u) {
|
||||
sample = sample_density_linear(vec3<f32>(xy_position, f32(z) + 0.5));
|
||||
}
|
||||
if (sample.y < 0.0) { radiance = vec3<f32>(1.0, 0.0, 1.0); transmittance = 0.0; break; }
|
||||
let density = max(0.0, sample.x) * params.material.x;
|
||||
let alpha = 1.0 - exp(-density * params.material.w * f32(stride));
|
||||
var scattering_color = params.color.rgb;
|
||||
if (params.material_grids.y > 0.5) {
|
||||
var color_sample = sample_color(vec3<i32>(xy, z));
|
||||
if (params.interpolation == 1u) { color_sample = sample_color_linear(vec3<f32>(xy_position, f32(z) + 0.5)); }
|
||||
if (color_sample.w >= 0.0 && color_sample.w > 0.5) { scattering_color = max(vec3<f32>(0.0), color_sample.xyz); }
|
||||
}
|
||||
var emitted = params.emission_color.rgb * params.material.y;
|
||||
if (params.material_grids.x > 0.5 && params.material.y > 0.0) {
|
||||
var temperature_sample = temperature_sample_density(vec3<i32>(xy, z));
|
||||
if (params.interpolation == 1u) { temperature_sample = temperature_sample_density_linear(vec3<f32>(xy_position, f32(z) + 0.5)); }
|
||||
if (temperature_sample.y > 0.5) { emitted += blackbody_color(temperature_sample.x * params.material_grids.w) * params.material.y; }
|
||||
}
|
||||
if (params.material_grids.z > 0.5 && params.material.y > 0.0) {
|
||||
var emission_sample = emission_sample_density(vec3<i32>(xy, z));
|
||||
if (params.interpolation == 1u) { emission_sample = emission_sample_density_linear(vec3<f32>(xy_position, f32(z) + 0.5)); }
|
||||
if (emission_sample.y > 0.5) { emitted += params.emission_color.rgb * max(0.0, emission_sample.x) * params.material.y; }
|
||||
}
|
||||
let source = scattering_color * (0.5 + 8.0 * phase) + emitted;
|
||||
radiance += transmittance * alpha * source;
|
||||
transmittance *= 1.0 - alpha;
|
||||
if (transmittance < 0.005) { break; }
|
||||
}
|
||||
pixels[id.y * params.width + id.x] = pack4x8unorm(vec4<f32>(clamp(radiance, vec3<f32>(0.0), vec3<f32>(1.0)), 1.0 - transmittance));
|
||||
}` });
|
||||
const pipeline = device.createComputePipeline({ layout: "auto", compute: { module, entryPoint: "main" } });
|
||||
const entries: Array<{ binding: number; resource: { buffer: GPUBuffer } }> = [
|
||||
{ binding: 0, resource: { buffer: uploaded.buffer } }, { binding: 1, resource: { buffer: output } }, { binding: 2, resource: { buffer: params } },
|
||||
{ binding: 3, resource: { buffer: uploaded.pageTable } },
|
||||
];
|
||||
if (materialGrids.temperature) entries.push({ binding: 4, resource: { buffer: materialGrids.temperature.buffer } });
|
||||
if (materialGrids.color) entries.push({ binding: 5, resource: { buffer: materialGrids.color.buffer } });
|
||||
if (materialGrids.emission) entries.push({ binding: 6, resource: { buffer: materialGrids.emission.buffer } });
|
||||
const bindGroup = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries });
|
||||
const encoder = device.createCommandEncoder();
|
||||
const pass = encoder.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, bindGroup); pass.dispatchWorkgroups(Math.ceil(width / 8), Math.ceil(height / 8)); pass.end();
|
||||
encoder.copyBufferToBuffer(output, 0, readback, 0, outputBytes);
|
||||
device.queue.submit([encoder.finish()]);
|
||||
await readback.mapAsync(GPUMapMode.READ);
|
||||
const pixels = new Uint8Array(readback.getMappedRange().slice(0));
|
||||
readback.unmap(); output.destroy(); readback.destroy(); params.destroy();
|
||||
return pixels;
|
||||
}
|
||||
Reference in New Issue
Block a user