1014 lines
49 KiB
TypeScript
1014 lines
49 KiB
TypeScript
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;
|
|
reason?: string;
|
|
maxStorageBufferBindingSize?: number;
|
|
maxBufferSize?: number;
|
|
maxStorageBuffersPerShaderStage?: number;
|
|
}
|
|
|
|
export interface NanoVDBWebGPUGrid {
|
|
buffer: GPUBuffer;
|
|
pageTable: GPUBuffer;
|
|
byteLength: number;
|
|
pageByteLength: number;
|
|
pageCount: number;
|
|
residentPageCount: number;
|
|
residentPageCapacity: number;
|
|
residentBytes: number;
|
|
maxResidentBytes: number;
|
|
evictionCount: number;
|
|
paged: boolean;
|
|
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;
|
|
}
|
|
|
|
export type NanoVDBViewAxis = "X" | "Y" | "Z";
|
|
|
|
export interface NanoVDBMaterialGridUploadsIR {
|
|
temperature?: NanoVDBWebGPUGrid;
|
|
color?: NanoVDBWebGPUGrid;
|
|
emission?: NanoVDBWebGPUGrid;
|
|
}
|
|
|
|
export interface NanoVDBGpuPageAllocatorStatsIR {
|
|
residentBytes: number;
|
|
maxResidentBytes: number;
|
|
residentPages: number;
|
|
evictions: number;
|
|
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 */`
|
|
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) { /* 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];
|
|
}
|
|
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);
|
|
}
|
|
`;
|
|
|
|
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}[`)
|
|
.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_page_table", "temperature");
|
|
const emissionTraversalWGSL = specializeFloatTraversal("emission", "emission_grid", "emission_page_table", "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)) { return 0u; }
|
|
if (params.color_paged == 0u) { return color_grid[byte_offset >> 2u]; }
|
|
if (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 = color_page_table[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, requiredStorageBuffers = 5): 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);
|
|
const maxStorageBuffersPerShaderStage = Number(adapter.limits.maxStorageBuffersPerShaderStage);
|
|
if (requiredBytes > maxStorageBufferBindingSize || requiredBytes > maxBufferSize || requiredStorageBuffers > maxStorageBuffersPerShaderStage) {
|
|
return { capability: { available: false, reason: "NanoVDB grid or material bindings exceed WebGPU adapter limits", maxStorageBufferBindingSize, maxBufferSize, maxStorageBuffersPerShaderStage } };
|
|
}
|
|
const device = await adapter.requestDevice({ requiredLimits: { maxStorageBufferBindingSize: requiredBytes, maxBufferSize: requiredBytes, maxStorageBuffersPerShaderStage: requiredStorageBuffers } });
|
|
return { capability: { available: true, maxStorageBufferBindingSize, maxBufferSize, maxStorageBuffersPerShaderStage }, 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 | GPUBufferUsage.COPY_SRC, mappedAtCreation: true });
|
|
new Uint32Array(pageTable.getMappedRange())[0] = 0;
|
|
pageTable.unmap();
|
|
let disposed = false;
|
|
return {
|
|
buffer,
|
|
pageTable,
|
|
byteLength: payload.byteLength,
|
|
pageByteLength: payload.byteLength,
|
|
pageCount: 1,
|
|
residentPageCount: 1,
|
|
residentPageCapacity: 1,
|
|
residentBytes: payload.byteLength,
|
|
maxResidentBytes: payload.byteLength,
|
|
evictionCount: 0,
|
|
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");
|
|
},
|
|
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: () => {
|
|
if (disposed) return;
|
|
disposed = true;
|
|
buffer.destroy();
|
|
pageTable.destroy();
|
|
},
|
|
};
|
|
}
|
|
|
|
export function uploadNanoVDBFloat32GridPaged(
|
|
device: GPUDevice,
|
|
payload: ArrayBuffer,
|
|
pageByteLength: number,
|
|
maxResidentBytes: number,
|
|
): NanoVDBWebGPUGrid {
|
|
if (payload.byteLength === 0 || payload.byteLength % 32 !== 0) {
|
|
throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: invalid paged grid layout");
|
|
}
|
|
const grid = createNanoVDBFloat32GridPaged(device, payload.byteLength, pageByteLength, maxResidentBytes);
|
|
const initialPages = Math.min(grid.pageCount, grid.residentPageCapacity);
|
|
for (let page = 0; page < initialPages; page++) {
|
|
grid.uploadPage(page, payload.slice(page * pageByteLength, Math.min(payload.byteLength, (page + 1) * pageByteLength)));
|
|
}
|
|
return grid;
|
|
}
|
|
|
|
export function createNanoVDBFloat32GridPaged(
|
|
device: GPUDevice,
|
|
byteLength: number,
|
|
pageByteLength: number,
|
|
maxResidentBytes: number,
|
|
): NanoVDBWebGPUGrid {
|
|
if (!Number.isSafeInteger(byteLength) || byteLength <= 0 || 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(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");
|
|
}
|
|
let buffer: GPUBuffer | undefined;
|
|
let pageTable: GPUBuffer | undefined;
|
|
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 | GPUBufferUsage.COPY_SRC, mappedAtCreation: true });
|
|
new Uint32Array(pageTable.getMappedRange()).fill(0xffffffff);
|
|
pageTable.unmap();
|
|
}
|
|
catch (error) {
|
|
pageTable?.destroy();
|
|
buffer?.destroy();
|
|
throw error;
|
|
}
|
|
const resident = new Map<number, number>();
|
|
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]));
|
|
};
|
|
const residentHasSlot = (slot: number): boolean => {
|
|
for (const current of resident.values()) if (current === slot) return true;
|
|
return false;
|
|
};
|
|
const evict = (pageIndex: number, replacement = false): void => {
|
|
if (!resident.delete(pageIndex)) return;
|
|
lastUsed.delete(pageIndex);
|
|
writePageTable(pageIndex, 0xffffffff);
|
|
if (replacement) evictions++;
|
|
};
|
|
const upload = (pageIndex: number, data?: ArrayBuffer): void => {
|
|
const expectedBytes = pageIndex === pageCount - 1 ? byteLength - pageIndex * pageByteLength : pageByteLength;
|
|
if (!Number.isSafeInteger(pageIndex) || pageIndex < 0 || pageIndex >= pageCount || !(data instanceof ArrayBuffer) || data.byteLength !== expectedBytes) {
|
|
throw new Error("NANOVDB_STREAM_INCOMPLETE: NanoVDB page is outside the virtual grid");
|
|
}
|
|
const existingSlot = resident.get(pageIndex);
|
|
let slot = existingSlot ?? [...Array(residentPageCount).keys()].find((candidate) => !residentHasSlot(candidate));
|
|
if (slot === undefined) {
|
|
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);
|
|
}
|
|
if (slot === undefined) throw new Error("NANOVDB_GPU_BUDGET_EXCEEDED: evicted NanoVDB page lost its physical slot");
|
|
device.queue.writeBuffer(buffer, slot * pageByteLength, data);
|
|
writePageTable(pageIndex, slot);
|
|
resident.set(pageIndex, slot);
|
|
lastUsed.set(pageIndex, ++clock);
|
|
};
|
|
return {
|
|
buffer,
|
|
pageTable,
|
|
byteLength,
|
|
pageByteLength,
|
|
pageCount,
|
|
get residentPageCount() { return resident.size; },
|
|
residentPageCapacity: residentPageCount,
|
|
get residentBytes() { return resident.size * pageByteLength; },
|
|
maxResidentBytes: residentPageCount * pageByteLength,
|
|
get evictionCount() { return evictions; },
|
|
paged: true,
|
|
get residentVirtualPages() { return [...resident.keys()].sort((a, b) => a - b); },
|
|
uploadPage: upload,
|
|
touchPage: (pageIndex) => {
|
|
if (!resident.has(pageIndex)) return false;
|
|
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: () => {
|
|
if (disposed) return;
|
|
disposed = true;
|
|
frameActive = false;
|
|
framePins.clear();
|
|
resident.clear();
|
|
lastUsed.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, requiredStorageBuffers = 5): Promise<GPUDevice> {
|
|
if (this.status === "disposed") throw new Error("VOLUME_SHADER_UNAVAILABLE: WebGPU session is disposed");
|
|
const probe = await probeNanoVDBWebGPU(requiredBytes, requiredStorageBuffers);
|
|
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, requiredStorageBuffers = 5): Promise<GPUDevice> {
|
|
this.device?.destroy();
|
|
return this.open(requiredBytes, requiredStorageBuffers);
|
|
}
|
|
|
|
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 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");
|
|
}
|
|
const offsets = new Uint32Array(byteOffsets);
|
|
const offsetBuffer = device.createBuffer({ size: offsets.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
|
|
device.queue.writeBuffer(offsetBuffer, 0, offsets);
|
|
const resultBytes = byteOffsets.length * 4;
|
|
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>;
|
|
@group(0) @binding(1) var<storage, read> offsets: array<u32>;
|
|
@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>;
|
|
${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]); }
|
|
}` });
|
|
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: 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();
|
|
pass.setPipeline(pipeline); pass.setBindGroup(0, bindGroup); pass.dispatchWorkgroups(Math.ceil(byteOffsets.length / 64)); pass.end();
|
|
encoder.copyBufferToBuffer(resultBuffer, 0, readback, 0, resultBytes);
|
|
device.queue.submit([encoder.finish()]);
|
|
await readback.mapAsync(GPUMapMode.READ);
|
|
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]>,
|
|
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));
|
|
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 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>;
|
|
@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>;
|
|
${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; }
|
|
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 } },
|
|
{ binding: 5, resource: { buffer: feedback.buffer } },
|
|
] });
|
|
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();
|
|
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 }));
|
|
}
|
|
|
|
export async function renderNanoVDBFloat32WebGPU(
|
|
device: GPUDevice,
|
|
uploaded: NanoVDBWebGPUGrid,
|
|
gridDefinition: NanoVDBGridIR,
|
|
material: NanoVDBMaterialIR,
|
|
width = 96,
|
|
height = 96,
|
|
materialGrids: NanoVDBMaterialGridUploadsIR = {},
|
|
viewAxis: NanoVDBViewAxis = "Z",
|
|
): 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[13] = viewAxis === "X" ? 0 : viewAxis === "Y" ? 1 : 2;
|
|
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);
|
|
const rayAxis = viewAxis === "X" ? 0 : viewAxis === "Y" ? 1 : 2;
|
|
f32.set([material.densityScale, material.emissionScale, material.anisotropy, Math.max(0.01, gridDefinition.voxelSize[rayAxis])], 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, view_axis: 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>;
|
|
@group(0) @binding(7) var<storage, read> temperature_page_table: array<u32>;
|
|
@group(0) @binding(8) var<storage, read> color_page_table: array<u32>;
|
|
@group(0) @binding(9) var<storage, read> emission_page_table: 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; }
|
|
var plane_min = params.index_min.xy;
|
|
var plane_max = params.index_max.xy;
|
|
var ray_min = params.index_min.z;
|
|
var ray_max = params.index_max.z;
|
|
if (params.view_axis == 0u) {
|
|
plane_min = params.index_min.yz; plane_max = params.index_max.yz;
|
|
ray_min = params.index_min.x; ray_max = params.index_max.x;
|
|
} else if (params.view_axis == 1u) {
|
|
plane_min = params.index_min.xz; plane_max = params.index_max.xz;
|
|
ray_min = params.index_min.y; ray_max = params.index_max.y;
|
|
}
|
|
let extent = vec2<f32>(plane_max - plane_min + vec2<i32>(1));
|
|
let uv = (vec2<f32>(id.xy) + vec2<f32>(0.5)) / vec2<f32>(f32(params.width), f32(params.height));
|
|
let plane_position = vec2<f32>(plane_min) + uv * extent - vec2<f32>(0.5);
|
|
let plane = vec2<i32>(round(plane_position));
|
|
let ray_count = max(1, ray_max - ray_min + 1);
|
|
let stride = max(1, (ray_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 ray = ray_min; ray <= ray_max; ray += stride) {
|
|
var coord = vec3<i32>(plane, ray);
|
|
var linear_coord = vec3<f32>(plane_position, f32(ray) + 0.5);
|
|
if (params.view_axis == 0u) {
|
|
coord = vec3<i32>(ray, plane.x, plane.y);
|
|
linear_coord = vec3<f32>(f32(ray) + 0.5, plane_position.x, plane_position.y);
|
|
} else if (params.view_axis == 1u) {
|
|
coord = vec3<i32>(plane.x, ray, plane.y);
|
|
linear_coord = vec3<f32>(plane_position.x, f32(ray) + 0.5, plane_position.y);
|
|
}
|
|
var sample = sample_density(coord);
|
|
if (params.interpolation == 1u) {
|
|
sample = sample_density_linear(linear_coord);
|
|
}
|
|
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(coord);
|
|
if (params.interpolation == 1u) { color_sample = sample_color_linear(linear_coord); }
|
|
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(coord);
|
|
if (params.interpolation == 1u) { temperature_sample = temperature_sample_density_linear(linear_coord); }
|
|
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(coord);
|
|
if (params.interpolation == 1u) { emission_sample = emission_sample_density_linear(linear_coord); }
|
|
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 } },
|
|
{ binding: 7, resource: { buffer: materialGrids.temperature.pageTable } },
|
|
);
|
|
if (materialGrids.color) entries.push(
|
|
{ binding: 5, resource: { buffer: materialGrids.color.buffer } },
|
|
{ binding: 8, resource: { buffer: materialGrids.color.pageTable } },
|
|
);
|
|
if (materialGrids.emission) entries.push(
|
|
{ binding: 6, resource: { buffer: materialGrids.emission.buffer } },
|
|
{ binding: 9, resource: { buffer: materialGrids.emission.pageTable } },
|
|
);
|
|
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;
|
|
}
|